[llvm] [NVPTX] Cleanup support for atomic loads/stores (PR #201468)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Jun 3 15:48:36 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Akshay Deodhar (akshayrdeodhar)
<details>
<summary>Changes</summary>
- Adds exhaustive python-scripted testing for codegen for loads and stores in NVPTX, for all valid combinations of datatype x ordering x scope x addrspace x volatility, in the style of atomicrmw/cmpxchg/fence.
- Migrates the emulation of seq_cst loads to AtomicExpandPass using bracketInstructionWithFences. For this, updates AtomicExpandPass to set load/store ordering after a fence split. Cleans up associated helper functions.
- Emulates formerly-unsupported atomic loads and stores using fences in AtomicExpandPass. This leads to all possible memory orders being emulated even for SM60, using membar as a fence.
- Relaxes the lowering of seq_cst stores to fence.sc; st.relaxed instead of fence.sc; st.release.
- Updates NVPTXAtomicLower to downgrade atomic local loads/stores to nonatomic, in a fashion similar to atomicrmw.
- Note that volatile loads are currently being lowered incorrectly- will be fixed in a followup PR.
---
Patch is 4.65 MiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/201468.diff
17 Files Affected:
- (modified) llvm/lib/CodeGen/AtomicExpandPass.cpp (+2-2)
- (modified) llvm/lib/Target/NVPTX/NVPTXAtomicLower.cpp (+26-5)
- (modified) llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp (+12-46)
- (modified) llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h (+1-1)
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+100-24)
- (modified) llvm/test/CodeGen/NVPTX/lit.local.cfg (+1-1)
- (added) llvm/test/CodeGen/NVPTX/load-sm60.ll (+20340)
- (added) llvm/test/CodeGen/NVPTX/load-sm70.ll (+19956)
- (added) llvm/test/CodeGen/NVPTX/load-sm90.ll (+22986)
- (modified) llvm/test/CodeGen/NVPTX/load-store-sm-70.ll (+180-180)
- (modified) llvm/test/CodeGen/NVPTX/load-store-sm-90.ll (+60-60)
- (added) llvm/test/CodeGen/NVPTX/load.py (+194)
- (added) llvm/test/CodeGen/NVPTX/store-sm60.ll (+19360)
- (added) llvm/test/CodeGen/NVPTX/store-sm70.ll (+19168)
- (added) llvm/test/CodeGen/NVPTX/store-sm90.ll (+21950)
- (added) llvm/test/CodeGen/NVPTX/store.py (+206)
- (modified) llvm/test/lit.cfg.py (+5-1)
``````````diff
diff --git a/llvm/lib/CodeGen/AtomicExpandPass.cpp b/llvm/lib/CodeGen/AtomicExpandPass.cpp
index 960d2492c2856..13ec35ca3f1cb 100644
--- a/llvm/lib/CodeGen/AtomicExpandPass.cpp
+++ b/llvm/lib/CodeGen/AtomicExpandPass.cpp
@@ -350,7 +350,7 @@ bool AtomicExpandImpl::processAtomicInstr(Instruction *I) {
}
MadeChange |= tryInsertFencesForAtomic(
- LI, isAcquireOrStronger(LI->getOrdering()), AtomicOrdering::Monotonic);
+ LI, isAcquireOrStronger(LI->getOrdering()), TLI->atomicOperationOrderAfterFenceSplit(LI));
MadeChange |= tryExpandAtomicLoad(LI);
return MadeChange;
@@ -373,7 +373,7 @@ bool AtomicExpandImpl::processAtomicInstr(Instruction *I) {
}
MadeChange |= tryInsertFencesForAtomic(
- SI, isReleaseOrStronger(SI->getOrdering()), AtomicOrdering::Monotonic);
+ SI, isReleaseOrStronger(SI->getOrdering()), TLI->atomicOperationOrderAfterFenceSplit(SI));
MadeChange |= tryExpandAtomicStore(SI);
return MadeChange;
diff --git a/llvm/lib/Target/NVPTX/NVPTXAtomicLower.cpp b/llvm/lib/Target/NVPTX/NVPTXAtomicLower.cpp
index 918daf6c04ecd..badb4ebb9f9fc 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAtomicLower.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAtomicLower.cpp
@@ -42,15 +42,36 @@ class NVPTXAtomicLower : public FunctionPass {
} // namespace
bool NVPTXAtomicLower::runOnFunction(Function &F) {
- SmallVector<AtomicRMWInst *> LocalMemoryAtomics;
- for (Instruction &I : instructions(F))
- if (AtomicRMWInst *RMWI = dyn_cast<AtomicRMWInst>(&I))
+ SmallVector<AtomicRMWInst *> LocalRMWs;
+ SmallVector<LoadInst *> LocalAtomicLoads;
+ SmallVector<StoreInst *> LocalAtomicStores;
+ // TODO: Handle cmpxchg
+ for (Instruction &I : instructions(F)) {
+ if (auto *RMWI = dyn_cast<AtomicRMWInst>(&I)) {
if (RMWI->getPointerAddressSpace() == ADDRESS_SPACE_LOCAL)
- LocalMemoryAtomics.push_back(RMWI);
+ LocalRMWs.push_back(RMWI);
+ } else if (auto *LI = dyn_cast<LoadInst>(&I)) {
+ if (LI->isAtomic() &&
+ LI->getPointerAddressSpace() == ADDRESS_SPACE_LOCAL)
+ LocalAtomicLoads.push_back(LI);
+ } else if (auto *SI = dyn_cast<StoreInst>(&I)) {
+ if (SI->isAtomic() &&
+ SI->getPointerAddressSpace() == ADDRESS_SPACE_LOCAL)
+ LocalAtomicStores.push_back(SI);
+ }
+ }
bool Changed = false;
- for (AtomicRMWInst *RMWI : LocalMemoryAtomics)
+ for (AtomicRMWInst *RMWI : LocalRMWs)
Changed |= lowerAtomicRMWInst(RMWI);
+ for (LoadInst *LI : LocalAtomicLoads) {
+ LI->setAtomic(AtomicOrdering::NotAtomic);
+ Changed = true;
+ }
+ for (StoreInst *SI : LocalAtomicStores) {
+ SI->setAtomic(AtomicOrdering::NotAtomic);
+ Changed = true;
+ }
return Changed;
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index ede1deb5400b0..7b16de900d11f 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -551,15 +551,8 @@ NVPTX::Scope NVPTXDAGToDAGISel::getAtomicScope(const MemSDNode *N) const {
namespace {
-struct OperationOrderings {
- NVPTX::Ordering InstructionOrdering, FenceOrdering;
- OperationOrderings(NVPTX::Ordering IO = NVPTX::Ordering::NotAtomic,
- NVPTX::Ordering FO = NVPTX::Ordering::NotAtomic)
- : InstructionOrdering(IO), FenceOrdering(FO) {}
-};
-
-static OperationOrderings
-getOperationOrderings(MemSDNode *N, const NVPTXSubtarget *Subtarget) {
+static NVPTX::Ordering
+getInstructionOrdering(MemSDNode *N, const NVPTXSubtarget *Subtarget) {
AtomicOrdering Ordering = N->getSuccessOrdering();
auto CodeAddrSpace = NVPTXDAGToDAGISel::getAddrSpace(N);
@@ -730,18 +723,8 @@ getOperationOrderings(MemSDNode *N, const NVPTXSubtarget *Subtarget) {
//
// This sets the ordering of the fence to SequentiallyConsistent, and
// sets the corresponding ordering for the instruction.
- NVPTX::Ordering InstrOrder;
- if (N->readMem())
- InstrOrder = NVPTX::Ordering::Acquire;
- else if (N->writeMem())
- InstrOrder = NVPTX::Ordering::Release;
- else
- report_fatal_error(
- formatv("NVPTX does not support SequentiallyConsistent Ordering on "
- "read-modify-writes yet: {}",
- N->getOperationName()));
- return OperationOrderings(InstrOrder,
- NVPTX::Ordering::SequentiallyConsistent);
+ llvm_unreachable("seq_cst ordering should not reach this point, "
+ "is emulated in AtomicExpandPass");
}
}
report_fatal_error(
@@ -888,14 +871,11 @@ static unsigned int getFenceOp(NVPTX::Ordering O, NVPTX::Scope S,
llvm_unreachable("unhandled ordering");
}
-// Returns Memory Order and Scope of a memory instruction, and
-// inserts any fence before the instruction that's required to
-// implement its memory ordering.
+// Returns Memory Order and Scope of a memory instruction.
std::pair<NVPTX::Ordering, NVPTX::Scope>
-NVPTXDAGToDAGISel::insertMemoryInstructionFence(SDLoc DL, SDValue &Chain,
- MemSDNode *N) {
- auto [InstructionOrdering, FenceOrdering] =
- getOperationOrderings(N, Subtarget);
+NVPTXDAGToDAGISel::getMemInstOrderingScope(MemSDNode *N) {
+ auto InstructionOrdering =
+ getInstructionOrdering(N, Subtarget);
auto Scope = getOperationScope(N, InstructionOrdering);
// Singlethread scope has no inter-thread synchronization requirements, so
@@ -907,20 +887,6 @@ NVPTXDAGToDAGISel::insertMemoryInstructionFence(SDLoc DL, SDValue &Chain,
InstructionOrdering != NVPTX::Ordering::Volatile)
return {NVPTX::Ordering::NotAtomic, Scope};
- // If a fence is required before the operation, insert it:
- switch (NVPTX::Ordering(FenceOrdering)) {
- case NVPTX::Ordering::NotAtomic:
- break;
- case NVPTX::Ordering::SequentiallyConsistent: {
- auto Op = getFenceOp(FenceOrdering, Scope, Subtarget);
- Chain = SDValue(CurDAG->getMachineNode(Op, DL, MVT::Other, Chain), 0);
- break;
- }
- default:
- report_fatal_error(
- formatv("Unexpected fence ordering: \"{}\".",
- OrderingToString(NVPTX::Ordering(FenceOrdering))));
- }
return {InstructionOrdering, Scope};
}
@@ -1124,7 +1090,7 @@ bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) {
SDLoc DL(LD);
SDValue Chain = N->getOperand(0);
- const auto [Ordering, Scope] = insertMemoryInstructionFence(DL, Chain, LD);
+ const auto [Ordering, Scope] = getMemInstOrderingScope(LD);
const unsigned FromTypeWidth = LD->getMemoryVT().getSizeInBits();
@@ -1203,7 +1169,7 @@ bool NVPTXDAGToDAGISel::tryLoadVector(SDNode *N) {
const MVT EltVT = LD->getSimpleValueType(0);
SDLoc DL(LD);
SDValue Chain = LD->getChain();
- const auto [Ordering, Scope] = insertMemoryInstructionFence(DL, Chain, LD);
+ const auto [Ordering, Scope] = getMemInstOrderingScope(LD);
// Type Setting: fromType + fromTypeWidth
//
@@ -1396,7 +1362,7 @@ bool NVPTXDAGToDAGISel::tryStore(SDNode *N) {
SDLoc DL(ST);
SDValue Chain = ST->getChain();
- const auto [Ordering, Scope] = insertMemoryInstructionFence(DL, Chain, ST);
+ const auto [Ordering, Scope] = getMemInstOrderingScope(ST);
// Vector Setting
const unsigned ToTypeWidth = ST->getMemoryVT().getSizeInBits();
@@ -1447,7 +1413,7 @@ bool NVPTXDAGToDAGISel::tryStoreVector(SDNode *N) {
SDLoc DL(ST);
SDValue Chain = ST->getChain();
- const auto [Ordering, Scope] = insertMemoryInstructionFence(DL, Chain, ST);
+ const auto [Ordering, Scope] = getMemInstOrderingScope(ST);
const unsigned NumElts = getStoreVectorNumElts(ST);
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
index fcb5700dcb6d4..f2d3272d1111a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
@@ -110,7 +110,7 @@ class LLVM_LIBRARY_VISIBILITY NVPTXDAGToDAGISel : public SelectionDAGISel {
// instruction, if needed to implement the instructions memory order. Required
// fences after the instruction need to be handled elsewhere.
std::pair<NVPTX::Ordering, NVPTX::Scope>
- insertMemoryInstructionFence(SDLoc DL, SDValue &Chain, MemSDNode *N);
+ getMemInstOrderingScope(MemSDNode *N);
NVPTX::Scope getOperationScope(MemSDNode *N, NVPTX::Ordering O) const;
public:
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index e13911f87eed5..f5e39bda0d358 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -7600,11 +7600,35 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
return AtomicExpansionKind::CmpXChg;
}
+// Atomic load lowering grid (rows: native memory-ordering support; columns:
+// LLVM ordering):
+//
+// seq_cst acquire
+// supported fence.sc; ld.acquire ld.acquire
+// unsupported fence.sc; ld.relaxed; fence.acquire ld.relaxed; fence.acquire
+//
+// Atomic store lowering grid (same rows; columns are the store-side orderings):
+//
+// seq_cst release
+// supported fence.sc; st.relaxed st.release
+// unsupported fence.sc; st.relaxed fence.release; st.relaxed
+//
+// "supported" = STI.hasMemoryOrdering() (sm_70+ with PTX 6.0+). monotonic and
+// weaker drop out of these grids -- they need no fences and pass through.
+// Stores never get a trailing fence here; their release-side ordering is
+// provided entirely by the leading fence (or by typed `st.release` when
+// available).
+
bool NVPTXTargetLowering::shouldInsertFencesForAtomic(
const Instruction *I) const {
- // This function returns true iff the operation is emulated using a CAS-loop,
- // or if it has the memory order seq_cst (which is not natively supported in
- // the PTX `atom` instruction).
+ // Returns true iff this op needs a leading and/or trailing fence emitted
+ // around it (see emitLeadingFence / emitTrailingFence). Fires when:
+ // - the operation is emulated using a CAS-loop, OR
+ // - it has memory order seq_cst (PTX has no seq_cst atom/ld/st), OR
+ // - it is a load that has an unsupported memory order
+ // (acquire / seq_cst on !hasMemoryOrdering targets), OR
+ // - it is a store that has an unsupported memory order
+ // (release / seq_cst on !hasMemoryOrdering targets).
//
// atomicrmw and cmpxchg instructions not efficiently supported by PTX
// are lowered to CAS emulation loops that preserve their memory order,
@@ -7617,6 +7641,9 @@ bool NVPTXTargetLowering::shouldInsertFencesForAtomic(
// and scope. Since PTX does not support seq_cst, we emulate it by lowering to
// a fence.sc followed by an atom according to the PTX atomics ABI
// https://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/atomic-abi.html
+ //
+ // Atomic loads with a memory order not supported on the SM can be
+ // inefficiently emulated by appending a `fence.acquire` after a relaxed load.
if (auto *CI = dyn_cast<AtomicCmpXchgInst>(I))
return (cast<IntegerType>(CI->getCompareOperand()->getType())
->getBitWidth() < STI.getMinCmpXchgSizeInBits()) ||
@@ -7624,11 +7651,29 @@ bool NVPTXTargetLowering::shouldInsertFencesForAtomic(
if (auto *RI = dyn_cast<AtomicRMWInst>(I))
return shouldExpandAtomicRMWInIR(RI) == AtomicExpansionKind::CmpXChg ||
RI->getOrdering() == AtomicOrdering::SequentiallyConsistent;
+ if (auto *LI = dyn_cast<LoadInst>(I))
+ return LI->getOrdering() == AtomicOrdering::SequentiallyConsistent ||
+ (!STI.hasMemoryOrdering() && isAcquireOrStronger(LI->getOrdering()));
+ if (auto *SI = dyn_cast<StoreInst>(I))
+ return SI->getOrdering() == AtomicOrdering::SequentiallyConsistent ||
+ (!STI.hasMemoryOrdering() && isReleaseOrStronger(SI->getOrdering()));
return false;
}
AtomicOrdering NVPTXTargetLowering::atomicOperationOrderAfterFenceSplit(
const Instruction *I) const {
+ // Returns the new AtomicOrdering for the original op once AtomicExpandPass
+ // has split it into leading/trailing fences + the op itself. Per the grids
+ // above:
+ // - cmpxchg / atomicrmw at seq_cst (not CAS-emulated): Acquire
+ // ("fence.sc; atom.<op>.acquire" pattern).
+ // - load at any ordering when hasMemoryOrdering: Acquire
+ // ("fence.sc; ld.acquire" for seq_cst; the acquire load is already typed
+ // on supported targets).
+ // - everything else (CAS-emulated ops, seq_cst on !hasMemoryOrdering,
+ // release-or-stronger store on !hasMemoryOrdering, etc.): Monotonic
+ // (i.e., relaxed) -- the surrounding fences carry the ordering.
+ //
// If the operation is emulated by a CAS-loop, we lower the instruction to
// atom.<op>.relaxed, since AtomicExpandPass will insert fences for enforcing
// the correct memory ordering around the CAS loop.
@@ -7652,10 +7697,13 @@ AtomicOrdering NVPTXTargetLowering::atomicOperationOrderAfterFenceSplit(
cast<IntegerType>(CI->getCompareOperand()->getType())->getBitWidth() >=
STI.getMinCmpXchgSizeInBits())
return AtomicOrdering::Acquire;
- else if (auto *RI = dyn_cast<AtomicRMWInst>(I);
+ if (auto *RI = dyn_cast<AtomicRMWInst>(I);
RI && RI->getOrdering() == AtomicOrdering::SequentiallyConsistent &&
shouldExpandAtomicRMWInIR(RI) == AtomicExpansionKind::None)
return AtomicOrdering::Acquire;
+ if (auto *LI = dyn_cast<LoadInst>(I);
+ LI && STI.hasMemoryOrdering())
+ return AtomicOrdering::Acquire;
return AtomicOrdering::Monotonic;
}
@@ -7663,23 +7711,42 @@ AtomicOrdering NVPTXTargetLowering::atomicOperationOrderAfterFenceSplit(
Instruction *NVPTXTargetLowering::emitLeadingFence(IRBuilderBase &Builder,
Instruction *Inst,
AtomicOrdering Ord) const {
+ // Emits the leading fence per the grids above:
+ // - cmpxchg / atomicrmw: fence.sc iff Ord == seq_cst, else fence.release
+ // if Ord is release-or-stronger.
+ // - load: fence.sc iff Ord == seq_cst.
+ // - store: same shape as cmpxchg/rmw (fence.sc or fence.release based on
+ // Ord being seq_cst vs release-or-stronger).
+ // Returns nullptr if no leading fence is needed for this Inst/Ord.
+ //
// prerequisite: shouldInsertFencesForAtomic() should have returned `true` for
// `Inst` before its memory order was modified. We cannot enforce this with an
// assert, because AtomicExpandPass will have modified the memory order
// between the initial call to shouldInsertFencesForAtomic() and the call to
// this function.
- if (!isa<AtomicCmpXchgInst>(Inst) && !isa<AtomicRMWInst>(Inst))
- return TargetLoweringBase::emitLeadingFence(Builder, Inst, Ord);
-
- // Specialize for cmpxchg and atomicrmw
auto SSID = getAtomicSyncScopeID(Inst);
assert(SSID.has_value() && "Expected an atomic operation");
-
- if (isReleaseOrStronger(Ord))
- return Builder.CreateFence(Ord == AtomicOrdering::SequentiallyConsistent
- ? AtomicOrdering::SequentiallyConsistent
- : AtomicOrdering::Release,
- SSID.value());
+ if (isa<AtomicCmpXchgInst>(Inst) || isa<AtomicRMWInst>(Inst)) {
+ // Specialize for cmpxchg and atomicrmw
+ if (isReleaseOrStronger(Ord))
+ return Builder.CreateFence(Ord == AtomicOrdering::SequentiallyConsistent
+ ? AtomicOrdering::SequentiallyConsistent
+ : AtomicOrdering::Release,
+ SSID.value());
+ } else if (auto *LI = dyn_cast<LoadInst>(Inst)) {
+ // volatile operations have sys scope by default
+ SSID = LI->isVolatile() ? Builder.getContext().getOrInsertSyncScopeID("") : SSID;
+ if (Ord == AtomicOrdering::SequentiallyConsistent)
+ return Builder.CreateFence(AtomicOrdering::SequentiallyConsistent, SSID.value());
+ } else if (auto *SI = dyn_cast<StoreInst>(Inst)) {
+ // volatile operations have sys scope by default
+ SSID = SI->isVolatile() ? Builder.getContext().getOrInsertSyncScopeID("") : SSID;
+ if (isReleaseOrStronger(Ord))
+ return Builder.CreateFence(Ord == AtomicOrdering::SequentiallyConsistent
+ ? AtomicOrdering::SequentiallyConsistent
+ : AtomicOrdering::Release,
+ SSID.value());
+ }
return nullptr;
}
@@ -7687,25 +7754,34 @@ Instruction *NVPTXTargetLowering::emitLeadingFence(IRBuilderBase &Builder,
Instruction *NVPTXTargetLowering::emitTrailingFence(IRBuilderBase &Builder,
Instruction *Inst,
AtomicOrdering Ord) const {
+ // Emits the trailing fence per the grids above:
+ // - cmpxchg / atomicrmw: fence.acquire iff Ord is acquire-or-stronger AND
+ // the op is being CAS-emulated (i.e., the inner CAS uses relaxed).
+ // - load: fence.acquire iff !hasMemoryOrdering (covers both the acquire
+ // and seq_cst cells of the load grid).
+ // - store: never (release-side ordering for stores comes from the leading
+ // fence or the typed `st.release` instruction).
+ // Returns nullptr if no trailing fence is needed for this Inst/Ord.
+ //
// prerequisite: shouldInsertFencesForAtomic() should have returned `true` for
// `Inst` before its memory order was modified. See `emitLeadingFence` for why
// this cannot be enforced with an assert. Specialize for cmpxchg and
// atomicrmw
auto *CI = dyn_cast<AtomicCmpXchgInst>(Inst);
auto *RI = dyn_cast<AtomicRMWInst>(Inst);
- if (!CI && !RI)
- return TargetLoweringBase::emitTrailingFence(Builder, Inst, Ord);
-
auto SSID = getAtomicSyncScopeID(Inst);
assert(SSID.has_value() && "Expected an atomic operation");
+ if (CI || RI) {
+ bool IsEmulated =
+ CI ? cast<IntegerType>(CI->getCompareOperand()->getType())
+ ->getBitWidth() < STI.getMinCmpXchgSizeInBits()
+ : shouldExpandAtomicRMWInIR(RI) == AtomicExpansionKind::CmpXChg;
- bool IsEmulated =
- CI ? cast<IntegerType>(CI->getCompareOperand()->getType())
- ->getBitWidth() < STI.getMinCmpXchgSizeInBits()
- : shouldExpandAtomicRMWInIR(RI) == AtomicExpansionKind::CmpXChg;
-
- if (isAcquireOrStronger(Ord) && IsEmulated)
- return Builder.CreateFence(AtomicOrdering::Acquire, SSID.value());
+ if (isAcquireOrStronger(Ord) && IsEmulated)
+ return Builder.CreateFence(AtomicOrdering::Acquire, SSID.value());
+ } else if (isa<LoadInst>(Inst) && !STI.hasMemoryOrdering()) {
+ return Builder.CreateFence(AtomicOrdering::Acquire, SSID.value());
+ }
return nullptr;
}
diff --git a/llvm/test/CodeGen/NVPTX/lit.local.cfg b/llvm/test/CodeGen/NVPTX/lit.local.cfg
index dede5b354bb85..344d17adc5594 100644
--- a/llvm/test/CodeGen/NVPTX/lit.local.cfg
+++ b/llvm/test/CodeGen/NVPTX/lit.local.cfg
@@ -1,4 +1,4 @@
if not "NVPTX" in config.root.targets:
config.unsupported = True
config.suffixes.add(".py")
-config.excludes = ["fence.py", "cmpxchg.py", "atomicrmw.py"]
+config.excludes = ["fence.py", "cmpxchg.py", "atomicrmw.py", "load.py", "store.py"]
diff --git a/llvm/test/CodeGen/NVPTX/load-sm60.ll b/llvm/test/CodeGen/NVPTX/load-sm60.ll
new file mode 100644
index 0000000000000..e39d262da43c8
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/load-sm60.ll
@@ -0,0 +1,20340 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 | FileCheck %s
+; RUN: %if ptxas-sm_60 && ptxas-isa-5.0 %{ llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 | %ptxas-verify -arch=sm_60 %}
+
+define i8 @notatomic_i8_generic(ptr %addr) {
+; CHECK-LABEL: notatomic_i8_generic(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [notatomic_i8_generic_param_0];
+; CHECK-NEXT: ld.b8 %r1, [%rd1];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %retval = load i8, ptr %addr
+ ret i8 %retval
+}
+
+define i8 @notat...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/201468
More information about the llvm-commits
mailing list