238 lines · plain
1// RUN: %clang_cc1 -x hip %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \2// RUN: -fcuda-is-device -target-cpu gfx906 -fnative-half-type \3// RUN: -fnative-half-arguments-and-returns | FileCheck -check-prefixes=FUN,CHECK,SAFEIR %s4 5// RUN: %clang_cc1 -x hip %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \6// RUN: -fcuda-is-device -target-cpu gfx906 -fnative-half-type \7// RUN: -fnative-half-arguments-and-returns -munsafe-fp-atomics | FileCheck -check-prefixes=FUN,CHECK,UNSAFEIR %s8 9// RUN: %clang_cc1 -x hip %s -O3 -S -o - -triple=amdgcn-amd-amdhsa \10// RUN: -fcuda-is-device -target-cpu gfx1100 -fnative-half-type \11// RUN: -fnative-half-arguments-and-returns | FileCheck -check-prefixes=FUN,SAFE %s12 13// RUN: %clang_cc1 -x hip %s -O3 -S -o - -triple=amdgcn-amd-amdhsa \14// RUN: -fcuda-is-device -target-cpu gfx942 -fnative-half-type \15// RUN: -fnative-half-arguments-and-returns -munsafe-fp-atomics \16// RUN: | FileCheck -check-prefixes=FUN,UNSAFE %s17 18// REQUIRES: amdgpu-registered-target19 20#include "Inputs/cuda.h"21#include <stdatomic.h>22 23__global__ void ffp1(float *p) {24 // FUN-LABEL: @_Z4ffp1Pf25 // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4, [[DEFMD:!amdgpu.no.fine.grained.memory ![0-9]+, !amdgpu.no.remote.memory ![0-9]+$]]26 // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4, [[DEFMD]]27 // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 4, [[DEFMD]]28 // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 4, [[DEFMD]]29 // SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE:[0-9]+]], [[DEFMD]]30 // SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]31 // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]32 // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]33 34 // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4, [[FADDMD:!amdgpu.no.fine.grained.memory ![0-9]+, !amdgpu.no.remote.memory ![0-9]+, !amdgpu.ignore.denormal.mode ![0-9]+$]]35 // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4, [[DEFMD:!amdgpu.no.fine.grained.memory ![0-9]+, !amdgpu.no.remote.memory ![0-9]+$]]36 // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 4, [[DEFMD]]37 // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 4, [[DEFMD]]38 // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE:[0-9]+]], [[FADDMD]]39 // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]40 // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]41 // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]42 43 // SAFE: global_atomic_add_f3244 // SAFE: global_atomic_cmpswap45 // SAFE: global_atomic_max46 // SAFE: global_atomic_min47 // SAFE: global_atomic_max48 // SAFE: global_atomic_min49 50 // UNSAFE: global_atomic_add_f3251 // UNSAFE: global_atomic_cmpswap52 // UNSAFE: global_atomic_cmpswap53 // UNSAFE: global_atomic_cmpswap54 // UNSAFE: global_atomic_cmpswap55 // UNSAFE: global_atomic_cmpswap56 57 __atomic_fetch_add(p, 1.0f, memory_order_relaxed);58 __atomic_fetch_sub(p, 1.0f, memory_order_relaxed);59 __atomic_fetch_max(p, 1.0f, memory_order_relaxed);60 __atomic_fetch_min(p, 1.0f, memory_order_relaxed);61 62 __hip_atomic_fetch_add(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);63 __hip_atomic_fetch_sub(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);64 __hip_atomic_fetch_max(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);65 __hip_atomic_fetch_min(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);66}67 68__global__ void ffp2(double *p) {69 // FUN-LABEL: @_Z4ffp2Pd70 // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, [[DEFMD]]71 // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, [[DEFMD]]72 // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, [[DEFMD]]73 // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, [[DEFMD]]74 // SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]75 // SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]76 // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]77 // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]78 79 // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, [[DEFMD]]80 // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, [[DEFMD]]81 // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, [[DEFMD]]82 // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, [[DEFMD]]83 // UNSAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]84 // UNSAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]85 // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]86 // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]87 88 // SAFE: global_atomic_cmpswap_b6489 // SAFE: global_atomic_cmpswap_b6490 // SAFE: global_atomic_cmpswap_b6491 // SAFE: global_atomic_cmpswap_b6492 // SAFE: global_atomic_cmpswap_b6493 // SAFE: global_atomic_cmpswap_b6494 95 // UNSAFE: global_atomic_add_f6496 // UNSAFE: global_atomic_cmpswap_x297 // UNSAFE: global_atomic_max_f6498 // UNSAFE: global_atomic_min_f6499 // UNSAFE: global_atomic_max_f64100 // UNSAFE: global_atomic_min_f64101 __atomic_fetch_add(p, 1.0, memory_order_relaxed);102 __atomic_fetch_sub(p, 1.0, memory_order_relaxed);103 __atomic_fetch_max(p, 1.0, memory_order_relaxed);104 __atomic_fetch_min(p, 1.0, memory_order_relaxed);105 __hip_atomic_fetch_add(p, 1.0, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);106 __hip_atomic_fetch_sub(p, 1.0, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);107 __hip_atomic_fetch_max(p, 1.0, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);108 __hip_atomic_fetch_min(p, 1.0, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);109}110 111// long double is the same as double for amdgcn.112__global__ void ffp3(long double *p) {113 // FUN-LABEL: @_Z4ffp3Pe114 // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, [[DEFMD]]115 // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, [[DEFMD]]116 // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, [[DEFMD]]117 // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, [[DEFMD]]118 // SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]119 // SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]120 // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]121 // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]122 123 // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, [[DEFMD]]124 // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, [[DEFMD]]125 // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, [[DEFMD]]126 // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, [[DEFMD]]127 // UNSAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]128 // UNSAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]129 // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]130 // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]131 132 // SAFE: global_atomic_cmpswap_b64133 // SAFE: global_atomic_cmpswap_b64134 // SAFE: global_atomic_cmpswap_b64135 // SAFE: global_atomic_cmpswap_b64136 // SAFE: global_atomic_cmpswap_b64137 138 // UNSAFE: global_atomic_cmpswap_x2139 // UNSAFE: global_atomic_max_f64140 // UNSAFE: global_atomic_min_f64141 // UNSAFE: global_atomic_max_f64142 // UNSAFE: global_atomic_min_f64143 __atomic_fetch_add(p, 1.0L, memory_order_relaxed);144 __atomic_fetch_sub(p, 1.0L, memory_order_relaxed);145 __atomic_fetch_max(p, 1.0L, memory_order_relaxed);146 __atomic_fetch_min(p, 1.0L, memory_order_relaxed);147 __hip_atomic_fetch_add(p, 1.0L, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);148 __hip_atomic_fetch_sub(p, 1.0L, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);149 __hip_atomic_fetch_max(p, 1.0L, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);150 __hip_atomic_fetch_min(p, 1.0L, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);151}152 153__device__ double ffp4(double *p, float f) {154 // FUN-LABEL: @_Z4ffp4Pdf155 // CHECK: fpext contract float {{.*}} to double156 // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, [[DEFMD]]157 // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, [[DEFMD]]158 159 // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]160 // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]161 __atomic_fetch_sub(p, f, memory_order_relaxed);162 return __hip_atomic_fetch_sub(p, f, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);163}164 165__device__ double ffp5(double *p, int i) {166 // FUN-LABEL: @_Z4ffp5Pdi167 // CHECK: sitofp i32 {{.*}} to double168 // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, [[DEFMD]]169 // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, [[DEFMD]]170 __atomic_fetch_sub(p, i, memory_order_relaxed);171 172 // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]173 // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]174 return __hip_atomic_fetch_sub(p, i, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);175}176 177__global__ void ffp6(_Float16 *p) {178 // FUN-LABEL: @_Z4ffp6PDF16179 // SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 2, [[DEFMD]]180 // SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 2, [[DEFMD]]181 // SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 2, [[DEFMD]]182 // SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 2, [[DEFMD]]183 // SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]184 // SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]185 // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]186 // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]187 188 // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 2, [[DEFMD]]189 // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 2, [[DEFMD]]190 // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 2, [[DEFMD]]191 // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 2, [[DEFMD]]192 // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]193 // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]194 // UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]195 // UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]]196 197 // SAFE: _Z4ffp6PDF16198 // SAFE: global_atomic_cmpswap199 // SAFE: global_atomic_cmpswap200 // SAFE: global_atomic_cmpswap201 // SAFE: global_atomic_cmpswap202 // SAFE: global_atomic_cmpswap203 // SAFE: global_atomic_cmpswap204 205 // UNSAFE: _Z4ffp6PDF16206 // UNSAFE: global_atomic_cmpswap207 // UNSAFE: global_atomic_cmpswap208 // UNSAFE: global_atomic_cmpswap209 // UNSAFE: global_atomic_cmpswap210 // UNSAFE: global_atomic_cmpswap211 // UNSAFE: global_atomic_cmpswap212 __atomic_fetch_add(p, 1.0, memory_order_relaxed);213 __atomic_fetch_sub(p, 1.0, memory_order_relaxed);214 __atomic_fetch_max(p, 1.0, memory_order_relaxed);215 __atomic_fetch_min(p, 1.0, memory_order_relaxed);216 217 __hip_atomic_fetch_add(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);218 __hip_atomic_fetch_sub(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);219 __hip_atomic_fetch_max(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);220 __hip_atomic_fetch_min(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);221}222 223// CHECK-LABEL: @_Z12test_cmpxchgPiii224// CHECK: cmpxchg ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} acquire acquire, align 4{{$}}225// CHECK: cmpxchg weak ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} acquire acquire, align 4{{$}}226// CHECK: cmpxchg ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} syncscope("workgroup") monotonic monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}227// CHECK: cmpxchg weak ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} syncscope("workgroup") monotonic monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}228__device__ int test_cmpxchg(int *ptr, int cmp, int desired) {229 bool flag = __atomic_compare_exchange(ptr, &cmp, &desired, 0, memory_order_acquire, memory_order_acquire);230 flag = __atomic_compare_exchange_n(ptr, &cmp, desired, 1, memory_order_acquire, memory_order_acquire);231 flag = __hip_atomic_compare_exchange_strong(ptr, &cmp, desired, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WORKGROUP);232 flag = __hip_atomic_compare_exchange_weak(ptr, &cmp, desired, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WORKGROUP);233 return flag;234}235 236// SAFEIR: ![[$NO_PRIVATE]] = !{i32 5, i32 6}237// UNSAFEIR: ![[$NO_PRIVATE]] = !{i32 5, i32 6}238