1 // REQUIRES
: amdgpu-registered-target
2 // RUN
: %clang_cc1 -O0 -cl-std
=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1030 \
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
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
11 void test_ds_addf_local
(__local float
*addr
, float x
){
13 *rtn
= __builtin_amdgcn_ds_atomic_fadd_f32
(addr, x
);