gandhi21299 updated this revision to Diff 363324.
gandhi21299 added a comment.
Herald added a subscriber: pengfei.
- transferred code over to AtomicExpandPass to be able to call
`getAnalysis<OptimizationRemarkEmitterWrapperPass>().getORE()` and simplify
code as `SITargetLowering::shouldExpandAtomicRMW()` is called there.
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
clang/test/CodeGenOpenCL/fp-atomics-optremarks-gfx90a.cl
llvm/lib/CodeGen/AtomicExpandPass.cpp
llvm/lib/Target/AMDGPU/SIISelLowering.cpp
llvm/test/CodeGen/AMDGPU/llc-pipeline.ll
llvm/test/CodeGen/X86/O0-pipeline.ll
llvm/test/CodeGen/X86/opt-pipeline.ll
Index: llvm/test/CodeGen/X86/opt-pipeline.ll
===================================================================
--- llvm/test/CodeGen/X86/opt-pipeline.ll
+++ llvm/test/CodeGen/X86/opt-pipeline.ll
@@ -16,15 +16,20 @@
; CHECK-NEXT: Target Pass Configuration
; CHECK-NEXT: Machine Module Information
; CHECK-NEXT: Target Transform Information
+; CHECK-NEXT: Profile summary info
; CHECK-NEXT: Type-Based Alias Analysis
; CHECK-NEXT: Scoped NoAlias Alias Analysis
; CHECK-NEXT: Assumption Cache Tracker
-; CHECK-NEXT: Profile summary info
; CHECK-NEXT: Create Garbage Collector Module Metadata
; CHECK-NEXT: Machine Branch Probability Analysis
; CHECK-NEXT: ModulePass Manager
; CHECK-NEXT: Pre-ISel Intrinsic Lowering
; CHECK-NEXT: FunctionPass Manager
+; CHECK-NEXT: Dominator Tree Construction
+; CHECK-NEXT: Natural Loop Information
+; CHECK-NEXT: Lazy Branch Probability Analysis
+; CHECK-NEXT: Lazy Block Frequency Analysis
+; CHECK-NEXT: Optimization Remark Emitter
; CHECK-NEXT: Expand Atomic instructions
; CHECK-NEXT: Lower AMX intrinsics
; CHECK-NEXT: Lower AMX type for load/store
Index: llvm/test/CodeGen/X86/O0-pipeline.ll
===================================================================
--- llvm/test/CodeGen/X86/O0-pipeline.ll
+++ llvm/test/CodeGen/X86/O0-pipeline.ll
@@ -10,13 +10,18 @@
; CHECK-NEXT: Target Pass Configuration
; CHECK-NEXT: Machine Module Information
; CHECK-NEXT: Target Transform Information
+; CHECK-NEXT: Profile summary info
; CHECK-NEXT: Create Garbage Collector Module Metadata
; CHECK-NEXT: Assumption Cache Tracker
-; CHECK-NEXT: Profile summary info
; CHECK-NEXT: Machine Branch Probability Analysis
; CHECK-NEXT: ModulePass Manager
; CHECK-NEXT: Pre-ISel Intrinsic Lowering
; CHECK-NEXT: FunctionPass Manager
+; CHECK-NEXT: Dominator Tree Construction
+; CHECK-NEXT: Natural Loop Information
+; CHECK-NEXT: Lazy Branch Probability Analysis
+; CHECK-NEXT: Lazy Block Frequency Analysis
+; CHECK-NEXT: Optimization Remark Emitter
; CHECK-NEXT: Expand Atomic instructions
; CHECK-NEXT: Lower AMX intrinsics
; CHECK-NEXT: Lower AMX type for load/store
Index: llvm/test/CodeGen/AMDGPU/llc-pipeline.ll
===================================================================
--- llvm/test/CodeGen/AMDGPU/llc-pipeline.ll
+++ llvm/test/CodeGen/AMDGPU/llc-pipeline.ll
@@ -44,6 +44,11 @@
; GCN-O0-NEXT: Lower OpenCL enqueued blocks
; GCN-O0-NEXT: Lower uses of LDS variables from non-kernel functions
; GCN-O0-NEXT: FunctionPass Manager
+; GCN-O0-NEXT: Dominator Tree Construction
+; GCN-O0-NEXT: Natural Loop Information
+; GCN-O0-NEXT: Lazy Branch Probability Analysis
+; GCN-O0-NEXT: Lazy Block Frequency Analysis
+; GCN-O0-NEXT: Optimization Remark Emitter
; GCN-O0-NEXT: Expand Atomic instructions
; GCN-O0-NEXT: Lower constant intrinsics
; GCN-O0-NEXT: Remove unreachable blocks from the CFG
@@ -180,6 +185,11 @@
; GCN-O1-NEXT: Lower uses of LDS variables from non-kernel functions
; GCN-O1-NEXT: FunctionPass Manager
; GCN-O1-NEXT: Infer address spaces
+; GCN-O1-NEXT: Dominator Tree Construction
+; GCN-O1-NEXT: Natural Loop Information
+; GCN-O1-NEXT: Lazy Branch Probability Analysis
+; GCN-O1-NEXT: Lazy Block Frequency Analysis
+; GCN-O1-NEXT: Optimization Remark Emitter
; GCN-O1-NEXT: Expand Atomic instructions
; GCN-O1-NEXT: AMDGPU Promote Alloca
; GCN-O1-NEXT: Dominator Tree Construction
@@ -431,6 +441,11 @@
; GCN-O1-OPTS-NEXT: Lower uses of LDS variables from non-kernel functions
; GCN-O1-OPTS-NEXT: FunctionPass Manager
; GCN-O1-OPTS-NEXT: Infer address spaces
+; GCN-O1-OPTS-NEXT: Dominator Tree Construction
+; GCN-O1-OPTS-NEXT: Natural Loop Information
+; GCN-O1-OPTS-NEXT: Lazy Branch Probability Analysis
+; GCN-O1-OPTS-NEXT: Lazy Block Frequency Analysis
+; GCN-O1-OPTS-NEXT: Optimization Remark Emitter
; GCN-O1-OPTS-NEXT: Expand Atomic instructions
; GCN-O1-OPTS-NEXT: AMDGPU Promote Alloca
; GCN-O1-OPTS-NEXT: Dominator Tree Construction
@@ -715,6 +730,11 @@
; GCN-O2-NEXT: Lower uses of LDS variables from non-kernel functions
; GCN-O2-NEXT: FunctionPass Manager
; GCN-O2-NEXT: Infer address spaces
+; GCN-O2-NEXT: Dominator Tree Construction
+; GCN-O2-NEXT: Natural Loop Information
+; GCN-O2-NEXT: Lazy Branch Probability Analysis
+; GCN-O2-NEXT: Lazy Block Frequency Analysis
+; GCN-O2-NEXT: Optimization Remark Emitter
; GCN-O2-NEXT: Expand Atomic instructions
; GCN-O2-NEXT: AMDGPU Promote Alloca
; GCN-O2-NEXT: Dominator Tree Construction
@@ -1001,6 +1021,11 @@
; GCN-O3-NEXT: Lower uses of LDS variables from non-kernel functions
; GCN-O3-NEXT: FunctionPass Manager
; GCN-O3-NEXT: Infer address spaces
+; GCN-O3-NEXT: Dominator Tree Construction
+; GCN-O3-NEXT: Natural Loop Information
+; GCN-O3-NEXT: Lazy Branch Probability Analysis
+; GCN-O3-NEXT: Lazy Block Frequency Analysis
+; GCN-O3-NEXT: Optimization Remark Emitter
; GCN-O3-NEXT: Expand Atomic instructions
; GCN-O3-NEXT: AMDGPU Promote Alloca
; GCN-O3-NEXT: Dominator Tree Construction
Index: llvm/lib/Target/AMDGPU/SIISelLowering.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -12099,16 +12099,16 @@
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")
+ if (!UnsafeFPAtomicFlag)
return AtomicExpansionKind::CmpXChg;
if (Subtarget->hasGFX90AInsts()) {
@@ -12126,8 +12126,9 @@
if (AS == AMDGPUAS::FLAT_ADDRESS)
return AtomicExpansionKind::CmpXChg;
- return RMW->use_empty() ? AtomicExpansionKind::None
- : AtomicExpansionKind::CmpXChg;
+ auto Kind = RMW->use_empty() ? AtomicExpansionKind::None
+ : AtomicExpansionKind::CmpXChg;
+ return Kind;
}
// DS FP atomics do repect the denormal mode, but the rounding mode is fixed
@@ -12137,12 +12138,10 @@
if (!Ty->isDoubleTy())
return AtomicExpansionKind::None;
- 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 Kind;
}
return AtomicExpansionKind::CmpXChg;
Index: llvm/lib/CodeGen/AtomicExpandPass.cpp
===================================================================
--- llvm/lib/CodeGen/AtomicExpandPass.cpp
+++ llvm/lib/CodeGen/AtomicExpandPass.cpp
@@ -17,6 +17,7 @@
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallVector.h"
+#include "llvm/Analysis/OptimizationRemarkEmitter.h"
#include "llvm/CodeGen/AtomicExpandUtils.h"
#include "llvm/CodeGen/RuntimeLibcalls.h"
#include "llvm/CodeGen/TargetLowering.h"
@@ -58,6 +59,7 @@
class AtomicExpand: public FunctionPass {
const TargetLowering *TLI = nullptr;
+ OptimizationRemarkEmitter *ORE;
public:
static char ID; // Pass identification, replacement for typeid
@@ -69,6 +71,9 @@
bool runOnFunction(Function &F) override;
private:
+ void getAnalysisUsage(AnalysisUsage &AU) const override;
+ bool emitAtomicExpansionRemarks(
+ AtomicRMWInst *RMW, TargetLowering::AtomicExpansionKind Kind);
bool bracketInstWithFences(Instruction *I, AtomicOrdering Order);
IntegerType *getCorrespondingIntegerType(Type *T, const DataLayout &DL);
LoadInst *convertAtomicLoadToIntegerType(LoadInst *LI);
@@ -165,11 +170,16 @@
Size <= TLI->getMaxAtomicSizeInBitsSupported() / 8;
}
+void AtomicExpand::getAnalysisUsage(AnalysisUsage &AU) const {
+ AU.addRequired<OptimizationRemarkEmitterWrapperPass>();
+}
+
bool AtomicExpand::runOnFunction(Function &F) {
auto *TPC = getAnalysisIfAvailable<TargetPassConfig>();
if (!TPC)
return false;
+ ORE = &getAnalysis<OptimizationRemarkEmitterWrapperPass>().getORE();
auto &TM = TPC->getTM<TargetMachine>();
if (!TM.getSubtargetImpl(F)->enableAtomicExpand())
return false;
@@ -569,8 +579,31 @@
}
}
+bool AtomicExpand::emitAtomicExpansionRemarks(
+ AtomicRMWInst *RMW, TargetLowering::AtomicExpansionKind Kind) {
+ bool UnsafeFPAtomicFlag = RMW->getFunction()
+ ->getFnAttribute("amdgpu-unsafe-fp-atomics")
+ .getValueAsBool();
+ if (Kind == TargetLowering::AtomicExpansionKind::CmpXChg) {
+ ORE->emit([&]() {
+ OptimizationRemark Remark(DEBUG_TYPE, "Passed", RMW->getFunction());
+ Remark << "A floating point atomic instruction was expanded into a CAS loop.";
+ return Remark;
+ });
+ } else if (Kind == TargetLowering::AtomicExpansionKind::None && UnsafeFPAtomicFlag) {
+ ORE->emit([&]() {
+ OptimizationRemark Remark(DEBUG_TYPE, "Passed", RMW->getFunction());
+ Remark << "An unsafe hardware instruction was generated.";
+ return Remark;
+ });
+ }
+ return false;
+}
+
bool AtomicExpand::tryExpandAtomicRMW(AtomicRMWInst *AI) {
- switch (TLI->shouldExpandAtomicRMWInIR(AI)) {
+ TargetLowering::AtomicExpansionKind Kind = TLI->shouldExpandAtomicRMWInIR(AI);
+ emitAtomicExpansionRemarks(AI, Kind);
+ switch (Kind) {
case TargetLoweringBase::AtomicExpansionKind::None:
return false;
case TargetLoweringBase::AtomicExpansionKind::LLSC: {
Index: clang/test/CodeGenOpenCL/fp-atomics-optremarks-gfx90a.cl
===================================================================
--- /dev/null
+++ clang/test/CodeGenOpenCL/fp-atomics-optremarks-gfx90a.cl
@@ -0,0 +1,28 @@
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \
+// RUN: -Rpass=atomic-expand -munsafe-fp-atomics %s -S -o - 2>&1 \
+// RUN: | FileCheck %s -check-prefix=GFX90A-HW
+
+typedef enum memory_order {
+ memory_order_relaxed = __ATOMIC_RELAXED,
+ memory_order_acquire = __ATOMIC_ACQUIRE,
+ memory_order_release = __ATOMIC_RELEASE,
+ memory_order_acq_rel = __ATOMIC_ACQ_REL,
+ memory_order_seq_cst = __ATOMIC_SEQ_CST
+} memory_order;
+
+typedef enum memory_scope {
+ memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
+ memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
+ memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
+ memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
+#if defined(cl_intel_subgroups) || defined(cl_khr_subgroups)
+ memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
+#endif
+} memory_scope;
+
+// remark: An unsafe hardware instruction was generated.
+// GFX90A-HW-LABEL: test_atomic_add
+// GFX90A-HW: global_atomic_add_f64
+float test_atomic_add(global atomic_double *d, double a) {
+ return __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
+}
Index: clang/test/CodeGenCUDA/fp-atomics-optremarks.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/fp-atomics-optremarks.cu
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN: -target-cpu gfx90a -Rpass=atomic-expand -S -o - 2>&1 | \
+// RUN: FileCheck %s --check-prefix=GFX90A-CAS
+
+// REQUIRES: amdgpu-registered-target
+
+#include "Inputs/cuda.h"
+#include <stdatomic.h>
+
+// GFX90A-CAS: A floating point atomic instruction was expanded into a CAS loop.
+// GFX90A-CAS-LABEL: _Z14atomic_add_casPf
+// 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);
+}
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
