// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
// RUN:   %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
// RUN:   -S -o - %s | FileCheck -check-prefix=GFX8 %s

// REQUIRES: amdgpu-registered-target

// CHECK-LABEL: test_fadd_local
// CHECK: = atomicrmw fadd ptr addrspace(3) %{{.+}}, float %{{.+}} monotonic, align 4
// GFX8-LABEL: test_fadd_local$local:
// GFX8: ds_add_rtn_f32 v{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
// GFX8: s_endpgm
kernel void test_fadd_local(__local float *ptr, float val){
    float *res;
    *res = __builtin_amdgcn_ds_atomic_fadd_f32(ptr, val);
}

// CHECK-LABEL: test_fadd_local_volatile
// CHECK: = atomicrmw volatile fadd ptr addrspace(3) %{{.+}}, float %{{.+}} monotonic, align 4
kernel void test_fadd_local_volatile(volatile __local float *ptr, float val){
    volatile float *res;
    *res = __builtin_amdgcn_ds_atomic_fadd_f32(ptr, val);
}