6e0b0038cd
Currently, for AMDGPU, when compiling for OpenCL, we unconditionally use `private` as the default address space. This is wrong for cases where the `generic` address space is available, and is corrected via this patch. In general, this AS map abuse is a bad hack and we should re-work it altogether, but at least after this patch we will stop being incorrect for e.g. OpenCL 2.0.
24 lines
979 B
Common Lisp
24 lines
979 B
Common Lisp
// 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);
|
|
}
|