xref: /llvm-project/clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl (revision 6e0b0038cd65ce726ce404305a06e1cf33e36cca)
1// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
2// RUN:   %s -emit-llvm -o - | FileCheck %s
3// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
4// RUN:   -S -o - %s | FileCheck -check-prefix=GFX8 %s
5
6// REQUIRES: amdgpu-registered-target
7
8// CHECK-LABEL: test_fadd_local
9// CHECK: = atomicrmw fadd ptr addrspace(3) %{{.+}}, float %{{.+}} monotonic, align 4
10// GFX8-LABEL: test_fadd_local$local:
11// GFX8: ds_add_rtn_f32 v{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
12// GFX8: s_endpgm
13kernel void test_fadd_local(__local float *ptr, float val){
14    float *res;
15    *res = __builtin_amdgcn_ds_atomic_fadd_f32(ptr, val);
16}
17
18// CHECK-LABEL: test_fadd_local_volatile
19// CHECK: = atomicrmw volatile fadd ptr addrspace(3) %{{.+}}, float %{{.+}} monotonic, align 4
20kernel void test_fadd_local_volatile(volatile __local float *ptr, float val){
21    volatile float *res;
22    *res = __builtin_amdgcn_ds_atomic_fadd_f32(ptr, val);
23}
24