llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Christian Sigg (chsigg) <details> <summary>Changes</summary> The !amdgpu.ignore.denormal.mode metadata tells the backend that an atomicrmw fadd need not honor the function's denormal mode, so a native atomic instruction whose denormal behavior is fixed in hardware may be used instead of a CAS loop. Nothing about that is AMDGPU specific: NVPTX has exactly the same problem with atom.add, whose FTZ behavior depends on the address space and cannot be controlled. Promote it to a target independent fixed metadata kind, !atomic.ignore.denormal.mode, and switch the AMDGPU, SPIR-V and OpenMP producers and consumers over to it. Document it in LangRef, and point AMDGPUUsage at that description rather than duplicating it. Existing IR keeps working: AutoUpgrade renames the metadata on atomicrmw instructions when parsing textual IR and when materializing bitcode. The upgrade is deliberately scoped to atomicrmw rather than being applied to every attachment of that name, since that is the only place the metadata was ever meaningful. Because bitcode can be materialized one function at a time, the bitcode side hooks into BitcodeReader::materialize() rather than a module-wide pass, which is the path clang's bitcode linking takes. This is not intended to change behavior for any existing target. --- Patch is 446.71 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/217585.diff 73 Files Affected: - (modified) clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp (+4-2) - (modified) clang/lib/CodeGen/Targets/AMDGPU.cpp (+2-1) - (modified) clang/lib/CodeGen/Targets/SPIR.cpp (+2-1) - (modified) clang/test/CodeGen/AMDGPU/amdgpu-atomic-float.c (+4-4) - (modified) clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu (+1-1) - (modified) clang/test/CodeGenCUDA/atomic-options.hip (+12-12) - (modified) clang/test/CodeGenHIP/amdgpu-flat-atomic-fadd.hip (+8-8) - (modified) clang/test/CodeGenHIP/amdgpu-global-atomic-fadd.hip (+4-4) - (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl (+2-2) - (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl (+1-1) - (modified) clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl (+1-1) - (modified) clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx942.cl (+2-2) - (modified) clang/test/OpenMP/amdgpu-unsafe-fp-atomics.cpp (+14-9) - (modified) flang/test/Driver/atomic-control-options.f90 (+2-2) - (modified) llvm/docs/AMDGPUUsage.rst (+14-10) - (modified) llvm/docs/LangRef.md (+64) - (modified) llvm/docs/ReleaseNotes.md (+5) - (modified) llvm/include/llvm/IR/AutoUpgrade.h (+9) - (modified) llvm/include/llvm/IR/FixedMetadataKinds.def (+1) - (modified) llvm/lib/AsmParser/LLParser.cpp (+1) - (modified) llvm/lib/Bitcode/Reader/BitcodeReader.cpp (+5) - (modified) llvm/lib/CodeGen/AtomicExpandPass.cpp (+1-1) - (modified) llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp (+3-3) - (modified) llvm/lib/IR/AutoUpgrade.cpp (+63-2) - (modified) llvm/lib/Target/AMDGPU/SIISelLowering.cpp (+1-1) - (modified) llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp (+3-2) - (added) llvm/test/Assembler/atomic-metadata-upgrade.ll (+41) - (added) llvm/test/Bitcode/Inputs/atomic-metadata-upgrade-caller.ll (+10) - (modified) llvm/test/Bitcode/amdgcn-atomic.ll (+2-2) - (modified) llvm/test/Bitcode/amdgpu-unsafe-fp-atomics-upgrade.ll (+8-8) - (added) llvm/test/Bitcode/atomic-metadata-upgrade.ll (+49) - (added) llvm/test/Bitcode/atomic-metadata-upgrade.ll.bc () - (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/global-atomic-fadd.f32-no-rtn.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/global-atomic-fadd.f32-rtn.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/a-v-flat-atomicrmw.ll (+24-24) - (modified) llvm/test/CodeGen/AMDGPU/a-v-global-atomicrmw.ll (+24-24) - (modified) llvm/test/CodeGen/AMDGPU/atomicrmw-expand.ll (+3-3) - (modified) llvm/test/CodeGen/AMDGPU/buffer-fat-pointer-atomicrmw-fadd.ll (+4-4) - (modified) llvm/test/CodeGen/AMDGPU/cgp-addressing-modes-gfx908.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/flat-atomicrmw-fadd.ll (+16-16) - (modified) llvm/test/CodeGen/AMDGPU/global-atomic-fadd.f32-no-rtn.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/global-atomic-fadd.f32-rtn.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/global-atomicrmw-fadd-wrong-subtarget.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/global-atomicrmw-fadd.ll (+10-10) - (modified) llvm/test/CodeGen/AMDGPU/global-atomics-fp-wrong-subtarget.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/global-saddr-atomics.gfx908.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/global_atomics_scan_fadd.ll (+16-16) - (modified) llvm/test/CodeGen/AMDGPU/global_atomics_scan_fmax.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/global_atomics_scan_fmin.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/local-atomicrmw-fadd.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/local-atomicrmw-fmax.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/local-atomicrmw-fmin.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/local-atomicrmw-fsub.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/ptradd-sdag-optimizations.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/shl_add_ptr_global.ll (+1-1) - (modified) llvm/test/CodeGen/SPIRV/amdgcnspirv-atomic-metadata-decoration.ll (+2-2) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-f32-agent.ll (+84-84) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-f32-system.ll (+78-78) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-f64-agent.ll (+44-44) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-f64-system.ll (+41-41) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-mmra.ll (+4-4) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-rmw-fadd-flat-specialization-preserve-name.ll (+1-1) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-rmw-fadd-flat-specialization.ll (+22-22) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-rmw-fadd.ll (+43-43) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-v2bf16-agent.ll (+30-30) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-v2bf16-system.ll (+28-28) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-v2f16-agent.ll (+34-34) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomic-v2f16-system.ll (+32-32) - (modified) llvm/test/Transforms/AtomicExpand/AMDGPU/expand-atomicrmw-flat-noalias-addrspace.ll (+5-5) - (modified) llvm/test/Transforms/InferAddressSpaces/AMDGPU/global-atomicrmw-fadd.ll (+2-2) - (modified) mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp (+2-1) - (modified) mlir/test/Target/LLVMIR/omptarget-atomic-update-control-options.mlir (+1-1) - (modified) mlir/test/Target/LLVMIR/rocdl.mlir (+1-1) ``````````diff diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 667c6508040ae..a4eb7ec126583 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -21,6 +21,7 @@ #include "llvm/IR/IntrinsicsAMDGPU.h" #include "llvm/IR/IntrinsicsR600.h" #include "llvm/IR/IntrinsicsSPIRV.h" +#include "llvm/IR/LLVMContext.h" #include "llvm/IR/MemoryModelRelaxationAnnotations.h" #include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Support/AtomicOrdering.h" @@ -2050,10 +2051,11 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, llvm::MDTuple *EmptyMD = MDNode::get(getLLVMContext(), {}); RMW->setMetadata("amdgpu.no.fine.grained.memory", EmptyMD); - // Most targets require "amdgpu.ignore.denormal.mode" to emit the native + // Most targets require "atomic.ignore.denormal.mode" to emit the native // instruction, but this only matters for float fadd. if (BinOp == llvm::AtomicRMWInst::FAdd && Val->getType()->isFloatTy()) - RMW->setMetadata("amdgpu.ignore.denormal.mode", EmptyMD); + RMW->setMetadata(llvm::LLVMContext::MD_atomic_ignore_denormal_mode, + EmptyMD); } return Builder.CreateBitCast(RMW, OrigTy); diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 07e2eac39305d..0b5ed1898f138 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -10,6 +10,7 @@ #include "TargetInfo.h" #include "clang/AST/DeclCXX.h" #include "llvm/ADT/StringExtras.h" +#include "llvm/IR/LLVMContext.h" #include "llvm/IR/MemoryModelRelaxationAnnotations.h" #include "llvm/Support/AMDGPUAddrSpace.h" @@ -583,7 +584,7 @@ void AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata( if (AO.getOption(clang::AtomicOptionKind::IgnoreDenormalMode) && RMW->getOperation() == llvm::AtomicRMWInst::FAdd && RMW->getType()->isFloatTy()) - RMW->setMetadata("amdgpu.ignore.denormal.mode", Empty); + RMW->setMetadata(llvm::LLVMContext::MD_atomic_ignore_denormal_mode, Empty); } bool AMDGPUTargetCodeGenInfo::shouldEmitStaticExternCAliases() const { diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp index 7f7f1a2a1fe8c..0b3bd5a4aab7f 100644 --- a/clang/lib/CodeGen/Targets/SPIR.cpp +++ b/clang/lib/CodeGen/Targets/SPIR.cpp @@ -12,6 +12,7 @@ #include "clang/AST/DeclCXX.h" #include "clang/Basic/LangOptions.h" #include "llvm/IR/DerivedTypes.h" +#include "llvm/IR/LLVMContext.h" #include <stdint.h> #include <utility> @@ -589,7 +590,7 @@ void SPIRVTargetCodeGenInfo::setTargetAtomicMetadata( if (AO.getOption(clang::AtomicOptionKind::IgnoreDenormalMode) && RMW->getOperation() == llvm::AtomicRMWInst::FAdd && RMW->getType()->isFloatTy()) - RMW->setMetadata("amdgpu.ignore.denormal.mode", Empty); + RMW->setMetadata(llvm::LLVMContext::MD_atomic_ignore_denormal_mode, Empty); } /// Construct a SPIR-V target extension type for the given OpenCL image type. diff --git a/clang/test/CodeGen/AMDGPU/amdgpu-atomic-float.c b/clang/test/CodeGen/AMDGPU/amdgpu-atomic-float.c index 8537fcd8f6ea2..f688ababfb8ba 100644 --- a/clang/test/CodeGen/AMDGPU/amdgpu-atomic-float.c +++ b/clang/test/CodeGen/AMDGPU/amdgpu-atomic-float.c @@ -13,7 +13,7 @@ // UNSAFE-LABEL: define dso_local float @test_float_post_inc( // UNSAFE-SAME: ) #[[ATTR0:[0-9]+]] { // UNSAFE-NEXT: [[ENTRY:.*:]] -// UNSAFE-NEXT: [[TMP0:%.*]] = atomicrmw fadd ptr addrspacecast (ptr addrspace(1) @test_float_post_inc.n to ptr), float 1.000000e+00 seq_cst, align 4, !amdgpu.no.fine.grained.memory [[META2:![0-9]+]], !amdgpu.no.remote.memory [[META2]], !amdgpu.ignore.denormal.mode [[META2]] +// UNSAFE-NEXT: [[TMP0:%.*]] = atomicrmw fadd ptr addrspacecast (ptr addrspace(1) @test_float_post_inc.n to ptr), float 1.000000e+00 seq_cst, align 4, !atomic.ignore.denormal.mode [[META2:![0-9]+]], !amdgpu.no.fine.grained.memory [[META2]], !amdgpu.no.remote.memory [[META2]] // UNSAFE-NEXT: ret float [[TMP0]] // // SAFE-SPIRV-LABEL: define spir_func float @test_float_post_inc( @@ -25,7 +25,7 @@ // UNSAFE-SPIRV-LABEL: define spir_func float @test_float_post_inc( // UNSAFE-SPIRV-SAME: ) addrspace(4) #[[ATTR0:[0-9]+]] { // UNSAFE-SPIRV-NEXT: [[ENTRY:.*:]] -// UNSAFE-SPIRV-NEXT: [[TMP0:%.*]] = atomicrmw fadd ptr addrspace(4) addrspacecast (ptr addrspace(1) @test_float_post_inc.n to ptr addrspace(4)), float 1.000000e+00 seq_cst, align 4, !amdgpu.no.fine.grained.memory [[META2:![0-9]+]], !amdgpu.no.remote.memory [[META2]], !amdgpu.ignore.denormal.mode [[META2]] +// UNSAFE-SPIRV-NEXT: [[TMP0:%.*]] = atomicrmw fadd ptr addrspace(4) addrspacecast (ptr addrspace(1) @test_float_post_inc.n to ptr addrspace(4)), float 1.000000e+00 seq_cst, align 4, !atomic.ignore.denormal.mode [[META2:![0-9]+]], !amdgpu.no.fine.grained.memory [[META2]], !amdgpu.no.remote.memory [[META2]] // UNSAFE-SPIRV-NEXT: ret float [[TMP0]] // float test_float_post_inc() @@ -82,7 +82,7 @@ float test_float_pre_dc() // UNSAFE-LABEL: define dso_local float @test_float_pre_inc( // UNSAFE-SAME: ) #[[ATTR0]] { // UNSAFE-NEXT: [[ENTRY:.*:]] -// UNSAFE-NEXT: [[TMP0:%.*]] = atomicrmw fadd ptr addrspacecast (ptr addrspace(1) @test_float_pre_inc.n to ptr), float 1.000000e+00 seq_cst, align 4, !amdgpu.no.fine.grained.memory [[META2]], !amdgpu.no.remote.memory [[META2]], !amdgpu.ignore.denormal.mode [[META2]] +// UNSAFE-NEXT: [[TMP0:%.*]] = atomicrmw fadd ptr addrspacecast (ptr addrspace(1) @test_float_pre_inc.n to ptr), float 1.000000e+00 seq_cst, align 4, !atomic.ignore.denormal.mode [[META2]], !amdgpu.no.fine.grained.memory [[META2]], !amdgpu.no.remote.memory [[META2]] // UNSAFE-NEXT: [[TMP1:%.*]] = fadd float [[TMP0]], 1.000000e+00 // UNSAFE-NEXT: ret float [[TMP1]] // @@ -96,7 +96,7 @@ float test_float_pre_dc() // UNSAFE-SPIRV-LABEL: define spir_func float @test_float_pre_inc( // UNSAFE-SPIRV-SAME: ) addrspace(4) #[[ATTR0]] { // UNSAFE-SPIRV-NEXT: [[ENTRY:.*:]] -// UNSAFE-SPIRV-NEXT: [[TMP0:%.*]] = atomicrmw fadd ptr addrspace(4) addrspacecast (ptr addrspace(1) @test_float_pre_inc.n to ptr addrspace(4)), float 1.000000e+00 seq_cst, align 4, !amdgpu.no.fine.grained.memory [[META2]], !amdgpu.no.remote.memory [[META2]], !amdgpu.ignore.denormal.mode [[META2]] +// UNSAFE-SPIRV-NEXT: [[TMP0:%.*]] = atomicrmw fadd ptr addrspace(4) addrspacecast (ptr addrspace(1) @test_float_pre_inc.n to ptr addrspace(4)), float 1.000000e+00 seq_cst, align 4, !atomic.ignore.denormal.mode [[META2]], !amdgpu.no.fine.grained.memory [[META2]], !amdgpu.no.remote.memory [[META2]] // UNSAFE-SPIRV-NEXT: [[TMP1:%.*]] = fadd float [[TMP0]], 1.000000e+00 // UNSAFE-SPIRV-NEXT: ret float [[TMP1]] // diff --git a/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu b/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu index 4255b28cb23eb..da4b1c31dace0 100644 --- a/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu +++ b/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu @@ -31,7 +31,7 @@ __global__ void ffp1(float *p) { // SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]] // SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], [[DEFMD]] - // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4, [[FADDMD:!amdgpu.no.fine.grained.memory ![0-9]+, !amdgpu.no.remote.memory ![0-9]+, !amdgpu.ignore.denormal.mode ![0-9]+$]] + // UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4, [[FADDMD:!atomic.ignore.denormal.mode ![0-9]+, !amdgpu.no.fine.grained.memory ![0-9]+, !amdgpu.no.remote.memory ![0-9]+$]] // UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4, [[DEFMD:!amdgpu.no.fine.grained.memory ![0-9]+, !amdgpu.no.remote.memory ![0-9]+$]] // UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 4, [[DEFMD]] // UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 4, [[DEFMD]] diff --git a/clang/test/CodeGenCUDA/atomic-options.hip b/clang/test/CodeGenCUDA/atomic-options.hip index 76b92e7154eb8..2bbd14c692e7c 100644 --- a/clang/test/CodeGenCUDA/atomic-options.hip +++ b/clang/test/CodeGenCUDA/atomic-options.hip @@ -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, !amdgpu.no.remote.memory [[META3:![0-9]+]], !amdgpu.ignore.denormal.mode [[META3]] +// 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: 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 @@ -89,7 +89,7 @@ // SPIRV-OPT-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8 // SPIRV-OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4 // SPIRV-OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4 -// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4:![0-9]+]], !amdgpu.ignore.denormal.mode [[META4]] +// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META4:![0-9]+]], !amdgpu.no.remote.memory [[META4]] // SPIRV-OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4 // SPIRV-OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4 // SPIRV-OPT-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, !amdgpu.no.remote.memory [[META3]], !amdgpu.ignore.denormal.mode [[META3]] +// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META3]], !amdgpu.no.remote.memory [[META3]] // 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 @@ -172,7 +172,7 @@ __device__ __host__ void test_default(float *a) { // SPIRV-OPT-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8 // SPIRV-OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4 // SPIRV-OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4 -// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]] +// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META4]], !amdgpu.no.remote.memory [[META4]] // SPIRV-OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4 // SPIRV-OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4 // SPIRV-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, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.ignore.denormal.mode [[META3]] +// 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: 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, !amdgpu.ignore.denormal.mode [[META3]] +// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META3]] // 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 @@ -241,7 +241,7 @@ __device__ __host__ void test_one(float *a) { // SPIRV-DEV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8 // SPIRV-DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4 // SPIRV-DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4 -// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]] +// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META4]], !amdgpu.no.fine.grained.memory [[META4]] // SPIRV-DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4 // SPIRV-DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4 // SPIRV-DEV-NEXT: ret void @@ -257,7 +257,7 @@ __device__ __host__ void test_one(float *a) { // SPIRV-OPT-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8 // SPIRV-OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4 // SPIRV-OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4 -// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META4]] +// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META4]] // SPIRV-OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4 // SPIRV-OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4 // SPIRV-OPT-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, !amdgpu.ignore.denormal.mode [[META3]] +// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META3]] // 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 @@ -427,7 +427,7 @@ __device__ __host__ void test_three(float *a) { // SPIRV-OPT-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8 // SPIRV-OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4 // SPIRV-OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4 -// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META4]] +// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META4]] // SPIRV-OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4 // SPIRV-OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4 // SPIRV-OPT-NEXT: ret void @@ -534,7 +534,7 @@ __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, !amdgpu.no.remote.memory [[META3]], !amdgpu.ignore.denormal.mode [[META3]] +// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META3]], !amdgpu.no.remote.memory [[META3]] // 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 @@ -614,7 +614,7 @@ __device__ __host__ void test_multiple_attrs(float *a) { // SPIRV-OPT-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8 // SPIRV-OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4 // SPIRV-OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4 -// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]] +// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META4]], !amdgpu.no.remote.memory [[META4]] // SPIRV-OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4 // SPIRV-OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4 // SPIRV-OPT-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8 diff --git a/clang/test/CodeGenHIP/amdgpu-flat-atomic-fadd.hip b/clang/test/CodeGenHIP/amdgpu-flat-atomic-fadd.hip index 67574ae6427f9..a8866961a92c9 100644 --- a/clang/test/CodeGenHIP/amdgpu-flat-atomic-fadd.hip +++ b/clang/test/CodeGenHIP/amdgpu-flat-atomic-fadd.hip @@ -25,7 +25,7 @@ __device__ double global_double; // CHECK-NEXT: store float [[VAL]], ptr [[VAL_ADDR_ASCAST]], align 4 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR_ASCAST]], align 8 // CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[VAL_ADDR_ASCAST]], align 4 -// CHECK-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3:![0-9]+]], !amdgpu.ignore.denormal.mode [[META3]] +// CHECK-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] syncscope("agent") monotonic, align 4, !atomic.ignore.denormal.mode [[META3:![0-9]+]], !amdgpu.no.fine.grained.memory [[META3]] // CHECK-NEXT: store float [[TMP2]], ptr [[RESULT_ASCAST]], align 4 // CHECK-NEXT: ret void // @@ -46,7 +46,7 @@ __device__ double global_double; // SPIRV-NEXT: [[TMP1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[PTR_ADDR_ASCAST]], align 8 // SPIRV-NEXT: [[TMP2:%.*]] = addrspacecast ptr addrspace(4) [[TMP1]] to ptr // SPIRV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(4) [[VAL_ADDR_ASCAST]], align 4 -// SPIRV-NEXT: [[TMP4:%.*]] = atomicrmw fadd ptr [[TMP2]], float [[TMP3]] syncscope("device") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META5:![0-9]+]], !amdgpu.ignore.denormal.mode [[META5]] +// SPIRV-NEXT: [[TMP4:%.*]] = atomicrmw fadd ptr [[TMP2]], float [[TMP3]] syncscope("device") monotonic, align 4, !amdgpu.no.fine.grained.memory [[META5:![0-9]+]], !atomic.ignore.denormal.mode [[META5]] // SPIRV-NEXT: store float [[TMP4]], ptr addrspace(4) [[RESULT_ASCAST]], align 4 // SPIRV-NEXT: br label %[[IF... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/217585 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
