gandhi21299 updated this revision to Diff 362093.
gandhi21299 added a comment.
- replaced `check-prefixes` to `check-prefix` in the test
- Function ffp1 emits a `flat_atomic_cmpswap` instead of `global_atomic_add`, I
am not sure if that is considered an unsafe hardware instruction. @arsenm ?
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D106891/new/
https://reviews.llvm.org/D106891
Files:
clang/test/CodeGenCUDA/fp-atomics-optremarks.cu
llvm/lib/Target/AMDGPU/SIISelLowering.cpp
llvm/lib/Target/AMDGPU/SIISelLowering.h
Index: llvm/lib/Target/AMDGPU/SIISelLowering.h
===================================================================
--- llvm/lib/Target/AMDGPU/SIISelLowering.h
+++ llvm/lib/Target/AMDGPU/SIISelLowering.h
@@ -30,6 +30,7 @@
class SITargetLowering final : public AMDGPUTargetLowering {
private:
const GCNSubtarget *Subtarget;
+ OptimizationRemarkEmitter *ORE;
public:
MVT getRegisterTypeForCallingConv(LLVMContext &Context,
Index: llvm/lib/Target/AMDGPU/SIISelLowering.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -19,6 +19,7 @@
#include "SIRegisterInfo.h"
#include "llvm/ADT/Statistic.h"
#include "llvm/Analysis/LegacyDivergenceAnalysis.h"
+#include "llvm/Analysis/OptimizationRemarkEmitter.h"
#include "llvm/BinaryFormat/ELF.h"
#include "llvm/CodeGen/Analysis.h"
#include "llvm/CodeGen/FunctionLoweringInfo.h"
@@ -12117,6 +12118,27 @@
return DenormMode == DenormalMode::getIEEE();
}
+static TargetLowering::AtomicExpansionKind
+atomicExpandReturn(OptimizationRemarkEmitter *ORE, AtomicRMWInst *RMW,
+ TargetLowering::AtomicExpansionKind Kind, bool UnsafeFlag) {
+ ORE = new OptimizationRemarkEmitter(RMW->getFunction());
+ if (Kind == TargetLowering::AtomicExpansionKind::CmpXChg) {
+ ORE->emit([&]() {
+ OptimizationRemark Remark(DEBUG_TYPE, "Passed", RMW->getFunction());
+ Remark << "An FP atomic instruction was expanded into a CAS loop.";
+ return Remark;
+ });
+ } else if (Kind == TargetLowering::AtomicExpansionKind::None && UnsafeFlag) {
+ ORE->emit([&]() {
+ OptimizationRemark Remark(DEBUG_TYPE, "Passed", RMW->getFunction());
+ Remark << "An unsafe hardware instruction was generated.";
+ return Remark;
+ });
+ }
+ delete ORE;
+ return Kind;
+}
+
TargetLowering::AtomicExpansionKind
SITargetLowering::shouldExpandAtomicRMWInIR(AtomicRMWInst *RMW) const {
switch (RMW->getOperation()) {
@@ -12132,35 +12154,42 @@
return AtomicExpansionKind::CmpXChg;
unsigned AS = RMW->getPointerAddressSpace();
-
+ bool UnsafeFPAtomicFlag = RMW->getFunction()
+ ->getFnAttribute("amdgpu-unsafe-fp-atomics")
+ .getValueAsBool();
if ((AS == AMDGPUAS::GLOBAL_ADDRESS || AS == AMDGPUAS::FLAT_ADDRESS) &&
Subtarget->hasAtomicFaddInsts()) {
// The amdgpu-unsafe-fp-atomics attribute enables generation of unsafe
// floating point atomic instructions. May generate more efficient code,
// but may not respect rounding and denormal modes, and may give incorrect
// results for certain memory destinations.
- if (RMW->getFunction()
- ->getFnAttribute("amdgpu-unsafe-fp-atomics")
- .getValueAsString() != "true")
- return AtomicExpansionKind::CmpXChg;
+ if (!UnsafeFPAtomicFlag)
+ return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::CmpXChg,
+ false);
+ atomicExpandReturn(ORE, RMW, AtomicExpansionKind::None, true);
if (Subtarget->hasGFX90AInsts()) {
if (Ty->isFloatTy() && AS == AMDGPUAS::FLAT_ADDRESS)
- return AtomicExpansionKind::CmpXChg;
+ return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::CmpXChg,
+ UnsafeFPAtomicFlag);
auto SSID = RMW->getSyncScopeID();
if (SSID == SyncScope::System ||
SSID == RMW->getContext().getOrInsertSyncScopeID("one-as"))
- return AtomicExpansionKind::CmpXChg;
+ return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::CmpXChg,
+ UnsafeFPAtomicFlag);
- return AtomicExpansionKind::None;
+ return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::None,
+ UnsafeFPAtomicFlag);
}
if (AS == AMDGPUAS::FLAT_ADDRESS)
- return AtomicExpansionKind::CmpXChg;
+ return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::CmpXChg,
+ UnsafeFPAtomicFlag);
- return RMW->use_empty() ? AtomicExpansionKind::None
- : AtomicExpansionKind::CmpXChg;
+ auto Kind = RMW->use_empty() ? AtomicExpansionKind::None
+ : AtomicExpansionKind::CmpXChg;
+ return atomicExpandReturn(ORE, RMW, Kind, UnsafeFPAtomicFlag);
}
// DS FP atomics do repect the denormal mode, but the rounding mode is fixed
@@ -12168,17 +12197,17 @@
// The only exception is DS_ADD_F64 which never flushes regardless of mode.
if (AS == AMDGPUAS::LOCAL_ADDRESS && Subtarget->hasLDSFPAtomics()) {
if (!Ty->isDoubleTy())
- return AtomicExpansionKind::None;
+ return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::None,
+ UnsafeFPAtomicFlag);
- return (fpModeMatchesGlobalFPAtomicMode(RMW) ||
- RMW->getFunction()
- ->getFnAttribute("amdgpu-unsafe-fp-atomics")
- .getValueAsString() == "true")
- ? AtomicExpansionKind::None
- : AtomicExpansionKind::CmpXChg;
+ auto Kind = (fpModeMatchesGlobalFPAtomicMode(RMW) || UnsafeFPAtomicFlag)
+ ? AtomicExpansionKind::None
+ : AtomicExpansionKind::CmpXChg;
+ return atomicExpandReturn(ORE, RMW, Kind, UnsafeFPAtomicFlag);
}
- return AtomicExpansionKind::CmpXChg;
+ return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::CmpXChg,
+ UnsafeFPAtomicFlag);
}
default:
break;
Index: clang/test/CodeGenCUDA/fp-atomics-optremarks.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/fp-atomics-optremarks.cu
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN: -target-cpu gfx90a -Rpass=.* -S -o - 2>&1 | \
+// RUN: FileCheck %s --check-prefix=GFX90A-CAS
+
+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN: -target-cpu gfx90a -munsafe-fp-atomics -Rpass=.* \
+// RUN: -S -o - 2>&1 | FileCheck %s --check-prefix=GFX90A-HW
+
+// REQUIRES: amdgpu-registered-target
+
+#include "Inputs/cuda.h"
+#include <stdatomic.h>
+
+// GFX90A-CAS: An FP atomic instruction was expanded into a CAS loop.
+// GFX90A-CAS: flat_atomic_cmpswap v0, v[2:3], v[4:5] glc
+// GFX90A-CAS: s_cbranch_execnz
+__device__ float atomic_add_cas(float *p) {
+ return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
+}
+
+// GFX90A-HW: An unsafe hardware instruction was generated.
+// GFX90A-HW: flat_atomic_cmpswap
+__device__ float ffp1(float *p) {
+ return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
