https://github.com/easyonaadit updated https://github.com/llvm/llvm-project/pull/230044
>From f071a2085f46f78b23fd32a210a00cdef0f3f376 Mon Sep 17 00:00:00 2001 From: Aaditya <[email protected]> Date: Thu, 8 Oct 2026 14:12:18 +0530 Subject: [PATCH] [Clang] Mark scoped_atomics with !noalias.addrspace(private) The HIP specification marks atomics on thread private memory as UB. Scoped atomics used within a HIP context are also considered UB, unless explicitley specified via a command line argument. These are now annotated with !noalias.addrspace(5) for amdgpus, to avoid an expensive runtime check. --- clang/include/clang/AST/Expr.h | 7 +- clang/include/clang/Basic/LangOptions.def | 1 + clang/include/clang/Basic/LangOptions.h | 2 +- clang/include/clang/Options/Options.td | 6 + clang/lib/CodeGen/Targets/AMDGPU.cpp | 8 +- clang/test/CodeGen/scoped-atomic-ops.c | 151 +++++++++--------- clang/test/CodeGenCUDA/atomic-options.hip | 42 ++--- .../CodeGenCUDA/private-atomics-undefined.hip | 27 ++++ 8 files changed, 143 insertions(+), 101 deletions(-) create mode 100644 clang/test/CodeGenCUDA/private-atomics-undefined.hip diff --git a/clang/include/clang/AST/Expr.h b/clang/include/clang/AST/Expr.h index a0324240ce34da..5642f26e0c50b0 100644 --- a/clang/include/clang/AST/Expr.h +++ b/clang/include/clang/AST/Expr.h @@ -7097,10 +7097,9 @@ class AtomicExpr : public Expr { Op <= AO__hip_atomic_store; } - /// Return true if atomics operations targeting allocations in private memory - /// are undefined. - bool threadPrivateMemoryAtomicsAreUndefined() const { - return isOpenCL() || isHIP(); + bool isScopedAtomic() const { + return getOp() >= AO__scoped_atomic_add_fetch && + getOp() <= AO__scoped_atomic_xor_fetch; } SourceLocation getBuiltinLoc() const { return BuiltinLoc; } diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index e0c29f7b0d791f..37db4f6039dabf 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -266,6 +266,7 @@ LANGOPT(GPUDeferDiag, 1, 0, NotCompatible, "defer host/device related diagnostic LANGOPT(GPUExcludeWrongSideOverloads, 1, 0, NotCompatible, "always exclude wrong side overloads in overloading resolution for CUDA/HIP") LANGOPT(OffloadViaLLVM, 1, 0, NotCompatible, "target LLVM/Offload as portable offloading runtime.") LANGOPT(CUDANVCCABI, 1, 0, NotCompatible, "emit the NVCC-compatible host registration ABI for CUDA") +LANGOPT(PrivateAtomicsUndefined, 1, 0, NotCompatible, "atomics on thread private memory are undefined behavior") LANGOPT(SYCLIsDevice , 1, 0, NotCompatible, "Generate code for SYCL device") LANGOPT(SYCLIsHost , 1, 0, NotCompatible, "SYCL host compilation") diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h index 7539e000d03f98..526960bcc44647 100644 --- a/clang/include/clang/Basic/LangOptions.h +++ b/clang/include/clang/Basic/LangOptions.h @@ -715,7 +715,7 @@ class LangOptions : public LangOptionsBase { bool threadPrivateMemoryAtomicsAreUndefined() const { // Should be false for OpenMP. // TODO: Should this be true for SYCL? - return OpenCL || CUDA; + return PrivateAtomicsUndefined; } /// Return the OpenCL C or C++ version as a VersionTuple. diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td index f3e50f90835327..93e27d3ab35e3c 100644 --- a/clang/include/clang/Options/Options.td +++ b/clang/include/clang/Options/Options.td @@ -1154,6 +1154,12 @@ defm offload_uniform_block : BoolFOption<"offload-uniform-block", NegFlag<SetFalse, [], [ClangOption, CC1Option], "Don't assume">, BothFlags<[], [ClangOption], " that kernels are launched with uniform block sizes (default true for CUDA/HIP and false otherwise)">>; +defm private_atomics_undefined : BoolFOption<"private-atomics-undefined", + LangOpts<"PrivateAtomicsUndefined">, Default<"LangOpts->CUDA || LangOpts->OpenCL">, + PosFlag<SetTrue, [], [ClangOption, CC1Option], "Assume">, + NegFlag<SetFalse, [], [ClangOption, CC1Option], "Don't assume">, + BothFlags<[], [ClangOption], " that atomics on thread private memory are undefined behavior (default true for HIP/OpenCL)">>; + def fcomplex_arithmetic_EQ : Joined<["-"], "fcomplex-arithmetic=">, Group<f_Group>, Visibility<[ClangOption, CC1Option, FlangOption, FC1Option]>, Values<"full,improved,promoted,basic">, NormalizedValuesScope<"LangOptions">, diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 0b5ed1898f1389..13dcf7262b7a7c 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -556,13 +556,19 @@ void AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata( // OpenCL and old style HIP atomics consider atomics targeting thread private // memory to be undefined. + // Scoped atomics targeting private memory, when used within a HIP context, + // are considered UB, unless explicitely opted in by the user via a command + // line flag(-f[no-]private-atomics-undefined). // // TODO: This is probably undefined for atomic load/store, but there's not // much direct codegen benefit to knowing this. if (((RMW && RMW->getPointerAddressSpace() == llvm::AMDGPUAS::FLAT_ADDRESS) || (CmpX && CmpX->getPointerAddressSpace() == llvm::AMDGPUAS::FLAT_ADDRESS)) && - AE && AE->threadPrivateMemoryAtomicsAreUndefined()) { + AE && + (AE->isOpenCL() || AE->isHIP() || + (AE->isScopedAtomic() && + CGF.getLangOpts().threadPrivateMemoryAtomicsAreUndefined()))) { llvm::MDBuilder MDHelper(CGF.getLLVMContext()); llvm::MDNode *ASRange = MDHelper.createRange( llvm::APInt(32, llvm::AMDGPUAS::PRIVATE_ADDRESS), diff --git a/clang/test/CodeGen/scoped-atomic-ops.c b/clang/test/CodeGen/scoped-atomic-ops.c index d0feb389256826..2682da9fd25661 100644 --- a/clang/test/CodeGen/scoped-atomic-ops.c +++ b/clang/test/CodeGen/scoped-atomic-ops.c @@ -775,7 +775,7 @@ void fi2b(int *i) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3:![0-9]+]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3:![0-9]+]], !amdgpu.no.fine.grained.memory [[META4:![0-9]+]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 @@ -783,7 +783,7 @@ void fi2b(int *i) { // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP7]], ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP9:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 @@ -791,7 +791,7 @@ void fi2b(int *i) { // AMDGCN_CL_20-NEXT: [[TMP10:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 // AMDGCN_CL_20-NEXT: [[TMP11:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP12]], ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP13:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP14:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 @@ -799,7 +799,7 @@ void fi2b(int *i) { // AMDGCN_CL_20-NEXT: [[TMP15:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 // AMDGCN_CL_20-NEXT: [[TMP16:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP17]], ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP18:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP19:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 @@ -807,7 +807,7 @@ void fi2b(int *i) { // AMDGCN_CL_20-NEXT: [[TMP20:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 // AMDGCN_CL_20-NEXT: [[TMP21:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP22]], ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP23:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP24:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 @@ -815,7 +815,7 @@ void fi2b(int *i) { // AMDGCN_CL_20-NEXT: [[TMP25:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 // AMDGCN_CL_20-NEXT: [[TMP26:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP27]], ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP28:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP29:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 @@ -823,7 +823,7 @@ void fi2b(int *i) { // AMDGCN_CL_20-NEXT: [[TMP30:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 // AMDGCN_CL_20-NEXT: [[TMP31:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP32]], ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP33:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP34:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 @@ -831,7 +831,7 @@ void fi2b(int *i) { // AMDGCN_CL_20-NEXT: [[TMP35:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 // AMDGCN_CL_20-NEXT: [[TMP36:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP37]], ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP38:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP39:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 @@ -1198,7 +1198,7 @@ void fi3a(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("agent") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 @@ -1206,7 +1206,7 @@ void fi3a(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] syncscope("agent") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP7]], ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP9:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 @@ -1214,7 +1214,7 @@ void fi3a(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP10:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 // AMDGCN_CL_20-NEXT: [[TMP11:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] syncscope("agent") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP12]], ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP13:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP14:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 @@ -1222,7 +1222,7 @@ void fi3a(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP15:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 // AMDGCN_CL_20-NEXT: [[TMP16:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] syncscope("agent") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP17]], ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP18:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP19:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 @@ -1230,7 +1230,7 @@ void fi3a(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP20:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 // AMDGCN_CL_20-NEXT: [[TMP21:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] syncscope("agent") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP22]], ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP23:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP24:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 @@ -1238,7 +1238,7 @@ void fi3a(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP25:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 // AMDGCN_CL_20-NEXT: [[TMP26:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] syncscope("agent") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP27]], ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP28:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP29:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 @@ -1246,7 +1246,7 @@ void fi3a(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP30:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 // AMDGCN_CL_20-NEXT: [[TMP31:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] syncscope("agent") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP32]], ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP33:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP34:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 @@ -1254,7 +1254,7 @@ void fi3a(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP35:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 // AMDGCN_CL_20-NEXT: [[TMP36:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] syncscope("agent") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP37]], ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP38:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP39:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 @@ -1621,7 +1621,7 @@ void fi3b(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("workgroup") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("workgroup") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 @@ -1629,7 +1629,7 @@ void fi3b(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] syncscope("workgroup") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] syncscope("workgroup") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP7]], ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP9:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 @@ -1637,7 +1637,7 @@ void fi3b(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP10:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 // AMDGCN_CL_20-NEXT: [[TMP11:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] syncscope("workgroup") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] syncscope("workgroup") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP12]], ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP13:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP14:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 @@ -1645,7 +1645,7 @@ void fi3b(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP15:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 // AMDGCN_CL_20-NEXT: [[TMP16:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] syncscope("workgroup") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] syncscope("workgroup") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP17]], ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP18:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP19:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 @@ -1653,7 +1653,7 @@ void fi3b(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP20:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 // AMDGCN_CL_20-NEXT: [[TMP21:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] syncscope("workgroup") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] syncscope("workgroup") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP22]], ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP23:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP24:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 @@ -1661,7 +1661,7 @@ void fi3b(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP25:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 // AMDGCN_CL_20-NEXT: [[TMP26:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] syncscope("workgroup") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] syncscope("workgroup") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP27]], ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP28:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP29:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 @@ -1669,7 +1669,7 @@ void fi3b(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP30:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 // AMDGCN_CL_20-NEXT: [[TMP31:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] syncscope("workgroup") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] syncscope("workgroup") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP32]], ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP33:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP34:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 @@ -1677,7 +1677,7 @@ void fi3b(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP35:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 // AMDGCN_CL_20-NEXT: [[TMP36:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] syncscope("workgroup") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] syncscope("workgroup") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP37]], ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP38:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP39:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 @@ -2044,7 +2044,7 @@ void fi3c(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("cluster") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("cluster") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 @@ -2052,7 +2052,7 @@ void fi3c(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] syncscope("cluster") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] syncscope("cluster") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP7]], ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP9:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 @@ -2060,7 +2060,7 @@ void fi3c(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP10:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 // AMDGCN_CL_20-NEXT: [[TMP11:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] syncscope("cluster") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] syncscope("cluster") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP12]], ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP13:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP14:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 @@ -2068,7 +2068,7 @@ void fi3c(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP15:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 // AMDGCN_CL_20-NEXT: [[TMP16:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] syncscope("cluster") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] syncscope("cluster") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP17]], ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP18:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP19:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 @@ -2076,7 +2076,7 @@ void fi3c(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP20:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 // AMDGCN_CL_20-NEXT: [[TMP21:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] syncscope("cluster") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] syncscope("cluster") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP22]], ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP23:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP24:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 @@ -2084,7 +2084,7 @@ void fi3c(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP25:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 // AMDGCN_CL_20-NEXT: [[TMP26:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] syncscope("cluster") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] syncscope("cluster") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP27]], ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP28:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP29:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 @@ -2092,7 +2092,7 @@ void fi3c(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP30:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 // AMDGCN_CL_20-NEXT: [[TMP31:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] syncscope("cluster") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] syncscope("cluster") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP32]], ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP33:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP34:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 @@ -2100,7 +2100,7 @@ void fi3c(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP35:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 // AMDGCN_CL_20-NEXT: [[TMP36:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] syncscope("cluster") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] syncscope("cluster") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP37]], ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP38:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP39:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 @@ -2467,7 +2467,7 @@ void fi3_clustr(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("wavefront") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("wavefront") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 @@ -2475,7 +2475,7 @@ void fi3_clustr(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] syncscope("wavefront") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] syncscope("wavefront") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP7]], ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP9:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 @@ -2483,7 +2483,7 @@ void fi3_clustr(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) // AMDGCN_CL_20-NEXT: [[TMP10:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 // AMDGCN_CL_20-NEXT: [[TMP11:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] syncscope("wavefront") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] syncscope("wavefront") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP12]], ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP13:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP14:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 @@ -2491,7 +2491,7 @@ void fi3_clustr(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) // AMDGCN_CL_20-NEXT: [[TMP15:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 // AMDGCN_CL_20-NEXT: [[TMP16:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] syncscope("wavefront") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] syncscope("wavefront") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP17]], ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP18:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP19:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 @@ -2499,7 +2499,7 @@ void fi3_clustr(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) // AMDGCN_CL_20-NEXT: [[TMP20:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 // AMDGCN_CL_20-NEXT: [[TMP21:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] syncscope("wavefront") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] syncscope("wavefront") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP22]], ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP23:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP24:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 @@ -2507,7 +2507,7 @@ void fi3_clustr(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) // AMDGCN_CL_20-NEXT: [[TMP25:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 // AMDGCN_CL_20-NEXT: [[TMP26:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] syncscope("wavefront") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] syncscope("wavefront") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP27]], ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP28:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP29:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 @@ -2515,7 +2515,7 @@ void fi3_clustr(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) // AMDGCN_CL_20-NEXT: [[TMP30:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 // AMDGCN_CL_20-NEXT: [[TMP31:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] syncscope("wavefront") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] syncscope("wavefront") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP32]], ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP33:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP34:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 @@ -2523,7 +2523,7 @@ void fi3_clustr(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) // AMDGCN_CL_20-NEXT: [[TMP35:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 // AMDGCN_CL_20-NEXT: [[TMP36:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] syncscope("wavefront") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] syncscope("wavefront") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP37]], ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP38:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP39:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 @@ -2890,7 +2890,7 @@ void fi3d(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("singlethread") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("singlethread") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 @@ -2898,7 +2898,7 @@ void fi3d(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] syncscope("singlethread") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw sub ptr [[TMP5]], i32 [[TMP6]] syncscope("singlethread") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP7]], ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP9:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 @@ -2906,7 +2906,7 @@ void fi3d(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP10:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 // AMDGCN_CL_20-NEXT: [[TMP11:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] syncscope("singlethread") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP12:%.*]] = atomicrmw and ptr [[TMP10]], i32 [[TMP11]] syncscope("singlethread") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP12]], ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP13:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // AMDGCN_CL_20-NEXT: [[TMP14:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 @@ -2914,7 +2914,7 @@ void fi3d(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP15:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 // AMDGCN_CL_20-NEXT: [[TMP16:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] syncscope("singlethread") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP17:%.*]] = atomicrmw or ptr [[TMP15]], i32 [[TMP16]] syncscope("singlethread") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP17]], ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP18:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // AMDGCN_CL_20-NEXT: [[TMP19:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 @@ -2922,7 +2922,7 @@ void fi3d(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP20:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 // AMDGCN_CL_20-NEXT: [[TMP21:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP7]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] syncscope("singlethread") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP22:%.*]] = atomicrmw xor ptr [[TMP20]], i32 [[TMP21]] syncscope("singlethread") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP22]], ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP23:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP8]], align 4 // AMDGCN_CL_20-NEXT: [[TMP24:%.*]] = load ptr, ptr addrspace(5) [[E_ADDR]], align 8 @@ -2930,7 +2930,7 @@ void fi3d(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP25:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 // AMDGCN_CL_20-NEXT: [[TMP26:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP9]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] syncscope("singlethread") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP27:%.*]] = atomicrmw nand ptr [[TMP25]], i32 [[TMP26]] syncscope("singlethread") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP27]], ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP28:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP10]], align 4 // AMDGCN_CL_20-NEXT: [[TMP29:%.*]] = load ptr, ptr addrspace(5) [[F_ADDR]], align 8 @@ -2938,7 +2938,7 @@ void fi3d(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP30:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 // AMDGCN_CL_20-NEXT: [[TMP31:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP11]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] syncscope("singlethread") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP32:%.*]] = atomicrmw min ptr [[TMP30]], i32 [[TMP31]] syncscope("singlethread") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP32]], ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP33:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP12]], align 4 // AMDGCN_CL_20-NEXT: [[TMP34:%.*]] = load ptr, ptr addrspace(5) [[G_ADDR]], align 8 @@ -2946,7 +2946,7 @@ void fi3d(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP35:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 // AMDGCN_CL_20-NEXT: [[TMP36:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP13]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] syncscope("singlethread") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP37:%.*]] = atomicrmw max ptr [[TMP35]], i32 [[TMP36]] syncscope("singlethread") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP37]], ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP38:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP14]], align 4 // AMDGCN_CL_20-NEXT: [[TMP39:%.*]] = load ptr, ptr addrspace(5) [[H_ADDR]], align 8 @@ -3210,7 +3210,7 @@ void fi3e(int *a, int *b, int *c, int *d, int *e, int *f, int *g, int *h) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[DESIRED]] to ptr // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP4]], 0 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = extractvalue { i32, i1 } [[TMP4]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP6]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -3330,7 +3330,7 @@ _Bool fi4a(int *i) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[DESIRED]] to ptr // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] syncscope("agent") acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] syncscope("agent") acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP4]], 0 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = extractvalue { i32, i1 } [[TMP4]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP6]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -3450,7 +3450,7 @@ _Bool fi4b(int *i) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[DESIRED]] to ptr // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] syncscope("workgroup") acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] syncscope("workgroup") acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP4]], 0 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = extractvalue { i32, i1 } [[TMP4]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP6]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -3570,7 +3570,7 @@ _Bool fi4c(int *i) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[DESIRED]] to ptr // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] syncscope("cluster") acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] syncscope("cluster") acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP4]], 0 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = extractvalue { i32, i1 } [[TMP4]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP6]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -3690,7 +3690,7 @@ _Bool fi4_clustr(int *i) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[DESIRED]] to ptr // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] syncscope("wavefront") acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] syncscope("wavefront") acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP4]], 0 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = extractvalue { i32, i1 } [[TMP4]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP6]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -3810,7 +3810,7 @@ _Bool fi4d(int *i) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[DESIRED]] to ptr // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] syncscope("singlethread") acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[TMP0]], i32 [[TMP2]], i32 [[TMP3]] syncscope("singlethread") acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP4]], 0 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = extractvalue { i32, i1 } [[TMP4]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP6]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -3928,7 +3928,7 @@ _Bool fi4e(int *i) { // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = extractvalue { i32, i1 } [[TMP3]], 0 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP3]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP5]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -4045,7 +4045,7 @@ _Bool fi5a(int *i) { // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] syncscope("agent") acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] syncscope("agent") acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = extractvalue { i32, i1 } [[TMP3]], 0 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP3]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP5]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -4162,7 +4162,7 @@ _Bool fi5b(int *i) { // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] syncscope("workgroup") acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] syncscope("workgroup") acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = extractvalue { i32, i1 } [[TMP3]], 0 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP3]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP5]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -4278,7 +4278,7 @@ _Bool fi5c(int *i) { // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] syncscope("cluster") acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] syncscope("cluster") acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = extractvalue { i32, i1 } [[TMP3]], 0 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP3]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP5]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -4394,7 +4394,7 @@ _Bool fi5_clustr(int *i) { // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] syncscope("wavefront") acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] syncscope("wavefront") acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = extractvalue { i32, i1 } [[TMP3]], 0 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP3]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP5]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -4510,7 +4510,7 @@ _Bool fi5d(int *i) { // AMDGCN_CL_20-NEXT: store i32 1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[CMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] syncscope("singlethread") acquire acquire, align 4 +// AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = cmpxchg weak ptr [[TMP0]], i32 [[TMP1]], i32 [[TMP2]] syncscope("singlethread") acquire acquire, align 4, !noalias.addrspace [[META3]] // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = extractvalue { i32, i1 } [[TMP3]], 0 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = extractvalue { i32, i1 } [[TMP3]], 1 // AMDGCN_CL_20-NEXT: br i1 [[TMP5]], label %[[CMPXCHG_CONTINUE:.*]], label %[[CMPXCHG_STORE_EXPECTED:.*]] @@ -4615,7 +4615,7 @@ _Bool fi5e(int *i) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = addrspacecast ptr addrspace(5) [[RET]] to ptr // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP4]], ptr [[TMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(5) [[RET]], align 4 // AMDGCN_CL_20-NEXT: ret i32 [[TMP5]] @@ -4689,7 +4689,7 @@ int fi6a(int *c, int *d) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = addrspacecast ptr addrspace(5) [[RET]] to ptr // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] syncscope("agent") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP4]], ptr [[TMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(5) [[RET]], align 4 // AMDGCN_CL_20-NEXT: ret i32 [[TMP5]] @@ -4763,7 +4763,7 @@ int fi6b(int *c, int *d) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = addrspacecast ptr addrspace(5) [[RET]] to ptr // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] syncscope("workgroup") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] syncscope("workgroup") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP4]], ptr [[TMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(5) [[RET]], align 4 // AMDGCN_CL_20-NEXT: ret i32 [[TMP5]] @@ -4837,7 +4837,7 @@ int fi6c(int *c, int *d) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = addrspacecast ptr addrspace(5) [[RET]] to ptr // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] syncscope("cluster") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] syncscope("cluster") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP4]], ptr [[TMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(5) [[RET]], align 4 // AMDGCN_CL_20-NEXT: ret i32 [[TMP5]] @@ -4911,7 +4911,7 @@ int fi6_clustr(int *c, int *d) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = addrspacecast ptr addrspace(5) [[RET]] to ptr // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] syncscope("wavefront") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] syncscope("wavefront") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP4]], ptr [[TMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(5) [[RET]], align 4 // AMDGCN_CL_20-NEXT: ret i32 [[TMP5]] @@ -4985,7 +4985,7 @@ int fi6d(int *c, int *d) { // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspace(5) [[D_ADDR]], align 8 // AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = addrspacecast ptr addrspace(5) [[RET]] to ptr // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] syncscope("singlethread") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = atomicrmw xchg ptr [[TMP0]], i32 [[TMP3]] syncscope("singlethread") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP4]], ptr [[TMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(5) [[RET]], align 4 // AMDGCN_CL_20-NEXT: ret i32 [[TMP5]] @@ -5055,7 +5055,7 @@ int fi6e(int *c, int *d) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i8 1, ptr addrspace(5) [[DOTATOMICTMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i8, ptr addrspace(5) [[DOTATOMICTMP]], align 1 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] monotonic, align 1, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] monotonic, align 1, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i8 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i8, ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[LOADEDV:%.*]] = icmp ne i8 [[TMP3]], 0 @@ -5125,7 +5125,7 @@ _Bool fi7a(_Bool *c) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i8 1, ptr addrspace(5) [[DOTATOMICTMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i8, ptr addrspace(5) [[DOTATOMICTMP]], align 1 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] syncscope("agent") monotonic, align 1, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] syncscope("agent") monotonic, align 1, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i8 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i8, ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[LOADEDV:%.*]] = icmp ne i8 [[TMP3]], 0 @@ -5195,7 +5195,7 @@ _Bool fi7b(_Bool *c) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i8 1, ptr addrspace(5) [[DOTATOMICTMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i8, ptr addrspace(5) [[DOTATOMICTMP]], align 1 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] syncscope("workgroup") monotonic, align 1, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] syncscope("workgroup") monotonic, align 1, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i8 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i8, ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[LOADEDV:%.*]] = icmp ne i8 [[TMP3]], 0 @@ -5265,7 +5265,7 @@ _Bool fi7c(_Bool *c) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i8 1, ptr addrspace(5) [[DOTATOMICTMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i8, ptr addrspace(5) [[DOTATOMICTMP]], align 1 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] syncscope("cluster") monotonic, align 1, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] syncscope("cluster") monotonic, align 1, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i8 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i8, ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[LOADEDV:%.*]] = icmp ne i8 [[TMP3]], 0 @@ -5335,7 +5335,7 @@ _Bool fi7_clustr(_Bool *c) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i8 1, ptr addrspace(5) [[DOTATOMICTMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i8, ptr addrspace(5) [[DOTATOMICTMP]], align 1 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] syncscope("wavefront") monotonic, align 1, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] syncscope("wavefront") monotonic, align 1, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i8 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i8, ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[LOADEDV:%.*]] = icmp ne i8 [[TMP3]], 0 @@ -5405,7 +5405,7 @@ _Bool fi7d(_Bool *c) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[C_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i8 1, ptr addrspace(5) [[DOTATOMICTMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i8, ptr addrspace(5) [[DOTATOMICTMP]], align 1 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] syncscope("singlethread") monotonic, align 1, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw xchg ptr [[TMP0]], i8 [[TMP1]] syncscope("singlethread") monotonic, align 1, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i8 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i8, ptr addrspace(5) [[ATOMIC_TEMP]], align 1 // AMDGCN_CL_20-NEXT: [[LOADEDV:%.*]] = icmp ne i8 [[TMP3]], 0 @@ -5493,7 +5493,7 @@ _Bool fi7e(_Bool *c) { // AMDGCN_CL_20-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 -1, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw uinc_wrap ptr [[TMP0]], i32 [[TMP1]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP2:%.*]] = atomicrmw uinc_wrap ptr [[TMP0]], i32 [[TMP1]] syncscope("agent") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // AMDGCN_CL_20-NEXT: [[TMP4:%.*]] = load ptr, ptr addrspace(5) [[B_ADDR]], align 8 @@ -5501,7 +5501,7 @@ _Bool fi7e(_Bool *c) { // AMDGCN_CL_20-NEXT: [[TMP5:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 // AMDGCN_CL_20-NEXT: store i32 -1, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 // AMDGCN_CL_20-NEXT: [[TMP6:%.*]] = load i32, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 -// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw udec_wrap ptr [[TMP5]], i32 [[TMP6]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// AMDGCN_CL_20-NEXT: [[TMP7:%.*]] = atomicrmw udec_wrap ptr [[TMP5]], i32 [[TMP6]] syncscope("agent") monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // AMDGCN_CL_20-NEXT: store i32 [[TMP7]], ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // AMDGCN_CL_20-NEXT: [[TMP9:%.*]] = load ptr, ptr addrspace(5) [[A_ADDR]], align 8 @@ -5573,7 +5573,8 @@ void fi8a(unsigned int *a, unsigned int *b) { //. // AMDGCN_CL_DEF: [[META2]] = !{} //. -// AMDGCN_CL_20: [[META3]] = !{} +// AMDGCN_CL_20: [[META3]] = !{i32 5, i32 6} +// AMDGCN_CL_20: [[META4]] = !{} //. //// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line: // AMDGCN: {{.*}} diff --git a/clang/test/CodeGenCUDA/atomic-options.hip b/clang/test/CodeGenCUDA/atomic-options.hip index 2bbd14c692e7cc..37b81b244f06b9 100644 --- a/clang/test/CodeGenCUDA/atomic-options.hip +++ b/clang/test/CodeGenCUDA/atomic-options.hip @@ -41,7 +41,7 @@ // DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3:![0-9]+]], !amdgpu.no.remote.memory [[META3]] +// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3:![0-9]+]], !amdgpu.no.fine.grained.memory [[META4:![0-9]+]], !amdgpu.no.remote.memory [[META4]] // DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: ret void @@ -57,7 +57,7 @@ // OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META3:![0-9]+]], !amdgpu.no.remote.memory [[META3]] +// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3:![0-9]+]], !atomic.ignore.denormal.mode [[META4:![0-9]+]], !amdgpu.no.remote.memory [[META4]] // OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: ret void @@ -124,7 +124,7 @@ __device__ __host__ void test_default(float *a) { // DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: ret void @@ -140,7 +140,7 @@ __device__ __host__ void test_default(float *a) { // OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META3]], !amdgpu.no.remote.memory [[META3]] +// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3]], !atomic.ignore.denormal.mode [[META4]], !amdgpu.no.remote.memory [[META4]] // OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: ret void @@ -209,7 +209,7 @@ __device__ __host__ void test_one(float *a) { // DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META3]], !amdgpu.no.fine.grained.memory [[META3]] +// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3]], !atomic.ignore.denormal.mode [[META4]], !amdgpu.no.fine.grained.memory [[META4]] // DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: ret void @@ -225,7 +225,7 @@ __device__ __host__ void test_one(float *a) { // OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META3]] +// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3]], !atomic.ignore.denormal.mode [[META4]] // OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: ret void @@ -294,7 +294,7 @@ __device__ __host__ void test_two(float *a) { // DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3]] +// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.remote.memory [[META4]] // DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: ret void @@ -310,7 +310,7 @@ __device__ __host__ void test_two(float *a) { // OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3]] +// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.remote.memory [[META4]] // OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: ret void @@ -379,7 +379,7 @@ __device__ __host__ void test_three(float *a) { // DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]] +// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]] // DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: ret void @@ -395,7 +395,7 @@ __device__ __host__ void test_three(float *a) { // OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META3]] +// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3]], !atomic.ignore.denormal.mode [[META4]] // OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: ret void @@ -494,25 +494,25 @@ __device__ __host__ void test_multiple_attrs(float *a) { // DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]] +// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]] // DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // DEV-NEXT: [[TMP4:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // DEV-NEXT: store float 2.000000e+00, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 // DEV-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 -// DEV-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4 +// DEV-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4, !noalias.addrspace [[META3]] // DEV-NEXT: store float [[TMP6]], ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // DEV-NEXT: [[TMP7:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // DEV-NEXT: [[TMP8:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // DEV-NEXT: store float 3.000000e+00, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 // DEV-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 -// DEV-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META3]] +// DEV-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !noalias.addrspace [[META3]], !amdgpu.no.remote.memory [[META4]] // DEV-NEXT: store float [[TMP10]], ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // DEV-NEXT: [[TMP11:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // DEV-NEXT: [[TMP12:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // DEV-NEXT: store float 4.000000e+00, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 // DEV-NEXT: [[TMP13:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 -// DEV-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !amdgpu.no.fine.grained.memory [[META3]] +// DEV-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]] // DEV-NEXT: store float [[TMP14]], ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // DEV-NEXT: [[TMP15:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // DEV-NEXT: ret void @@ -534,25 +534,25 @@ __device__ __host__ void test_multiple_attrs(float *a) { // OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4 // OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4 -// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META3]], !amdgpu.no.remote.memory [[META3]] +// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !noalias.addrspace [[META3]], !atomic.ignore.denormal.mode [[META4]], !amdgpu.no.remote.memory [[META4]] // OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4 // OPT-NEXT: [[TMP4:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // OPT-NEXT: store float 2.000000e+00, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 // OPT-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP1]], align 4 -// OPT-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4 +// OPT-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4, !noalias.addrspace [[META3]] // OPT-NEXT: store float [[TMP6]], ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // OPT-NEXT: [[TMP7:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP2]], align 4 // OPT-NEXT: [[TMP8:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // OPT-NEXT: store float 3.000000e+00, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 // OPT-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP3]], align 4 -// OPT-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META3]] +// OPT-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !noalias.addrspace [[META3]], !amdgpu.no.remote.memory [[META4]] // OPT-NEXT: store float [[TMP10]], ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // OPT-NEXT: [[TMP11:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP4]], align 4 // OPT-NEXT: [[TMP12:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8 // OPT-NEXT: store float 4.000000e+00, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 // OPT-NEXT: [[TMP13:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP5]], align 4 -// OPT-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !amdgpu.no.fine.grained.memory [[META3]] +// OPT-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !noalias.addrspace [[META3]], !amdgpu.no.fine.grained.memory [[META4]] // OPT-NEXT: store float [[TMP14]], ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // OPT-NEXT: [[TMP15:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP6]], align 4 // OPT-NEXT: ret void @@ -669,9 +669,11 @@ template<typename T> __device__ __host__ void test_template(T *a) { template __device__ __host__ void test_template<float>(float *a); //. -// DEV: [[META3]] = !{} +// DEV: [[META3]] = !{i32 5, i32 6} +// DEV: [[META4]] = !{} //. -// OPT: [[META3]] = !{} +// OPT: [[META3]] = !{i32 5, i32 6} +// OPT: [[META4]] = !{} //. // SPIRV-DEV: [[META4]] = !{} //. diff --git a/clang/test/CodeGenCUDA/private-atomics-undefined.hip b/clang/test/CodeGenCUDA/private-atomics-undefined.hip new file mode 100644 index 00000000000000..6a5c84d33a4af3 --- /dev/null +++ b/clang/test/CodeGenCUDA/private-atomics-undefined.hip @@ -0,0 +1,27 @@ +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm -o - -fcuda-is-device %s \ +// RUN: | FileCheck --check-prefix=DEFAULT %s +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm -o - -fcuda-is-device \ +// RUN: -fno-private-atomics-undefined %s \ +// RUN: | FileCheck --check-prefix=NO-PRIV-UNDEF %s + +#include "Inputs/cuda.h" + +// DEFAULT-LABEL: define {{.*}}@_Z18test_opencl_atomicPU7_Atomici +// DEFAULT: atomicrmw add +// DEFAULT-SAME: !noalias.addrspace +// +// NO-PRIV-UNDEF-LABEL: define {{.*}}@_Z18test_opencl_atomicPU7_Atomici +// NO-PRIV-UNDEF: !noalias.addrspace +__device__ void test_opencl_atomic(_Atomic int *a) { + __opencl_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __OPENCL_MEMORY_SCOPE_WORK_GROUP); +} + +// DEFAULT-LABEL: define {{.*}}@_Z18test_scoped_atomicPi +// DEFAULT: atomicrmw add +// DEFAULT-SAME: !noalias.addrspace +// +// NO-PRIV-UNDEF-NOT: !noalias.addrspace +// NO-PRIV-UNDEF: ret void +__device__ void test_scoped_atomic(int *a) { + __scoped_atomic_fetch_add(a, 1, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
