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