AMDGPU: Add subtarget feature for global atomic fadd denormal support

[arsenm]

Not sure what the behavior for gfx90a is. The SPG says it always flushes.
The instruction documentation says it does not.

[arsenm]

• AMDGPU: Remove flat/global fmin/fmax intrinsics #105642
• AMDGPU: Remove global/flat atomic fadd intrinics #97051
• AMDGPU: Remove flat/global atomic fadd v2bf16 intrinsics #97050
• clang/AMDGPU: Emit atomicrmw for flat/global atomic min/max f64 builtins #96876
• clang/AMDGPU: Emit atomicrmw for global/flat fadd v2bf16 builtins #96875
• clang/AMDGPU: Emit atomicrmw from flat_atomic_{f32|f64} builtins #96874
• clang/AMDGPU: Emit atomicrmw from {global|flat}_atomic_fadd_v2f16 builtins #96873
• clang/AMDGPU: Emit atomicrmw for __builtin_amdgcn_global_atomic_fadd_{f32|f64} #96872
• AMDGPU: Handle new atomicrmw metadata for fadd case #96760
• AMDGPU: Handle remote/fine-grained memory in atomicrmw fmin/fmax lowering #96759
• AMDGPU: Add subtarget feature for memory atomic fadd f64 #96444
• AMDGPU: Add subtarget feature for global atomic fadd denormal support #96443 👈
• AMDGPU: Add a subtarget feature for fine-grained remote memory support #96442
• AMDGPU: Handle legal v2bf16 atomicrmw fadd for gfx12 #95930
• AMDGPU: Legalize v2f16 atomicrmw fadd for buffer fat pointers #95929
• main

This stack of pull requests is managed by Graphite. Learn more about stacking.

Join @arsenm and the rest of your teammates on Graphite

[llvmbot]

@llvm/pr-subscribers-backend-amdgpu

Author: Matt Arsenault (arsenm)

Changes

Not sure what the behavior for gfx90a is. The SPG says it always flushes.
The instruction documentation says it does not.

---

Full diff: https://github.com/llvm/llvm-project/pull/96443.diff

2 Files Affected:

• (modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+11-2)
• (modified) llvm/lib/Target/AMDGPU/GCNSubtarget.h (+7)

diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index 7ff861f5b144d..5f798b4391704 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -788,6 +788,13 @@ def FeatureFlatAtomicFaddF32Inst
   "Has flat_atomic_add_f32 instruction"
 >;
 
+def FeatureMemoryAtomicFaddF32DenormalSupport
+  : SubtargetFeature<"memory-atomic-fadd-f32-denormal-support",
+  "HasAtomicMemoryAtomicFaddF32DenormalSupport",
+  "true",
+  "global/flat/buffer atomic fadd for float supports denormal handling"
+>;
+
 def FeatureAgentScopeFineGrainedRemoteMemoryAtomics
   : SubtargetFeature<"agent-scope-fine-grained-remote-memory-atomics",
   "HasAgentScopeFineGrainedRemoteMemoryAtomics",
@@ -1425,7 +1432,8 @@ def FeatureISAVersion9_4_Common : FeatureSet<
    FeatureKernargPreload,
    FeatureAtomicFMinFMaxF64GlobalInsts,
    FeatureAtomicFMinFMaxF64FlatInsts,
-   FeatureAgentScopeFineGrainedRemoteMemoryAtomics
+   FeatureAgentScopeFineGrainedRemoteMemoryAtomics,
+   FeatureMemoryAtomicFaddF32DenormalSupport
    ]>;
 
 def FeatureISAVersion9_4_0 : FeatureSet<
@@ -1628,7 +1636,8 @@ def FeatureISAVersion12 : FeatureSet<
    FeatureVGPRSingleUseHintInsts,
    FeatureScalarDwordx3Loads,
    FeatureDPPSrc1SGPR,
-   FeatureMaxHardClauseLength32]>;
+   FeatureMaxHardClauseLength32,
+   FeatureMemoryAtomicFaddF32DenormalSupport]>;
 
 def FeatureISAVersion12_Generic: FeatureSet<
   !listconcat(FeatureISAVersion12.Features,
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.h b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
index c40efbdcf7f0b..674d84422538f 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.h
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
@@ -167,6 +167,7 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
   bool HasAtomicFlatPkAdd16Insts = false;
   bool HasAtomicFaddRtnInsts = false;
   bool HasAtomicFaddNoRtnInsts = false;
+  bool HasAtomicMemoryAtomicFaddF32DenormalSupport = false;
   bool HasAtomicBufferGlobalPkAddF16NoRtnInsts = false;
   bool HasAtomicBufferGlobalPkAddF16Insts = false;
   bool HasAtomicCSubNoRtnInsts = false;
@@ -872,6 +873,12 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
 
   bool hasFlatAtomicFaddF32Inst() const { return HasFlatAtomicFaddF32Inst; }
 
+  /// \return true if the target's flat, global, and buffer atomic fadd for
+  /// float supports denormal handling.
+  bool hasMemoryAtomicFaddF32DenormalSupport() const {
+    return HasAtomicMemoryAtomicFaddF32DenormalSupport;
+  }
+
   /// \return true if atomic operations targeting fine-grained memory work
   /// correctly at device scope, in allocations in host or peer PCIe device
   /// memory.

[rampitec]

It is worse than that. It behaves differently depending on where atomic is executed. There is no single answer if this instruction supports denorms or not.

[arsenm]

It is worse than that. It behaves differently depending on where atomic is executed. There is no single answer if this instruction supports denorms or not.

That doesn't matter. The flat case that sometimes flushes is just a no. Flushing is never a guarantee, we only need to know a flush may happen

[arsenm]

Merge activity

• Jul 10, 8:37 AM EDT: @arsenm started a stack merge that includes this pull request via Graphite.
• Jul 10, 8:45 AM EDT: Graphite rebased this pull request as part of a merge.
• Jul 10, 8:47 AM EDT: @arsenm merged this pull request with Graphite.