| // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 |
| // REQUIRES: amdgpu-registered-target |
| // RUN: %clang_cc1 -triple amdgpu9.42-amd-amdhsa -emit-llvm -fcuda-is-device %s -o - | FileCheck %s |
| // RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm -fcuda-is-device %s -o - | FileCheck --check-prefix=SPIRV %s |
| |
| #define __device__ __attribute__((device)) |
| |
| __device__ float global_float; |
| __device__ double global_double; |
| |
| // CHECK-LABEL: define dso_local void @_Z33test_global_atomic_fadd_f32_validPff( |
| // CHECK-SAME: ptr noundef [[PTR:%.*]], float noundef [[VAL:%.*]]) #[[ATTR0:[0-9]+]] { |
| // CHECK-NEXT: [[ENTRY:.*:]] |
| // CHECK-NEXT: [[PTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) |
| // CHECK-NEXT: [[VAL_ADDR:%.*]] = alloca float, align 4, addrspace(5) |
| // CHECK-NEXT: [[RESULT:%.*]] = alloca float, align 4, addrspace(5) |
| // CHECK-NEXT: [[PTR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[PTR_ADDR]] to ptr |
| // CHECK-NEXT: [[VAL_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VAL_ADDR]] to ptr |
| // CHECK-NEXT: [[RESULT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RESULT]] to ptr |
| // CHECK-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR_ASCAST]], align 8 |
| // CHECK-NEXT: store float [[VAL]], ptr [[VAL_ADDR_ASCAST]], align 4 |
| // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR_ASCAST]], align 8 |
| // CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(1) |
| // CHECK-NEXT: [[TMP2:%.*]] = load float, ptr [[VAL_ADDR_ASCAST]], align 4 |
| // CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fadd ptr addrspace(1) [[TMP1]], float [[TMP2]] syncscope("agent") monotonic, align 4, !atomic.ignore.denormal.mode [[META3:![0-9]+]], !amdgpu.no.fine.grained.memory [[META3]] |
| // CHECK-NEXT: store float [[TMP3]], ptr [[RESULT_ASCAST]], align 4 |
| // CHECK-NEXT: [[TMP4:%.*]] = load float, ptr [[VAL_ADDR_ASCAST]], align 4 |
| // CHECK-NEXT: [[TMP5:%.*]] = atomicrmw fadd ptr addrspace(1) @global_float, float [[TMP4]] syncscope("agent") monotonic, align 4, !atomic.ignore.denormal.mode [[META3]], !amdgpu.no.fine.grained.memory [[META3]] |
| // CHECK-NEXT: store float [[TMP5]], ptr [[RESULT_ASCAST]], align 4 |
| // CHECK-NEXT: ret void |
| // |
| // SPIRV-LABEL: define spir_func void @_Z33test_global_atomic_fadd_f32_validPff( |
| // SPIRV-SAME: ptr addrspace(4) noundef [[PTR:%.*]], float noundef [[VAL:%.*]]) addrspace(4) #[[ATTR0:[0-9]+]] { |
| // SPIRV-NEXT: [[ENTRY:.*:]] |
| // SPIRV-NEXT: [[PTR_ADDR:%.*]] = alloca ptr addrspace(4), align 8 |
| // SPIRV-NEXT: [[VAL_ADDR:%.*]] = alloca float, align 4 |
| // SPIRV-NEXT: [[RESULT:%.*]] = alloca float, align 4 |
| // SPIRV-NEXT: [[PTR_ADDR_ASCAST:%.*]] = addrspacecast ptr [[PTR_ADDR]] to ptr addrspace(4) |
| // SPIRV-NEXT: [[VAL_ADDR_ASCAST:%.*]] = addrspacecast ptr [[VAL_ADDR]] to ptr addrspace(4) |
| // SPIRV-NEXT: [[RESULT_ASCAST:%.*]] = addrspacecast ptr [[RESULT]] to ptr addrspace(4) |
| // SPIRV-NEXT: store ptr addrspace(4) [[PTR]], ptr addrspace(4) [[PTR_ADDR_ASCAST]], align 8 |
| // SPIRV-NEXT: store float [[VAL]], ptr addrspace(4) [[VAL_ADDR_ASCAST]], align 4 |
| // SPIRV-NEXT: [[TMP0:%.*]] = call addrspace(4) i1 @llvm.spv.named.boolean.spec.constant(i32 -1, i1 false, metadata [[META4:![0-9]+]]) |
| // SPIRV-NEXT: br i1 [[TMP0]], label %[[IF_THEN:.*]], label %[[IF_END:.*]] |
| // SPIRV: [[IF_THEN]]: |
| // SPIRV-NEXT: [[TMP1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[PTR_ADDR_ASCAST]], align 8 |
| // SPIRV-NEXT: [[TMP2:%.*]] = addrspacecast ptr addrspace(4) [[TMP1]] to ptr addrspace(1) |
| // SPIRV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(4) [[VAL_ADDR_ASCAST]], align 4 |
| // SPIRV-NEXT: [[TMP4:%.*]] = atomicrmw fadd ptr addrspace(1) [[TMP2]], float [[TMP3]] syncscope("device") monotonic, align 4, !atomic.ignore.denormal.mode [[META5:![0-9]+]], !amdgpu.no.fine.grained.memory [[META5]] |
| // SPIRV-NEXT: store float [[TMP4]], ptr addrspace(4) [[RESULT_ASCAST]], align 4 |
| // SPIRV-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(4) [[VAL_ADDR_ASCAST]], align 4 |
| // SPIRV-NEXT: [[TMP6:%.*]] = atomicrmw fadd ptr addrspace(1) @global_float, float [[TMP5]] syncscope("device") monotonic, align 4, !atomic.ignore.denormal.mode [[META5]], !amdgpu.no.fine.grained.memory [[META5]] |
| // SPIRV-NEXT: store float [[TMP6]], ptr addrspace(4) [[RESULT_ASCAST]], align 4 |
| // SPIRV-NEXT: br label %[[IF_END]] |
| // SPIRV: [[IF_END]]: |
| // SPIRV-NEXT: ret void |
| // |
| __device__ void test_global_atomic_fadd_f32_valid(float *ptr, float val) { |
| float result; |
| if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_global_atomic_fadd_f32)) { |
| result = __builtin_amdgcn_global_atomic_fadd_f32(ptr, val); |
| result = __builtin_amdgcn_global_atomic_fadd_f32(&global_float, val); |
| } |
| } |
| |
| // CHECK-LABEL: define dso_local void @_Z33test_global_atomic_fadd_f64_validPdd( |
| // CHECK-SAME: ptr noundef [[PTR:%.*]], double noundef [[VAL:%.*]]) #[[ATTR0]] { |
| // CHECK-NEXT: [[ENTRY:.*:]] |
| // CHECK-NEXT: [[PTR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5) |
| // CHECK-NEXT: [[VAL_ADDR:%.*]] = alloca double, align 8, addrspace(5) |
| // CHECK-NEXT: [[RESULT:%.*]] = alloca double, align 8, addrspace(5) |
| // CHECK-NEXT: [[PTR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[PTR_ADDR]] to ptr |
| // CHECK-NEXT: [[VAL_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VAL_ADDR]] to ptr |
| // CHECK-NEXT: [[RESULT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RESULT]] to ptr |
| // CHECK-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR_ASCAST]], align 8 |
| // CHECK-NEXT: store double [[VAL]], ptr [[VAL_ADDR_ASCAST]], align 8 |
| // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR_ASCAST]], align 8 |
| // CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(1) |
| // CHECK-NEXT: [[TMP2:%.*]] = load double, ptr [[VAL_ADDR_ASCAST]], align 8 |
| // CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fadd ptr addrspace(1) [[TMP1]], double [[TMP2]] syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory [[META3]] |
| // CHECK-NEXT: store double [[TMP3]], ptr [[RESULT_ASCAST]], align 8 |
| // CHECK-NEXT: [[TMP4:%.*]] = load double, ptr [[VAL_ADDR_ASCAST]], align 8 |
| // CHECK-NEXT: [[TMP5:%.*]] = atomicrmw fadd ptr addrspace(1) @global_double, double [[TMP4]] syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory [[META3]] |
| // CHECK-NEXT: store double [[TMP5]], ptr [[RESULT_ASCAST]], align 8 |
| // CHECK-NEXT: ret void |
| // |
| // SPIRV-LABEL: define spir_func void @_Z33test_global_atomic_fadd_f64_validPdd( |
| // SPIRV-SAME: ptr addrspace(4) noundef [[PTR:%.*]], double noundef [[VAL:%.*]]) addrspace(4) #[[ATTR0]] { |
| // SPIRV-NEXT: [[ENTRY:.*:]] |
| // SPIRV-NEXT: [[PTR_ADDR:%.*]] = alloca ptr addrspace(4), align 8 |
| // SPIRV-NEXT: [[VAL_ADDR:%.*]] = alloca double, align 8 |
| // SPIRV-NEXT: [[RESULT:%.*]] = alloca double, align 8 |
| // SPIRV-NEXT: [[PTR_ADDR_ASCAST:%.*]] = addrspacecast ptr [[PTR_ADDR]] to ptr addrspace(4) |
| // SPIRV-NEXT: [[VAL_ADDR_ASCAST:%.*]] = addrspacecast ptr [[VAL_ADDR]] to ptr addrspace(4) |
| // SPIRV-NEXT: [[RESULT_ASCAST:%.*]] = addrspacecast ptr [[RESULT]] to ptr addrspace(4) |
| // SPIRV-NEXT: store ptr addrspace(4) [[PTR]], ptr addrspace(4) [[PTR_ADDR_ASCAST]], align 8 |
| // SPIRV-NEXT: store double [[VAL]], ptr addrspace(4) [[VAL_ADDR_ASCAST]], align 8 |
| // SPIRV-NEXT: [[TMP0:%.*]] = call addrspace(4) i1 @llvm.spv.named.boolean.spec.constant(i32 -1, i1 false, metadata [[META6:![0-9]+]]) |
| // SPIRV-NEXT: br i1 [[TMP0]], label %[[IF_THEN:.*]], label %[[IF_END:.*]] |
| // SPIRV: [[IF_THEN]]: |
| // SPIRV-NEXT: [[TMP1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[PTR_ADDR_ASCAST]], align 8 |
| // SPIRV-NEXT: [[TMP2:%.*]] = addrspacecast ptr addrspace(4) [[TMP1]] to ptr addrspace(1) |
| // SPIRV-NEXT: [[TMP3:%.*]] = load double, ptr addrspace(4) [[VAL_ADDR_ASCAST]], align 8 |
| // SPIRV-NEXT: [[TMP4:%.*]] = atomicrmw fadd ptr addrspace(1) [[TMP2]], double [[TMP3]] syncscope("device") monotonic, align 8, !amdgpu.no.fine.grained.memory [[META5]] |
| // SPIRV-NEXT: store double [[TMP4]], ptr addrspace(4) [[RESULT_ASCAST]], align 8 |
| // SPIRV-NEXT: [[TMP5:%.*]] = load double, ptr addrspace(4) [[VAL_ADDR_ASCAST]], align 8 |
| // SPIRV-NEXT: [[TMP6:%.*]] = atomicrmw fadd ptr addrspace(1) @global_double, double [[TMP5]] syncscope("device") monotonic, align 8, !amdgpu.no.fine.grained.memory [[META5]] |
| // SPIRV-NEXT: store double [[TMP6]], ptr addrspace(4) [[RESULT_ASCAST]], align 8 |
| // SPIRV-NEXT: br label %[[IF_END]] |
| // SPIRV: [[IF_END]]: |
| // SPIRV-NEXT: ret void |
| // |
| __device__ void test_global_atomic_fadd_f64_valid(double *ptr, double val) { |
| double result; |
| if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_global_atomic_fadd_f64)) { |
| result = __builtin_amdgcn_global_atomic_fadd_f64(ptr, val); |
| result = __builtin_amdgcn_global_atomic_fadd_f64(&global_double, val); |
| } |
| } |
| //. |
| // CHECK: [[META3]] = !{} |
| //. |
| // SPIRV: [[META4]] = !{!"has.atomic-fadd-rtn-insts"} |
| // SPIRV: [[META5]] = !{} |
| // SPIRV: [[META6]] = !{!"has.gfx90a-insts"} |
| //. |
