1// REQUIRES: amdgpu-registered-target 2// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1030 \ 3// RUN: -S -o - %s 4// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1030 \ 5// RUN: -S -o - %s | FileCheck -check-prefix=GFX1030 %s 6 7// CHECK-LABEL: test_ds_addf_local 8// CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %{{.*}}, float %{{.*}}, 9// GFX1030-LABEL: test_ds_addf_local$local 10// GFX1030: ds_add_rtn_f32 11void test_ds_addf_local(__local float *addr, float x){ 12 float *rtn; 13 *rtn = __builtin_amdgcn_ds_atomic_fadd_f32(addr, x); 14} 15