[Mlir-commits] [clang] [flang] [llvm] [mlir] [IR] Generalize !amdgpu.ignore.denormal.mode into !atomic.ignore.denormal.mode (PR #217585)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Thu Aug 20 09:59:14 PDT 2026
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
More information about the Mlir-commits
mailing list