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

   1 // RUN: %clang_cc1 -no-opaque-pointers -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
   2 // RUN:   %s -S -emit-llvm -o - | FileCheck %s
   3 // RUN: %clang_cc1 -no-opaque-pointers -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: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %{{.*}}, float %{{.*}}, i32 0, i32 0, i1 false)
  10 // GFX8-LABEL: test_fadd_local$local:
  11 // GFX8: ds_add_rtn_f32 v2, v0, v1
  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 }