AMDGPU: Add subtarget feature for memory atomic fadd f64

[arsenm]

No description provided.

[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

---

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

3 Files Affected:

• (modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+9-1)
• (modified) llvm/lib/Target/AMDGPU/GCNSubtarget.h (+7)
• (modified) llvm/lib/Target/AMDGPU/SIISelLowering.cpp (+1-1)

diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index 5f798b4391704..fe3cd75d81009 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 FeatureFlatBufferGlobalAtomicFaddF64Inst
+  : SubtargetFeature<"flat-buffer-global-fadd-f64-inst",
+  "HasFlatBufferGlobalAtomicFaddF64Inst",
+  "true",
+  "Has flat, buffer, and global instructions for f64 atomic fadd"
+>;
+
 def FeatureMemoryAtomicFaddF32DenormalSupport
   : SubtargetFeature<"memory-atomic-fadd-f32-denormal-support",
   "HasAtomicMemoryAtomicFaddF32DenormalSupport",
@@ -1388,7 +1395,8 @@ def FeatureISAVersion9_0_A : FeatureSet<
      FeatureBackOffBarrier,
      FeatureKernargPreload,
      FeatureAtomicFMinFMaxF64GlobalInsts,
-     FeatureAtomicFMinFMaxF64FlatInsts
+     FeatureAtomicFMinFMaxF64FlatInsts,
+     FeatureFlatBufferGlobalAtomicFaddF64Inst
      ])>;
 
 def FeatureISAVersion9_0_C : FeatureSet<
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.h b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
index 674d84422538f..922435c5efaa6 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.h
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
@@ -174,6 +174,7 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
   bool HasAtomicGlobalPkAddBF16Inst = false;
   bool HasAtomicBufferPkAddBF16Inst = false;
   bool HasFlatAtomicFaddF32Inst = false;
+  bool HasFlatBufferGlobalAtomicFaddF64Inst = false;
   bool HasDefaultComponentZero = false;
   bool HasAgentScopeFineGrainedRemoteMemoryAtomics = false;
   bool HasDefaultComponentBroadcast = false;
@@ -873,6 +874,12 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
 
   bool hasFlatAtomicFaddF32Inst() const { return HasFlatAtomicFaddF32Inst; }
 
+  /// \return true if the target has flat, global, and buffer atomic fadd for
+  /// double.
+  bool hasFlatBufferGlobalAtomicFaddF64Inst() const {
+    return HasFlatBufferGlobalAtomicFaddF64Inst;
+  }
+
   /// \return true if the target's flat, global, and buffer atomic fadd for
   /// float supports denormal handling.
   bool hasMemoryAtomicFaddF32DenormalSupport() const {
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index eec750e5b8251..6b5ba160d6402 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -16028,7 +16028,7 @@ SITargetLowering::shouldExpandAtomicRMWInIR(AtomicRMWInst *RMW) const {
       return AtomicExpansionKind::CmpXChg;
 
     // global and flat atomic fadd f64: gfx90a, gfx940.
-    if (Subtarget->hasGFX90AInsts() && Ty->isDoubleTy())
+    if (Subtarget->hasFlatBufferGlobalAtomicFaddF64Inst() && Ty->isDoubleTy())
       return ReportUnsafeHWInst(AtomicExpansionKind::None);
 
     if (AS != AMDGPUAS::FLAT_ADDRESS && Ty->isFloatTy()) {

[rampitec]

Use it in a predicate when defining pseudos?

[arsenm]

Merge activity

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