24 lines · plain
1// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \2// RUN: %s -emit-llvm -o - | FileCheck %s3// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \4// RUN: -S -o - %s | FileCheck -check-prefix=GFX8 %s5 6// REQUIRES: amdgpu-registered-target7 8// CHECK-LABEL: test_fadd_local9// CHECK: = atomicrmw fadd ptr addrspace(3) %{{.+}}, float %{{.+}} monotonic, align 410// GFX8-LABEL: test_fadd_local$local:11// GFX8: ds_add_rtn_f32 v{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}12// GFX8: s_endpgm13kernel 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_volatile19// CHECK: = atomicrmw volatile fadd ptr addrspace(3) %{{.+}}, float %{{.+}} monotonic, align 420kernel 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