clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl

   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
  13 kernel 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
  20 kernel 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 }