[llvm] [SDAG][NVPTX] Support cache hints on atomic operations (PR #214918)
Yonah Goldberg via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 25 16:49:17 PDT 2026
https://github.com/YonahGoldberg updated https://github.com/llvm/llvm-project/pull/214918
>From 2bcb75aacaa01565c941be5c783666e43003faad Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Sat, 8 Aug 2026 02:52:44 +0000
Subject: [PATCH 1/6] atomic cache hints
---
.../SelectionDAG/SelectionDAGBuilder.cpp | 17 +-
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 116 +++++----
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 111 ++++++---
llvm/test/CodeGen/NVPTX/cache-hint-atomics.ll | 58 +++--
.../CodeGen/NVPTX/cache-hint-sm-version.ll | 232 ++++++++++++++++--
5 files changed, 422 insertions(+), 112 deletions(-)
diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
index a2b9a14ed9543..df0c442781a7d 100644
--- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
@@ -5254,9 +5254,11 @@ void SelectionDAGBuilder::visitAtomicCmpXchg(const AtomicCmpXchgInst &I) {
auto Flags = TLI.getAtomicMemOperandFlags(I, DAG.getDataLayout());
MachineFunction &MF = DAG.getMachineFunction();
+ const MDNode *MemCacheHint = getMemCacheHintMetadata(I);
MachineMemOperand *MMO = MF.getMachineMemOperand(
MachinePointerInfo(I.getPointerOperand()), Flags, MemVT.getStoreSize(),
- I.getAlign(), MMOMetadata(), SSID, SuccessOrdering, FailureOrdering);
+ I.getAlign(), MMOMetadata(AAMDNodes(), /*Ranges=*/nullptr, MemCacheHint),
+ SSID, SuccessOrdering, FailureOrdering);
SDValue L = DAG.getAtomicCmpSwap(ISD::ATOMIC_CMP_SWAP_WITH_SUCCESS,
dl, MemVT, VTs, InChain,
@@ -5325,9 +5327,11 @@ void SelectionDAGBuilder::visitAtomicRMW(const AtomicRMWInst &I) {
auto Flags = TLI.getAtomicMemOperandFlags(I, DAG.getDataLayout());
MachineFunction &MF = DAG.getMachineFunction();
+ const MDNode *MemCacheHint = getMemCacheHintMetadata(I);
MachineMemOperand *MMO = MF.getMachineMemOperand(
MachinePointerInfo(I.getPointerOperand()), Flags, MemVT.getStoreSize(),
- I.getAlign(), MMOMetadata(), SSID, Ordering);
+ I.getAlign(), MMOMetadata(AAMDNodes(), /*Ranges=*/nullptr, MemCacheHint),
+ SSID, Ordering);
SDValue L =
DAG.getAtomic(NT, dl, MemVT, InChain,
@@ -5372,9 +5376,11 @@ void SelectionDAGBuilder::visitAtomicLoad(const LoadInst &I) {
auto Flags = TLI.getLoadMemOperandFlags(I, DAG.getDataLayout(), AC, LibInfo);
const MDNode *Ranges = getRangeMetadata(I);
+ const MDNode *MemCacheHint = getMemCacheHintMetadata(I);
MachineMemOperand *MMO = DAG.getMachineFunction().getMachineMemOperand(
MachinePointerInfo(I.getPointerOperand()), Flags, MemVT.getStoreSize(),
- I.getAlign(), MMOMetadata(AAMDNodes(), Ranges), SSID, Order);
+ I.getAlign(), MMOMetadata(AAMDNodes(), Ranges, MemCacheHint), SSID,
+ Order);
InChain = TLI.prepareVolatileOrAtomicLoad(InChain, dl, DAG);
@@ -5409,9 +5415,12 @@ void SelectionDAGBuilder::visitAtomicStore(const StoreInst &I) {
auto Flags = TLI.getStoreMemOperandFlags(I, DAG.getDataLayout());
MachineFunction &MF = DAG.getMachineFunction();
+ const MDNode *MemCacheHint =
+ getMemCacheHintMetadata(I, I.getPointerOperandIndex());
MachineMemOperand *MMO = MF.getMachineMemOperand(
MachinePointerInfo(I.getPointerOperand()), Flags, MemVT.getStoreSize(),
- I.getAlign(), MMOMetadata(), SSID, Ordering);
+ I.getAlign(), MMOMetadata(AAMDNodes(), /*Ranges=*/nullptr, MemCacheHint),
+ SSID, Ordering);
SDValue Val = getValue(I.getValueOperand());
if (Val.getValueType() != MemVT)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index b3bb96c202c55..c26c943913ccb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -72,14 +72,21 @@ struct NVPTXScopes {
LLVMContext *Context = nullptr;
};
+enum class NVPTXMemCacheHintInstruction { Ld, St, Atom };
+
struct NVPTXMemCacheHintAccess {
+ NVPTXMemCacheHintInstruction Instruction;
NVPTX::AddressSpace AddrSpace;
- bool IsLoad;
unsigned NumElts;
unsigned EltWidth;
bool IsVolatile;
};
+struct NVPTXMemCacheHintOperands {
+ SDValue EvictionAndPrefetchHint;
+ SDValue CachePolicyReg;
+};
+
class NVPTXDAGToDAGISel : public SelectionDAGISel {
const NVPTXTargetMachine &TM;
@@ -149,7 +156,7 @@ class NVPTXDAGToDAGISel : public SelectionDAGISel {
// dropped. If L2::cache_hint is active, returns the hint with
// L2CacheHintBit set and a register containing the 64-bit cache policy
// value. Otherwise returns NOREG for the policy operand.
- std::pair<unsigned, SDValue>
+ NVPTXMemCacheHintOperands
getMemCacheHintOperands(const MemSDNode *N, NVPTXMemCacheHintAccess Access,
const SDLoc &DL);
@@ -1196,6 +1203,21 @@ static bool isGlobalOrGeneric(NVPTX::AddressSpace AddrSpace) {
AddrSpace == NVPTX::AddressSpace::Generic;
}
+static bool isLdOrSt(NVPTXMemCacheHintAccess Access) {
+ return Access.Instruction == NVPTXMemCacheHintInstruction::Ld ||
+ Access.Instruction == NVPTXMemCacheHintInstruction::St;
+}
+
+static bool isL1EvictionSupported(const NVPTXSubtarget &Subtarget,
+ NVPTX::L1Eviction Eviction,
+ NVPTXMemCacheHintAccess Access) {
+ if (Eviction == NVPTX::L1Eviction::Normal)
+ return true;
+
+ return isLdOrSt(Access) && !Access.IsVolatile &&
+ Subtarget.hasL1EvictionHint();
+}
+
static bool isL2PrefetchSupported(const NVPTXSubtarget &Subtarget,
NVPTX::L2Prefetch Prefetch,
NVPTXMemCacheHintAccess Access) {
@@ -1203,14 +1225,14 @@ static bool isL2PrefetchSupported(const NVPTXSubtarget &Subtarget,
case NVPTX::L2Prefetch::None:
return true;
case NVPTX::L2Prefetch::Bytes64:
- return Access.IsLoad && isGlobalOrGeneric(Access.AddrSpace) &&
- Subtarget.hasL2Prefetch64B();
+ return Access.Instruction == NVPTXMemCacheHintInstruction::Ld &&
+ isGlobalOrGeneric(Access.AddrSpace) && Subtarget.hasL2Prefetch64B();
case NVPTX::L2Prefetch::Bytes128:
- return Access.IsLoad && isGlobalOrGeneric(Access.AddrSpace) &&
- Subtarget.hasL2Prefetch128B();
+ return Access.Instruction == NVPTXMemCacheHintInstruction::Ld &&
+ isGlobalOrGeneric(Access.AddrSpace) && Subtarget.hasL2Prefetch128B();
case NVPTX::L2Prefetch::Bytes256:
- return Access.IsLoad && isGlobalOrGeneric(Access.AddrSpace) &&
- Subtarget.hasL2Prefetch256B();
+ return Access.Instruction == NVPTXMemCacheHintInstruction::Ld &&
+ isGlobalOrGeneric(Access.AddrSpace) && Subtarget.hasL2Prefetch256B();
}
llvm_unreachable("Unexpected L2 prefetch hint");
}
@@ -1221,22 +1243,28 @@ static bool isL2EvictionSupported(const NVPTXSubtarget &Subtarget,
if (Eviction == NVPTX::L2Eviction::Normal)
return true;
- return Subtarget.hasL2EvictionHint() && isGlobalOrGeneric(Access.AddrSpace) &&
- !Access.IsVolatile &&
+ return isLdOrSt(Access) && !Access.IsVolatile &&
+ Subtarget.hasL2EvictionHint() && isGlobalOrGeneric(Access.AddrSpace) &&
((Access.NumElts == 8 && Access.EltWidth == 32) ||
(Access.NumElts == 4 && Access.EltWidth == 64));
}
-std::pair<unsigned, SDValue> NVPTXDAGToDAGISel::getMemCacheHintOperands(
+static bool isCachePolicySupported(const NVPTXSubtarget &Subtarget,
+ NVPTXMemCacheHintAccess Access) {
+ return !Access.IsVolatile && isGlobalOrGeneric(Access.AddrSpace) &&
+ Subtarget.hasL2CacheHint();
+}
+
+NVPTXMemCacheHintOperands NVPTXDAGToDAGISel::getMemCacheHintOperands(
const MemSDNode *N, NVPTXMemCacheHintAccess Access, const SDLoc &DL) {
LLVMContext &Ctx = *CurDAG->getContext();
const MDNode *Node = N->getMemCacheHint();
SDValue PolicyReg = CurDAG->getRegister(NVPTX::NoRegister, MVT::i64);
if (!Node)
- return {0, PolicyReg};
+ return {getI32Imm(0, DL), PolicyReg};
if (Node->getNumOperands() == 0) {
emitInvalidMemCacheHint(Ctx, "empty hint node");
- return {0, PolicyReg};
+ return {getI32Imm(0, DL), PolicyReg};
}
NVPTX::L1Eviction L1 = NVPTX::L1Eviction::Normal;
@@ -1252,7 +1280,7 @@ std::pair<unsigned, SDValue> NVPTXDAGToDAGISel::getMemCacheHintOperands(
if (KeyStr == "nvvm.l1_eviction") {
auto ParsedL1 =
parseMemCacheHintStringValue(Ctx, KeyStr, Value, parseL1Eviction);
- if (ParsedL1 && !Access.IsVolatile && Subtarget->hasL1EvictionHint())
+ if (ParsedL1 && isL1EvictionSupported(*Subtarget, *ParsedL1, Access))
L1 = *ParsedL1;
continue;
}
@@ -1279,8 +1307,7 @@ std::pair<unsigned, SDValue> NVPTXDAGToDAGISel::getMemCacheHintOperands(
if (!ValCI)
emitInvalidMemCacheHint(
Ctx, "'nvvm.l2_cache_hint' expects an integer value");
- else if (isGlobalOrGeneric(Access.AddrSpace) && !Access.IsVolatile &&
- Subtarget->hasL2CacheHint())
+ else if (isCachePolicySupported(*Subtarget, Access))
CachePolicy = ValCI->getZExtValue();
continue;
}
@@ -1297,7 +1324,7 @@ std::pair<unsigned, SDValue> NVPTXDAGToDAGISel::getMemCacheHintOperands(
Bitfield::set<NVPTX::L2CacheHintBit>(EvictionAndPrefetchHint, true);
}
- return {EvictionAndPrefetchHint, PolicyReg};
+ return {getI32Imm(EvictionAndPrefetchHint, DL), PolicyReg};
}
bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) {
@@ -1345,7 +1372,7 @@ bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) {
const auto [Base, Offset] = selectADDR(N->getOperand(1), CurDAG);
const auto [EvictionAndPrefetchHint, PolicyReg] = getMemCacheHintOperands(
LD,
- {CodeAddrSpace, /*IsLoad=*/true,
+ {NVPTXMemCacheHintInstruction::Ld, CodeAddrSpace,
/*NumElts=*/1, /*EltWidth=*/FromTypeWidth, LD->isVolatile()},
DL);
@@ -1358,7 +1385,7 @@ bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) {
getI32Imm(UsedBytesMask, DL),
Base,
Offset,
- getI32Imm(EvictionAndPrefetchHint, DL),
+ EvictionAndPrefetchHint,
PolicyReg,
Chain};
@@ -1425,7 +1452,7 @@ bool NVPTXDAGToDAGISel::tryLoadVector(SDNode *N) {
const auto [EvictionAndPrefetchHint, PolicyReg] =
getMemCacheHintOperands(LD,
- {CodeAddrSpace, /*IsLoad=*/true,
+ {NVPTXMemCacheHintInstruction::Ld, CodeAddrSpace,
/*NumElts=*/LD->getNumValues() - 1,
/*EltWidth=*/FromTypeWidth, LD->isVolatile()},
DL);
@@ -1438,7 +1465,7 @@ bool NVPTXDAGToDAGISel::tryLoadVector(SDNode *N) {
getI32Imm(UsedBytesMask, DL),
Base,
Offset,
- getI32Imm(EvictionAndPrefetchHint, DL),
+ EvictionAndPrefetchHint,
PolicyReg,
Chain};
@@ -1493,18 +1520,17 @@ bool NVPTXDAGToDAGISel::tryLDG(MemSDNode *LD) {
ExtensionType != ISD::NON_EXTLOAD));
const auto [Base, Offset] = selectADDR(LD->getOperand(1), CurDAG);
- const auto [EvictionAndPrefetchHint, PolicyReg] =
- getMemCacheHintOperands(LD,
- {NVPTX::AddressSpace::Global,
- /*IsLoad=*/true, LD->getNumValues() - 1,
- FromTypeWidth, LD->isVolatile()},
- DL);
+ const auto [EvictionAndPrefetchHint, PolicyReg] = getMemCacheHintOperands(
+ LD,
+ {NVPTXMemCacheHintInstruction::Ld, NVPTX::AddressSpace::Global,
+ LD->getNumValues() - 1, FromTypeWidth, LD->isVolatile()},
+ DL);
SDValue Ops[] = {getI32Imm(FromType, DL),
getI32Imm(FromTypeWidth, DL),
getI32Imm(UsedBytesMask, DL),
Base,
Offset,
- getI32Imm(EvictionAndPrefetchHint, DL),
+ EvictionAndPrefetchHint,
PolicyReg,
LD->getChain()};
@@ -1618,7 +1644,7 @@ bool NVPTXDAGToDAGISel::tryStore(SDNode *N) {
// Extract eviction/prefetch hint and cache policy register.
const auto [EvictionAndPrefetchHint, PolicyReg] = getMemCacheHintOperands(
ST,
- {CodeAddrSpace, /*IsLoad=*/false,
+ {NVPTXMemCacheHintInstruction::St, CodeAddrSpace,
/*NumElts=*/1, /*EltWidth=*/ToTypeWidth, ST->isVolatile()},
DL);
@@ -1629,7 +1655,7 @@ bool NVPTXDAGToDAGISel::tryStore(SDNode *N) {
getI32Imm(ToTypeWidth, DL),
Base,
Offset,
- getI32Imm(EvictionAndPrefetchHint, DL),
+ EvictionAndPrefetchHint,
PolicyReg,
Chain};
@@ -1679,15 +1705,14 @@ bool NVPTXDAGToDAGISel::tryStoreVector(SDNode *N) {
// Extract eviction/prefetch hint and cache policy register.
const auto [EvictionAndPrefetchHint, PolicyReg] = getMemCacheHintOperands(
ST,
- {CodeAddrSpace, /*IsLoad=*/false, /*NumElts=*/NumElts,
- /*EltWidth=*/ToTypeWidth, ST->isVolatile()},
+ {NVPTXMemCacheHintInstruction::St, CodeAddrSpace,
+ /*NumElts=*/NumElts, /*EltWidth=*/ToTypeWidth, ST->isVolatile()},
DL);
const auto [Base, Offset] = selectADDR(Addr, CurDAG);
Ops.append({getI32Imm(Ordering, DL), getI32Imm(Scope, DL),
getI32Imm(CodeAddrSpace, DL), getI32Imm(ToTypeWidth, DL), Base,
- Offset, getI32Imm(EvictionAndPrefetchHint, DL), PolicyReg,
- Chain});
+ Offset, EvictionAndPrefetchHint, PolicyReg, Chain});
const MVT::SimpleValueType EltVT =
ST->getOperand(1).getSimpleValueType().SimpleTy;
@@ -2308,14 +2333,23 @@ void NVPTXDAGToDAGISel::selectAtomicSwap128(SDNode *N) {
const SDValue Chain = N->getOperand(0);
const auto [Base, Offset] = selectADDR(N->getOperand(1), CurDAG);
- SmallVector<SDValue, 5> Ops{Base, Offset};
+ SmallVector<SDValue, 10> Ops{Base, Offset};
Ops.append(N->op_begin() + 2, N->op_end());
- Ops.append({
- getI32Imm(getMemOrder(AN), dl),
- getI32Imm(getAtomicScope(AN), dl),
- getI32Imm(getAddrSpace(AN), dl),
- Chain,
- });
+ Ops.append({getI32Imm(getMemOrder(AN), dl), getI32Imm(getAtomicScope(AN), dl),
+ getI32Imm(getAddrSpace(AN), dl)});
+
+ if (N->getOpcode() == NVPTXISD::ATOMIC_SWAP_B128) {
+ unsigned EltWidth = AN->getMemoryVT().getFixedSizeInBits();
+ NVPTXMemCacheHintAccess Access{NVPTXMemCacheHintInstruction::Atom,
+ getAddrSpace(AN),
+ /*NumElts=*/1, EltWidth, AN->isVolatile()};
+ const auto [EvictionAndPrefetchHint, CachePolicyReg] =
+ getMemCacheHintOperands(AN, Access, dl);
+ Ops.push_back(EvictionAndPrefetchHint);
+ Ops.push_back(CachePolicyReg);
+ }
+
+ Ops.push_back(Chain);
assert(N->getOpcode() == NVPTXISD::ATOMIC_CMP_SWAP_B128 ||
N->getOpcode() == NVPTXISD::ATOMIC_SWAP_B128);
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 83caf03942b26..2e11871cefa1d 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2527,32 +2527,65 @@ class AtomicSDNodeXForm<SDNode atomic> {
SDNodeXForm GetAddSp = SDNodeXForm<atomic, [{
return getI32Imm(getAddrSpace(cast<MemSDNode>(N)), SDLoc(N));
}]>;
+
+ SDNodeXForm GetEvictionAndPrefetchHint = SDNodeXForm<atomic, [{
+ auto *MN = cast<MemSDNode>(N);
+ unsigned EltWidth = MN->getMemoryVT().getFixedSizeInBits();
+ NVPTXMemCacheHintAccess Access{
+ NVPTXMemCacheHintInstruction::Atom, getAddrSpace(MN),
+ /*NumElts=*/1, EltWidth, MN->isVolatile()};
+ return getMemCacheHintOperands(MN, Access, SDLoc(N))
+ .EvictionAndPrefetchHint;
+ }]>;
+
+ SDNodeXForm GetCachePolicy = SDNodeXForm<atomic, [{
+ auto *MN = cast<MemSDNode>(N);
+ unsigned EltWidth = MN->getMemoryVT().getFixedSizeInBits();
+ NVPTXMemCacheHintAccess Access{
+ NVPTXMemCacheHintInstruction::Atom, getAddrSpace(MN),
+ /*NumElts=*/1, EltWidth, MN->isVolatile()};
+ return getMemCacheHintOperands(MN, Access, SDLoc(N)).CachePolicyReg;
+ }]>;
}
-multiclass F_ATOMIC_2<RegTyInfo t, SDPatternOperator op, string op_str, SDNode atomic, list<Predicate> preds = []> {
- defvar asm_str = "atom${sem:sem}${scope:scope}${addsp:addsp}" # "." # op_str;
+multiclass F_ATOMIC_2<RegTyInfo t, SDPatternOperator op, string op_str,
+ string type_str, SDNode atomic,
+ list<Predicate> preds = []> {
+ defvar asm_str = "atom${sem:sem}${scope:scope}${addsp:addsp}." # op_str #
+ "${evictionAndPrefetchHint:l2}." # type_str #
+ " \t$dst, [$addr], $b${policy};";
let mayLoad = 1, mayStore = 1, hasSideEffects = 1 in {
- def _r : BasicFlagsNVPTXInst<(outs t.RC:$dst),
- (ins ADDR:$addr, t.RC:$b),
- (ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp),
- asm_str>,
+ def _r : NVPTXInst<
+ (outs t.RC:$dst),
+ (ins ADDR:$addr, t.RC:$b, AtomicCode:$sem, AtomicCode:$scope,
+ AtomicCode:$addsp,
+ EvictionAndPrefetchHint:$evictionAndPrefetchHint,
+ CachePolicy:$policy), asm_str>,
Requires<preds>;
if t.SupportsImm then
- def _i : BasicFlagsNVPTXInst<(outs t.RC:$dst),
- (ins ADDR:$addr, t.Imm:$b),
- (ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp),
- asm_str>,
+ def _i : NVPTXInst<
+ (outs t.RC:$dst),
+ (ins ADDR:$addr, t.Imm:$b, AtomicCode:$sem, AtomicCode:$scope,
+ AtomicCode:$addsp,
+ EvictionAndPrefetchHint:$evictionAndPrefetchHint,
+ CachePolicy:$policy), asm_str>,
Requires<preds>;
}
defvar XForm = AtomicSDNodeXForm<atomic>;
def : Pat<(op:$this addr:$addr, t.Ty:$b),
- (!cast<Instruction>(NAME # _r) ADDR:$addr, t.Ty:$b, (XForm.GetSem $this), (XForm.GetScope $this), (XForm.GetAddSp $this))>;
+ (!cast<Instruction>(NAME # _r) ADDR:$addr, t.Ty:$b, (XForm.GetSem $this),
+ (XForm.GetScope $this), (XForm.GetAddSp $this),
+ (XForm.GetEvictionAndPrefetchHint $this),
+ (XForm.GetCachePolicy $this))>;
if t.SupportsImm then
def : Pat<(op:$this addr:$addr, (t.Ty t.ImmNode:$b)),
- (!cast<Instruction>(NAME # _i) ADDR:$addr, (t.Ty t.ImmNode:$b), (XForm.GetSem $this), (XForm.GetScope $this), (XForm.GetAddSp $this))>;
+ (!cast<Instruction>(NAME # _i) ADDR:$addr, (t.Ty t.ImmNode:$b),
+ (XForm.GetSem $this), (XForm.GetScope $this), (XForm.GetAddSp $this),
+ (XForm.GetEvictionAndPrefetchHint $this),
+ (XForm.GetCachePolicy $this))>;
}
multiclass F_ATOMIC_3<RegTyInfo t, string op_str, SDPatternOperator op, SDNode atomic> {
@@ -2598,29 +2631,29 @@ multiclass F_ATOMIC_3<RegTyInfo t, string op_str, SDPatternOperator op, SDNode a
defm atomic_load_fadd : binary_atomic_op_fp<atomic_load_fadd>;
// atom_add
-defm INT_PTX_ATOM_ADD_32 : F_ATOMIC_2<I32RT, atomic_load_add_i32, "add.u32", atomic_load_add>;
-defm INT_PTX_ATOM_ADD_64 : F_ATOMIC_2<I64RT, atomic_load_add_i64, "add.u64", atomic_load_add>;
+defm INT_PTX_ATOM_ADD_32 : F_ATOMIC_2<I32RT, atomic_load_add_i32, "add", "u32", atomic_load_add>;
+defm INT_PTX_ATOM_ADD_64 : F_ATOMIC_2<I64RT, atomic_load_add_i64, "add", "u64", atomic_load_add>;
-defm INT_PTX_ATOM_ADD_F16 : F_ATOMIC_2<F16RT, atomic_load_fadd, "add.noftz.f16", atomic_load_fadd, [SM70, hasPTX<63>]>;
-defm INT_PTX_ATOM_ADD_BF16 : F_ATOMIC_2<BF16RT, atomic_load_fadd, "add.noftz.bf16", atomic_load_fadd, [SM90, hasPTX<78>]>;
-defm INT_PTX_ATOM_ADD_F32 : F_ATOMIC_2<F32RT, atomic_load_fadd, "add.f32", atomic_load_fadd>;
-defm INT_PTX_ATOM_ADD_F64 : F_ATOMIC_2<F64RT, atomic_load_fadd, "add.f64", atomic_load_fadd, [hasAtomAddF64]>;
+defm INT_PTX_ATOM_ADD_F16 : F_ATOMIC_2<F16RT, atomic_load_fadd, "add.noftz", "f16", atomic_load_fadd, [SM70, hasPTX<63>]>;
+defm INT_PTX_ATOM_ADD_BF16 : F_ATOMIC_2<BF16RT, atomic_load_fadd, "add.noftz", "bf16", atomic_load_fadd, [SM90, hasPTX<78>]>;
+defm INT_PTX_ATOM_ADD_F32 : F_ATOMIC_2<F32RT, atomic_load_fadd, "add", "f32", atomic_load_fadd>;
+defm INT_PTX_ATOM_ADD_F64 : F_ATOMIC_2<F64RT, atomic_load_fadd, "add", "f64", atomic_load_fadd, [hasAtomAddF64]>;
// atom_swap
-defm INT_PTX_ATOM_SWAP_32 : F_ATOMIC_2<I32RT, atomic_swap_i32, "exch.b32", atomic_swap>;
-defm INT_PTX_ATOM_SWAP_64 : F_ATOMIC_2<I64RT, atomic_swap_i64, "exch.b64", atomic_swap>;
+defm INT_PTX_ATOM_SWAP_32 : F_ATOMIC_2<I32RT, atomic_swap_i32, "exch", "b32", atomic_swap>;
+defm INT_PTX_ATOM_SWAP_64 : F_ATOMIC_2<I64RT, atomic_swap_i64, "exch", "b64", atomic_swap>;
// atom_max
-defm INT_PTX_ATOMIC_MAX_32 : F_ATOMIC_2<I32RT, atomic_load_max_i32, "max.s32", atomic_load_max>;
-defm INT_PTX_ATOMIC_MAX_64 : F_ATOMIC_2<I64RT, atomic_load_max_i64, "max.s64", atomic_load_max, [SM32]>;
-defm INT_PTX_ATOMIC_UMAX_32 : F_ATOMIC_2<I32RT, atomic_load_umax_i32, "max.u32", atomic_load_umax>;
-defm INT_PTX_ATOMIC_UMAX_64 : F_ATOMIC_2<I64RT, atomic_load_umax_i64, "max.u64", atomic_load_umax, [SM32]>;
+defm INT_PTX_ATOMIC_MAX_32 : F_ATOMIC_2<I32RT, atomic_load_max_i32, "max", "s32", atomic_load_max>;
+defm INT_PTX_ATOMIC_MAX_64 : F_ATOMIC_2<I64RT, atomic_load_max_i64, "max", "s64", atomic_load_max, [SM32]>;
+defm INT_PTX_ATOMIC_UMAX_32 : F_ATOMIC_2<I32RT, atomic_load_umax_i32, "max", "u32", atomic_load_umax>;
+defm INT_PTX_ATOMIC_UMAX_64 : F_ATOMIC_2<I64RT, atomic_load_umax_i64, "max", "u64", atomic_load_umax, [SM32]>;
// atom_min
-defm INT_PTX_ATOMIC_MIN_32 : F_ATOMIC_2<I32RT, atomic_load_min_i32, "min.s32", atomic_load_min>;
-defm INT_PTX_ATOMIC_MIN_64 : F_ATOMIC_2<I64RT, atomic_load_min_i64, "min.s64", atomic_load_min, [SM32]>;
-defm INT_PTX_ATOMIC_UMIN_32 : F_ATOMIC_2<I32RT, atomic_load_umin_i32, "min.u32", atomic_load_umin>;
-defm INT_PTX_ATOMIC_UMIN_64 : F_ATOMIC_2<I64RT, atomic_load_umin_i64, "min.u64", atomic_load_umin, [SM32]>;
+defm INT_PTX_ATOMIC_MIN_32 : F_ATOMIC_2<I32RT, atomic_load_min_i32, "min", "s32", atomic_load_min>;
+defm INT_PTX_ATOMIC_MIN_64 : F_ATOMIC_2<I64RT, atomic_load_min_i64, "min", "s64", atomic_load_min, [SM32]>;
+defm INT_PTX_ATOMIC_UMIN_32 : F_ATOMIC_2<I32RT, atomic_load_umin_i32, "min", "u32", atomic_load_umin>;
+defm INT_PTX_ATOMIC_UMIN_64 : F_ATOMIC_2<I64RT, atomic_load_umin_i64, "min", "u64", atomic_load_umin, [SM32]>;
// NOTE: The semantics for atomicrmw fmin (and fmax) upholds LangRef
// requirements. The LangRef requires the semantics of fmin/fmax to follow
@@ -2633,20 +2666,20 @@ defm INT_PTX_ATOMIC_UMIN_64 : F_ATOMIC_2<I64RT, atomic_load_umin_i64, "min.u64",
// [3] https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions-min
// atom_inc atom_dec
-defm INT_PTX_ATOM_INC_32 : F_ATOMIC_2<I32RT, atomic_load_uinc_wrap_i32, "inc.u32", atomic_load_uinc_wrap>;
-defm INT_PTX_ATOM_DEC_32 : F_ATOMIC_2<I32RT, atomic_load_udec_wrap_i32, "dec.u32", atomic_load_udec_wrap>;
+defm INT_PTX_ATOM_INC_32 : F_ATOMIC_2<I32RT, atomic_load_uinc_wrap_i32, "inc", "u32", atomic_load_uinc_wrap>;
+defm INT_PTX_ATOM_DEC_32 : F_ATOMIC_2<I32RT, atomic_load_udec_wrap_i32, "dec", "u32", atomic_load_udec_wrap>;
// atom_and
-defm INT_PTX_ATOM_AND_32 : F_ATOMIC_2<I32RT, atomic_load_and_i32, "and.b32", atomic_load_and>;
-defm INT_PTX_ATOM_AND_64 : F_ATOMIC_2<I64RT, atomic_load_and_i64, "and.b64", atomic_load_and, [SM32]>;
+defm INT_PTX_ATOM_AND_32 : F_ATOMIC_2<I32RT, atomic_load_and_i32, "and", "b32", atomic_load_and>;
+defm INT_PTX_ATOM_AND_64 : F_ATOMIC_2<I64RT, atomic_load_and_i64, "and", "b64", atomic_load_and, [SM32]>;
// atom_or
-defm INT_PTX_ATOM_OR_32 : F_ATOMIC_2<I32RT, atomic_load_or_i32, "or.b32", atomic_load_or>;
-defm INT_PTX_ATOM_OR_64 : F_ATOMIC_2<I64RT, atomic_load_or_i64, "or.b64", atomic_load_or, [SM32]>;
+defm INT_PTX_ATOM_OR_32 : F_ATOMIC_2<I32RT, atomic_load_or_i32, "or", "b32", atomic_load_or>;
+defm INT_PTX_ATOM_OR_64 : F_ATOMIC_2<I64RT, atomic_load_or_i64, "or", "b64", atomic_load_or, [SM32]>;
// atom_xor
-defm INT_PTX_ATOM_XOR_32 : F_ATOMIC_2<I32RT, atomic_load_xor_i32, "xor.b32", atomic_load_xor>;
-defm INT_PTX_ATOM_XOR_64 : F_ATOMIC_2<I64RT, atomic_load_xor_i64, "xor.b64", atomic_load_xor, [SM32]>;
+defm INT_PTX_ATOM_XOR_32 : F_ATOMIC_2<I32RT, atomic_load_xor_i32, "xor", "b32", atomic_load_xor>;
+defm INT_PTX_ATOM_XOR_64 : F_ATOMIC_2<I64RT, atomic_load_xor_i64, "xor", "b64", atomic_load_xor, [SM32]>;
// Define atom.cas for all combinations of size x addrspace x memory order
// supported in PTX *and* on the hardware.
@@ -2677,11 +2710,13 @@ let mayLoad = true, mayStore = true, hasSideEffects = true,
NVPTXInst<
(outs B64:$dst0, B64:$dst1),
(ins ADDR:$addr, B64:$amt0, B64:$amt1,
- AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp),
+ AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp,
+ EvictionAndPrefetchHint:$evictionAndPrefetchHint,
+ CachePolicy:$policy),
"{{\n\t"
".reg .b128 amt, dst;\n\t"
"mov.b128 amt, {$amt0, $amt1};\n\t"
- "atom${sem:sem}${scope:scope}${addsp:addsp}.exch.b128 dst, [$addr], amt;\n\t"
+ "atom${sem:sem}${scope:scope}${addsp:addsp}.exch${evictionAndPrefetchHint:l2}.b128 dst, [$addr], amt${policy};\n\t"
"mov.b128 {$dst0, $dst1}, dst;\n\t"
"}}">;
}
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-atomics.ll b/llvm/test/CodeGen/NVPTX/cache-hint-atomics.ll
index dab37a6b7510b..e768374ed6033 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-atomics.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-atomics.ll
@@ -1,36 +1,48 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --filter "^\s*(?:mov\.b64|ld(?:\.[A-Za-z0-9_:]+)*\.global|st(?:\.[A-Za-z0-9_:]+)*\.global|atom(?:\.[A-Za-z0-9_:]+)*\.global)" --version 6
-; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 | FileCheck %s
-; RUN: %if ptxas-sm_80 && ptxas-isa-7.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 | %ptxas-verify -arch=sm_80 %}
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --filter "^\s*(?:mov\.b64|ld(?:\.[A-Za-z0-9_:]+)*\.global|st(?:\.[A-Za-z0-9_:]+)*\.global|atom(?:\.[A-Za-z0-9_:]+)*(?:\.global|\.shared)|\.reg \.b128)" --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | FileCheck %s
+; RUN: %if ptxas-sm_100 && ptxas-isa-8.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | %ptxas-verify -arch=sm_100 %}
-; !mem.cache_hint is legal on atomic IR memory instructions, but
-; SelectionDAGBuilder currently drops it for atomic loads, stores, RMWs, and
-; cmpxchg.
-
-; TODO: This should eventually lower to ld.acquire.sys.global.L1::evict_first.b32.
define i32 @atomic_load_l1_hint(ptr addrspace(1) %p) {
; CHECK-LABEL: atomic_load_l1_hint(
-; CHECK: ld.acquire.sys.global.b32 %r1, [%rd1];
+; CHECK: ld.acquire.sys.global.L1::evict_first.b32 %r1, [%rd1];
%v = load atomic i32, ptr addrspace(1) %p acquire, align 4, !mem.cache_hint !0
ret i32 %v
}
-; TODO: This should eventually lower to st.release.sys.global.L2::cache_hint.b32.
define void @atomic_store_l2_cache_policy(ptr addrspace(1) %p, i32 %v) {
; CHECK-LABEL: atomic_store_l2_cache_policy(
-; CHECK: st.release.sys.global.b32 [%rd1], %r1;
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: st.release.sys.global.L2::cache_hint.b32 [%rd1], %r1, %rd2;
store atomic i32 %v, ptr addrspace(1) %p release, align 4, !mem.cache_hint !2
ret void
}
-; TODO: This should eventually lower to atom.relaxed.sys.global.add.L2::cache_hint.u32.
define i32 @atomicrmw_add_l2_cache_policy(ptr addrspace(1) %p, i32 %v) {
; CHECK-LABEL: atomicrmw_add_l2_cache_policy(
-; CHECK: atom.relaxed.sys.global.add.u32 %r2, [%rd1], %r1;
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: atom.relaxed.sys.global.add.L2::cache_hint.u32 %r2, [%rd1], %r1, %rd2;
%old = atomicrmw add ptr addrspace(1) %p, i32 %v monotonic, align 4, !mem.cache_hint !1
ret i32 %old
}
+; PTX atom only supports L2::cache_hint. Other cache hints must be dropped.
+define i32 @atomicrmw_add_drops_unsupported_hints(ptr addrspace(1) %p, i32 %v) {
+; CHECK-LABEL: atomicrmw_add_drops_unsupported_hints(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: atom.relaxed.sys.global.add.L2::cache_hint.u32 %r2, [%rd1], %r1, %rd2;
+ %old = atomicrmw add ptr addrspace(1) %p, i32 %v monotonic, align 4, !mem.cache_hint !3
+ ret i32 %old
+}
+
; PTX atom.cas does not have a .level::cache_hint operand.
+define i32 @atomicrmw_xchg_l2_cache_policy(ptr addrspace(1) %p) {
+; CHECK-LABEL: atomicrmw_xchg_l2_cache_policy(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: atom.relaxed.sys.global.exch.L2::cache_hint.b32 %r1, [%rd1], 42, %rd2;
+ %old = atomicrmw xchg ptr addrspace(1) %p, i32 42 monotonic, align 4, !mem.cache_hint !1
+ ret i32 %old
+}
+
define i32 @cmpxchg_l2_cache_policy(ptr addrspace(1) %p, i32 %cmp, i32 %new) {
; CHECK-LABEL: cmpxchg_l2_cache_policy(
; CHECK: atom.relaxed.sys.global.cas.b32 %r3, [%rd1], %r1, %r2;
@@ -39,9 +51,27 @@ define i32 @cmpxchg_l2_cache_policy(ptr addrspace(1) %p, i32 %cmp, i32 %new) {
ret i32 %old
}
+; L2::cache_hint is not supported on shared memory.
+define i32 @atomicrmw_shared_l2_cache_policy(ptr addrspace(3) %p, i32 %v) {
+; CHECK-LABEL: atomicrmw_shared_l2_cache_policy(
+; CHECK: atom.relaxed.sys.shared.add.u32 %r2, [%rd1], %r1;
+ %old = atomicrmw add ptr addrspace(3) %p, i32 %v monotonic, align 4, !mem.cache_hint !1
+ ret i32 %old
+}
+
+define i128 @atomicrmw_xchg_l2_cache_policy_b128(ptr addrspace(1) %p, i128 %v) {
+; CHECK-LABEL: atomicrmw_xchg_l2_cache_policy_b128(
+; CHECK: mov.b64 %rd2, 12345;
+; CHECK: .reg .b128 amt, dst;
+; CHECK: atom.relaxed.sys.global.exch.L2::cache_hint.b128 dst, [%rd1], amt, %rd2;
+ %old = atomicrmw xchg ptr addrspace(1) %p, i128 %v monotonic, align 16, !mem.cache_hint !1
+ ret i128 %old
+}
+
!0 = !{i32 0, !10}
!1 = !{i32 0, !11}
!2 = !{i32 1, !11}
-
+!3 = !{i32 0, !12}
!10 = !{!"nvvm.l1_eviction", !"first"}
!11 = !{!"nvvm.l2_cache_hint", i64 12345}
+!12 = !{!"nvvm.l1_eviction", !"first", !"nvvm.l2_eviction", !"last", !"nvvm.l2_prefetch_size", !"128B", !"nvvm.l2_cache_hint", i64 12345}
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-sm-version.ll b/llvm/test/CodeGen/NVPTX/cache-hint-sm-version.ll
index 511a0fd4a41ca..eacd27f39f6f0 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-sm-version.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-sm-version.ll
@@ -26,23 +26,39 @@
define i32 @test_load_l1_first(ptr addrspace(1) %p) {
; SM60-LABEL: test_load_l1_first(
; SM60: ld.global.b32 %r1, [%rd1];
+; SM60: ld.volatile.global.b32 %r2, [%rd1];
;
; SM70-PTX73-LABEL: test_load_l1_first(
; SM70-PTX73: ld.global.b32 %r1, [%rd1];
+; SM70-PTX73: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM70-LABEL: test_load_l1_first(
; SM70: ld.global.L1::evict_first.b32 %r1, [%rd1];
+; SM70: ld.relaxed.sys.global.L1::evict_first.b32 %r2, [%rd1];
;
; SM75-LABEL: test_load_l1_first(
; SM75: ld.global.L1::evict_first.b32 %r1, [%rd1];
+; SM75: ld.relaxed.sys.global.L1::evict_first.b32 %r2, [%rd1];
;
; SM80PLUS-LABEL: test_load_l1_first(
; SM80PLUS: ld.global.L1::evict_first.b32 %r1, [%rd1];
+; SM80PLUS: ld.relaxed.sys.global.L1::evict_first.b32 %r2, [%rd1];
;
; SM80-PTX70-LABEL: test_load_l1_first(
; SM80-PTX70: ld.global.b32 %r1, [%rd1];
+; SM80-PTX70: ld.relaxed.sys.global.b32 %r2, [%rd1];
+;
+; SM100-PTX86-LABEL: test_load_l1_first(
+; SM100-PTX86: ld.global.L1::evict_first.b32 %r1, [%rd1];
+; SM100-PTX86: ld.relaxed.sys.global.L1::evict_first.b32 %r2, [%rd1];
+;
+; SM100-PTX88-LABEL: test_load_l1_first(
+; SM100-PTX88: ld.global.L1::evict_first.b32 %r1, [%rd1];
+; SM100-PTX88: ld.relaxed.sys.global.L1::evict_first.b32 %r2, [%rd1];
%v = load i32, ptr addrspace(1) %p, !mem.cache_hint !0
- ret i32 %v
+ %atomic = load atomic i32, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !0
+ %sum = add i32 %v, %atomic
+ ret i32 %sum
}
;-----------------------------------------------------------------------------
@@ -51,30 +67,53 @@ define i32 @test_load_l1_first(ptr addrspace(1) %p) {
define <8 x i32> @test_load_l2_last(ptr addrspace(1) %p) {
; SM60-LABEL: test_load_l2_last(
-; SM60: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; SM60: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];
+; SM60: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];
+; SM60: ld.volatile.global.b32 %r9, [%rd1];
;
; SM70-PTX73-LABEL: test_load_l2_last(
-; SM70-PTX73: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; SM70-PTX73: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];
+; SM70-PTX73: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];
+; SM70-PTX73: ld.relaxed.sys.global.b32 %r9, [%rd1];
;
; SM70-LABEL: test_load_l2_last(
-; SM70: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; SM70: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];
+; SM70: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];
+; SM70: ld.relaxed.sys.global.b32 %r9, [%rd1];
;
; SM75-LABEL: test_load_l2_last(
-; SM75: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; SM75: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];
+; SM75: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];
+; SM75: ld.relaxed.sys.global.b32 %r9, [%rd1];
;
; SM80PLUS-LABEL: test_load_l2_last(
-; SM80PLUS: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; SM80PLUS: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];
+; SM80PLUS: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];
+; SM80PLUS: ld.relaxed.sys.global.b32 %r9, [%rd1];
;
; SM80-PTX70-LABEL: test_load_l2_last(
-; SM80-PTX70: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; SM80-PTX70: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];
+; SM80-PTX70: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];
+; SM80-PTX70: ld.relaxed.sys.global.b32 %r9, [%rd1];
;
; SM100-PTX86-LABEL: test_load_l2_last(
-; SM100-PTX86: ld.global.v2.b64 {%rd2, %rd3}, [%rd1];
+; SM100-PTX86: ld.global.v2.b64 {%rd2, %rd3}, [%rd1+16];
+; SM100-PTX86: ld.global.v2.b64 {%rd4, %rd5}, [%rd1];
+; SM100-PTX86: mov.b64 {%r1, %r2}, %rd4;
+; SM100-PTX86: ld.relaxed.sys.global.b32 %r3, [%rd1];
+; SM100-PTX86: mov.b64 %rd6, {%r4, %r2};
;
; SM100-PTX88-LABEL: test_load_l2_last(
; SM100-PTX88: ld.global.L2::evict_last.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];
+; SM100-PTX88: mov.b64 {%r1, %r2}, %rd2;
+; SM100-PTX88: ld.relaxed.sys.global.b32 %r3, [%rd1];
+; SM100-PTX88: mov.b64 %rd6, {%r4, %r2};
%v = load <8 x i32>, ptr addrspace(1) %p, align 32, !mem.cache_hint !1
- ret <8 x i32> %v
+ %atomic = load atomic i32, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !1
+ %v0 = extractelement <8 x i32> %v, i64 0
+ %sum = add i32 %v0, %atomic
+ %result = insertelement <8 x i32> %v, i32 %sum, i64 0
+ ret <8 x i32> %result
}
;-----------------------------------------------------------------------------
@@ -86,23 +125,39 @@ define <8 x i32> @test_load_l2_last(ptr addrspace(1) %p) {
define i32 @test_load_prefetch_64(ptr addrspace(1) %p) {
; SM60-LABEL: test_load_prefetch_64(
; SM60: ld.global.b32 %r1, [%rd1];
+; SM60: ld.volatile.global.b32 %r2, [%rd1];
;
; SM70-PTX73-LABEL: test_load_prefetch_64(
; SM70-PTX73: ld.global.b32 %r1, [%rd1];
+; SM70-PTX73: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM70-LABEL: test_load_prefetch_64(
; SM70: ld.global.b32 %r1, [%rd1];
+; SM70: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM75-LABEL: test_load_prefetch_64(
; SM75: ld.global.L2::64B.b32 %r1, [%rd1];
+; SM75: ld.relaxed.sys.global.L2::64B.b32 %r2, [%rd1];
;
; SM80PLUS-LABEL: test_load_prefetch_64(
; SM80PLUS: ld.global.L2::64B.b32 %r1, [%rd1];
+; SM80PLUS: ld.relaxed.sys.global.L2::64B.b32 %r2, [%rd1];
;
; SM80-PTX70-LABEL: test_load_prefetch_64(
; SM80-PTX70: ld.global.b32 %r1, [%rd1];
+; SM80-PTX70: ld.relaxed.sys.global.b32 %r2, [%rd1];
+;
+; SM100-PTX86-LABEL: test_load_prefetch_64(
+; SM100-PTX86: ld.global.L2::64B.b32 %r1, [%rd1];
+; SM100-PTX86: ld.relaxed.sys.global.L2::64B.b32 %r2, [%rd1];
+;
+; SM100-PTX88-LABEL: test_load_prefetch_64(
+; SM100-PTX88: ld.global.L2::64B.b32 %r1, [%rd1];
+; SM100-PTX88: ld.relaxed.sys.global.L2::64B.b32 %r2, [%rd1];
%v = load i32, ptr addrspace(1) %p, !mem.cache_hint !6
- ret i32 %v
+ %atomic = load atomic i32, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !6
+ %sum = add i32 %v, %atomic
+ ret i32 %sum
}
;-----------------------------------------------------------------------------
@@ -114,23 +169,39 @@ define i32 @test_load_prefetch_64(ptr addrspace(1) %p) {
define i32 @test_load_prefetch_128(ptr addrspace(1) %p) {
; SM60-LABEL: test_load_prefetch_128(
; SM60: ld.global.b32 %r1, [%rd1];
+; SM60: ld.volatile.global.b32 %r2, [%rd1];
;
; SM70-PTX73-LABEL: test_load_prefetch_128(
; SM70-PTX73: ld.global.b32 %r1, [%rd1];
+; SM70-PTX73: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM70-LABEL: test_load_prefetch_128(
; SM70: ld.global.b32 %r1, [%rd1];
+; SM70: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM75-LABEL: test_load_prefetch_128(
; SM75: ld.global.L2::128B.b32 %r1, [%rd1];
+; SM75: ld.relaxed.sys.global.L2::128B.b32 %r2, [%rd1];
;
; SM80PLUS-LABEL: test_load_prefetch_128(
; SM80PLUS: ld.global.L2::128B.b32 %r1, [%rd1];
+; SM80PLUS: ld.relaxed.sys.global.L2::128B.b32 %r2, [%rd1];
;
; SM80-PTX70-LABEL: test_load_prefetch_128(
; SM80-PTX70: ld.global.b32 %r1, [%rd1];
+; SM80-PTX70: ld.relaxed.sys.global.b32 %r2, [%rd1];
+;
+; SM100-PTX86-LABEL: test_load_prefetch_128(
+; SM100-PTX86: ld.global.L2::128B.b32 %r1, [%rd1];
+; SM100-PTX86: ld.relaxed.sys.global.L2::128B.b32 %r2, [%rd1];
+;
+; SM100-PTX88-LABEL: test_load_prefetch_128(
+; SM100-PTX88: ld.global.L2::128B.b32 %r1, [%rd1];
+; SM100-PTX88: ld.relaxed.sys.global.L2::128B.b32 %r2, [%rd1];
%v = load i32, ptr addrspace(1) %p, !mem.cache_hint !2
- ret i32 %v
+ %atomic = load atomic i32, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !2
+ %sum = add i32 %v, %atomic
+ ret i32 %sum
}
;-----------------------------------------------------------------------------
@@ -142,23 +213,39 @@ define i32 @test_load_prefetch_128(ptr addrspace(1) %p) {
define i32 @test_load_prefetch_256(ptr addrspace(1) %p) {
; SM60-LABEL: test_load_prefetch_256(
; SM60: ld.global.b32 %r1, [%rd1];
+; SM60: ld.volatile.global.b32 %r2, [%rd1];
;
; SM70-PTX73-LABEL: test_load_prefetch_256(
; SM70-PTX73: ld.global.b32 %r1, [%rd1];
+; SM70-PTX73: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM70-LABEL: test_load_prefetch_256(
; SM70: ld.global.b32 %r1, [%rd1];
+; SM70: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM75-LABEL: test_load_prefetch_256(
; SM75: ld.global.b32 %r1, [%rd1];
+; SM75: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM80PLUS-LABEL: test_load_prefetch_256(
; SM80PLUS: ld.global.L2::256B.b32 %r1, [%rd1];
+; SM80PLUS: ld.relaxed.sys.global.L2::256B.b32 %r2, [%rd1];
;
; SM80-PTX70-LABEL: test_load_prefetch_256(
; SM80-PTX70: ld.global.b32 %r1, [%rd1];
+; SM80-PTX70: ld.relaxed.sys.global.b32 %r2, [%rd1];
+;
+; SM100-PTX86-LABEL: test_load_prefetch_256(
+; SM100-PTX86: ld.global.L2::256B.b32 %r1, [%rd1];
+; SM100-PTX86: ld.relaxed.sys.global.L2::256B.b32 %r2, [%rd1];
+;
+; SM100-PTX88-LABEL: test_load_prefetch_256(
+; SM100-PTX88: ld.global.L2::256B.b32 %r1, [%rd1];
+; SM100-PTX88: ld.relaxed.sys.global.L2::256B.b32 %r2, [%rd1];
%v = load i32, ptr addrspace(1) %p, !mem.cache_hint !7
- ret i32 %v
+ %atomic = load atomic i32, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !7
+ %sum = add i32 %v, %atomic
+ ret i32 %sum
}
;-----------------------------------------------------------------------------
@@ -171,24 +258,73 @@ define i32 @test_load_prefetch_256(ptr addrspace(1) %p) {
define i32 @test_load_cache_hint(ptr addrspace(1) %p) {
; SM60-LABEL: test_load_cache_hint(
; SM60: ld.global.b32 %r1, [%rd1];
+; SM60: ld.volatile.global.b32 %r2, [%rd1];
;
; SM70-PTX73-LABEL: test_load_cache_hint(
; SM70-PTX73: ld.global.b32 %r1, [%rd1];
+; SM70-PTX73: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM70-LABEL: test_load_cache_hint(
; SM70: ld.global.b32 %r1, [%rd1];
+; SM70: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM75-LABEL: test_load_cache_hint(
; SM75: ld.global.b32 %r1, [%rd1];
+; SM75: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM80PLUS-LABEL: test_load_cache_hint(
; SM80PLUS: mov.b64 %rd2, 12345;
; SM80PLUS: ld.global.L2::cache_hint.b32 %r1, [%rd1], %rd2;
+; SM80PLUS: ld.relaxed.sys.global.L2::cache_hint.b32 %r2, [%rd1], %rd2;
;
; SM80-PTX70-LABEL: test_load_cache_hint(
; SM80-PTX70: ld.global.b32 %r1, [%rd1];
+; SM80-PTX70: ld.relaxed.sys.global.b32 %r2, [%rd1];
+;
+; SM100-PTX86-LABEL: test_load_cache_hint(
+; SM100-PTX86: mov.b64 %rd2, 12345;
+; SM100-PTX86: ld.global.L2::cache_hint.b32 %r1, [%rd1], %rd2;
+; SM100-PTX86: ld.relaxed.sys.global.L2::cache_hint.b32 %r2, [%rd1], %rd2;
+;
+; SM100-PTX88-LABEL: test_load_cache_hint(
+; SM100-PTX88: mov.b64 %rd2, 12345;
+; SM100-PTX88: ld.global.L2::cache_hint.b32 %r1, [%rd1], %rd2;
+; SM100-PTX88: ld.relaxed.sys.global.L2::cache_hint.b32 %r2, [%rd1], %rd2;
%v = load i32, ptr addrspace(1) %p, !mem.cache_hint !3
- ret i32 %v
+ %atomic = load atomic i32, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !3
+ %sum = add i32 %v, %atomic
+ ret i32 %sum
+}
+
+define i32 @test_atomicrmw_cache_hint(ptr addrspace(1) %p, i32 %v) {
+; SM60-LABEL: test_atomicrmw_cache_hint(
+; SM60: atom.sys.global.add.u32 %r2, [%rd1], %r1;
+;
+; SM70-PTX73-LABEL: test_atomicrmw_cache_hint(
+; SM70-PTX73: atom.relaxed.sys.global.add.u32 %r2, [%rd1], %r1;
+;
+; SM70-LABEL: test_atomicrmw_cache_hint(
+; SM70: atom.relaxed.sys.global.add.u32 %r2, [%rd1], %r1;
+;
+; SM75-LABEL: test_atomicrmw_cache_hint(
+; SM75: atom.relaxed.sys.global.add.u32 %r2, [%rd1], %r1;
+;
+; SM80PLUS-LABEL: test_atomicrmw_cache_hint(
+; SM80PLUS: mov.b64 %rd2, 12345;
+; SM80PLUS: atom.relaxed.sys.global.add.L2::cache_hint.u32 %r2, [%rd1], %r1, %rd2;
+;
+; SM80-PTX70-LABEL: test_atomicrmw_cache_hint(
+; SM80-PTX70: atom.relaxed.sys.global.add.u32 %r2, [%rd1], %r1;
+;
+; SM100-PTX86-LABEL: test_atomicrmw_cache_hint(
+; SM100-PTX86: mov.b64 %rd2, 12345;
+; SM100-PTX86: atom.relaxed.sys.global.add.L2::cache_hint.u32 %r2, [%rd1], %r1, %rd2;
+;
+; SM100-PTX88-LABEL: test_atomicrmw_cache_hint(
+; SM100-PTX88: mov.b64 %rd2, 12345;
+; SM100-PTX88: atom.relaxed.sys.global.add.L2::cache_hint.u32 %r2, [%rd1], %r1, %rd2;
+ %old = atomicrmw add ptr addrspace(1) %p, i32 %v monotonic, align 4, !mem.cache_hint !3
+ ret i32 %old
}
;-----------------------------------------------------------------------------
@@ -201,24 +337,42 @@ define i32 @test_load_cache_hint(ptr addrspace(1) %p) {
define i32 @test_load_cache_hint_with_l1(ptr addrspace(1) %p) {
; SM60-LABEL: test_load_cache_hint_with_l1(
; SM60: ld.global.b32 %r1, [%rd1];
+; SM60: ld.volatile.global.b32 %r2, [%rd1];
;
; SM70-PTX73-LABEL: test_load_cache_hint_with_l1(
; SM70-PTX73: ld.global.b32 %r1, [%rd1];
+; SM70-PTX73: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM70-LABEL: test_load_cache_hint_with_l1(
; SM70: ld.global.L1::evict_first.b32 %r1, [%rd1];
+; SM70: ld.relaxed.sys.global.L1::evict_first.b32 %r2, [%rd1];
;
; SM75-LABEL: test_load_cache_hint_with_l1(
; SM75: ld.global.L1::evict_first.b32 %r1, [%rd1];
+; SM75: ld.relaxed.sys.global.L1::evict_first.b32 %r2, [%rd1];
;
; SM80PLUS-LABEL: test_load_cache_hint_with_l1(
; SM80PLUS: mov.b64 %rd2, 12345;
; SM80PLUS: ld.global.L1::evict_first.L2::cache_hint.b32 %r1, [%rd1], %rd2;
+; SM80PLUS: ld.relaxed.sys.global.L1::evict_first.L2::cache_hint.b32 %r2, [%rd1], %rd2;
;
; SM80-PTX70-LABEL: test_load_cache_hint_with_l1(
; SM80-PTX70: ld.global.b32 %r1, [%rd1];
+; SM80-PTX70: ld.relaxed.sys.global.b32 %r2, [%rd1];
+;
+; SM100-PTX86-LABEL: test_load_cache_hint_with_l1(
+; SM100-PTX86: mov.b64 %rd2, 12345;
+; SM100-PTX86: ld.global.L1::evict_first.L2::cache_hint.b32 %r1, [%rd1], %rd2;
+; SM100-PTX86: ld.relaxed.sys.global.L1::evict_first.L2::cache_hint.b32 %r2, [%rd1], %rd2;
+;
+; SM100-PTX88-LABEL: test_load_cache_hint_with_l1(
+; SM100-PTX88: mov.b64 %rd2, 12345;
+; SM100-PTX88: ld.global.L1::evict_first.L2::cache_hint.b32 %r1, [%rd1], %rd2;
+; SM100-PTX88: ld.relaxed.sys.global.L1::evict_first.L2::cache_hint.b32 %r2, [%rd1], %rd2;
%v = load i32, ptr addrspace(1) %p, !mem.cache_hint !4
- ret i32 %v
+ %atomic = load atomic i32, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !4
+ %sum = add i32 %v, %atomic
+ ret i32 %sum
}
;-----------------------------------------------------------------------------
@@ -231,23 +385,39 @@ define i32 @test_load_cache_hint_with_l1(ptr addrspace(1) %p) {
define i32 @test_load_prefetch_with_l1(ptr addrspace(1) %p) {
; SM60-LABEL: test_load_prefetch_with_l1(
; SM60: ld.global.b32 %r1, [%rd1];
+; SM60: ld.volatile.global.b32 %r2, [%rd1];
;
; SM70-PTX73-LABEL: test_load_prefetch_with_l1(
; SM70-PTX73: ld.global.b32 %r1, [%rd1];
+; SM70-PTX73: ld.relaxed.sys.global.b32 %r2, [%rd1];
;
; SM70-LABEL: test_load_prefetch_with_l1(
; SM70: ld.global.L1::evict_first.b32 %r1, [%rd1];
+; SM70: ld.relaxed.sys.global.L1::evict_first.b32 %r2, [%rd1];
;
; SM75-LABEL: test_load_prefetch_with_l1(
; SM75: ld.global.L1::evict_first.L2::128B.b32 %r1, [%rd1];
+; SM75: ld.relaxed.sys.global.L1::evict_first.L2::128B.b32 %r2, [%rd1];
;
; SM80PLUS-LABEL: test_load_prefetch_with_l1(
; SM80PLUS: ld.global.L1::evict_first.L2::128B.b32 %r1, [%rd1];
+; SM80PLUS: ld.relaxed.sys.global.L1::evict_first.L2::128B.b32 %r2, [%rd1];
;
; SM80-PTX70-LABEL: test_load_prefetch_with_l1(
; SM80-PTX70: ld.global.b32 %r1, [%rd1];
+; SM80-PTX70: ld.relaxed.sys.global.b32 %r2, [%rd1];
+;
+; SM100-PTX86-LABEL: test_load_prefetch_with_l1(
+; SM100-PTX86: ld.global.L1::evict_first.L2::128B.b32 %r1, [%rd1];
+; SM100-PTX86: ld.relaxed.sys.global.L1::evict_first.L2::128B.b32 %r2, [%rd1];
+;
+; SM100-PTX88-LABEL: test_load_prefetch_with_l1(
+; SM100-PTX88: ld.global.L1::evict_first.L2::128B.b32 %r1, [%rd1];
+; SM100-PTX88: ld.relaxed.sys.global.L1::evict_first.L2::128B.b32 %r2, [%rd1];
%v = load i32, ptr addrspace(1) %p, !mem.cache_hint !8
- ret i32 %v
+ %atomic = load atomic i32, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !8
+ %sum = add i32 %v, %atomic
+ ret i32 %sum
}
;-----------------------------------------------------------------------------
@@ -257,23 +427,40 @@ define i32 @test_load_prefetch_with_l1(ptr addrspace(1) %p) {
define void @test_store_cache_hint(ptr addrspace(1) %p, i32 %v) {
; SM60-LABEL: test_store_cache_hint(
; SM60: st.global.b32 [%rd1], %r1;
+; SM60: st.volatile.global.b32 [%rd1], %r1;
;
; SM70-PTX73-LABEL: test_store_cache_hint(
; SM70-PTX73: st.global.b32 [%rd1], %r1;
+; SM70-PTX73: st.relaxed.sys.global.b32 [%rd1], %r1;
;
; SM70-LABEL: test_store_cache_hint(
; SM70: st.global.b32 [%rd1], %r1;
+; SM70: st.relaxed.sys.global.b32 [%rd1], %r1;
;
; SM75-LABEL: test_store_cache_hint(
; SM75: st.global.b32 [%rd1], %r1;
+; SM75: st.relaxed.sys.global.b32 [%rd1], %r1;
;
; SM80PLUS-LABEL: test_store_cache_hint(
; SM80PLUS: mov.b64 %rd2, 12345;
; SM80PLUS: st.global.L2::cache_hint.b32 [%rd1], %r1, %rd2;
+; SM80PLUS: st.relaxed.sys.global.L2::cache_hint.b32 [%rd1], %r1, %rd2;
;
; SM80-PTX70-LABEL: test_store_cache_hint(
; SM80-PTX70: st.global.b32 [%rd1], %r1;
+; SM80-PTX70: st.relaxed.sys.global.b32 [%rd1], %r1;
+;
+; SM100-PTX86-LABEL: test_store_cache_hint(
+; SM100-PTX86: mov.b64 %rd2, 12345;
+; SM100-PTX86: st.global.L2::cache_hint.b32 [%rd1], %r1, %rd2;
+; SM100-PTX86: st.relaxed.sys.global.L2::cache_hint.b32 [%rd1], %r1, %rd2;
+;
+; SM100-PTX88-LABEL: test_store_cache_hint(
+; SM100-PTX88: mov.b64 %rd2, 12345;
+; SM100-PTX88: st.global.L2::cache_hint.b32 [%rd1], %r1, %rd2;
+; SM100-PTX88: st.relaxed.sys.global.L2::cache_hint.b32 [%rd1], %r1, %rd2;
store i32 %v, ptr addrspace(1) %p, !mem.cache_hint !5
+ store atomic i32 %v, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !5
ret void
}
@@ -284,22 +471,37 @@ define void @test_store_cache_hint(ptr addrspace(1) %p, i32 %v) {
define void @test_store_l1_no_allocate(ptr addrspace(1) %p, i32 %v) {
; SM60-LABEL: test_store_l1_no_allocate(
; SM60: st.global.b32 [%rd1], %r1;
+; SM60: st.volatile.global.b32 [%rd1], %r1;
;
; SM70-PTX73-LABEL: test_store_l1_no_allocate(
; SM70-PTX73: st.global.b32 [%rd1], %r1;
+; SM70-PTX73: st.relaxed.sys.global.b32 [%rd1], %r1;
;
; SM70-LABEL: test_store_l1_no_allocate(
; SM70: st.global.L1::no_allocate.b32 [%rd1], %r1;
+; SM70: st.relaxed.sys.global.L1::no_allocate.b32 [%rd1], %r1;
;
; SM75-LABEL: test_store_l1_no_allocate(
; SM75: st.global.L1::no_allocate.b32 [%rd1], %r1;
+; SM75: st.relaxed.sys.global.L1::no_allocate.b32 [%rd1], %r1;
;
; SM80PLUS-LABEL: test_store_l1_no_allocate(
; SM80PLUS: st.global.L1::no_allocate.b32 [%rd1], %r1;
+; SM80PLUS: st.relaxed.sys.global.L1::no_allocate.b32 [%rd1], %r1;
;
; SM80-PTX70-LABEL: test_store_l1_no_allocate(
; SM80-PTX70: st.global.b32 [%rd1], %r1;
+; SM80-PTX70: st.relaxed.sys.global.b32 [%rd1], %r1;
+;
+; SM100-PTX86-LABEL: test_store_l1_no_allocate(
+; SM100-PTX86: st.global.L1::no_allocate.b32 [%rd1], %r1;
+; SM100-PTX86: st.relaxed.sys.global.L1::no_allocate.b32 [%rd1], %r1;
+;
+; SM100-PTX88-LABEL: test_store_l1_no_allocate(
+; SM100-PTX88: st.global.L1::no_allocate.b32 [%rd1], %r1;
+; SM100-PTX88: st.relaxed.sys.global.L1::no_allocate.b32 [%rd1], %r1;
store i32 %v, ptr addrspace(1) %p, !mem.cache_hint !9
+ store atomic i32 %v, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !9
ret void
}
>From 799c9e33898bfe31c8febf9f32239fe9c0068152 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Sat, 8 Aug 2026 02:56:56 +0000
Subject: [PATCH 2/6] small change
---
llvm/test/CodeGen/NVPTX/cache-hint-atomics.ll | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-atomics.ll b/llvm/test/CodeGen/NVPTX/cache-hint-atomics.ll
index e768374ed6033..44f2645c14569 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-atomics.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-atomics.ll
@@ -46,7 +46,7 @@ define i32 @atomicrmw_xchg_l2_cache_policy(ptr addrspace(1) %p) {
define i32 @cmpxchg_l2_cache_policy(ptr addrspace(1) %p, i32 %cmp, i32 %new) {
; CHECK-LABEL: cmpxchg_l2_cache_policy(
; CHECK: atom.relaxed.sys.global.cas.b32 %r3, [%rd1], %r1, %r2;
- %pair = cmpxchg ptr addrspace(1) %p, i32 %cmp, i32 %new monotonic monotonic, align 4, !mem.cache_hint !1
+ %pair = cmpxchg ptr addrspace(1) %p, i32 %cmp, i32 %new monotonic monotonic, align 4, !mem.cache_hint !3
%old = extractvalue { i32, i1 } %pair, 0
ret i32 %old
}
>From ed1db39fd5e5ed5bfce6bb78b096c09463d229a0 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Mon, 24 Aug 2026 20:26:36 +0000
Subject: [PATCH 3/6] fix emitting two cache hint warnings
---
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 21 ++++++++++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 40 ++++++++-----------
llvm/test/CodeGen/NVPTX/cache-hint-invalid.ll | 30 ++++++++++++++
3 files changed, 67 insertions(+), 24 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index cc293ee25ff40..490059b2e53f9 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -148,6 +148,9 @@ class NVPTXDAGToDAGISel : public SelectionDAGISel {
NVPTX::Scope getAtomicScope(const MemSDNode *N) const;
bool SelectADDR(SDValue Addr, SDValue &Base, SDValue &Offset);
+ bool SelectAtomicADDR(SDNode *Root, SDValue Addr, SDValue &Base,
+ SDValue &Offset, SDValue &EvictionAndPrefetchHint,
+ SDValue &CachePolicyReg);
SDValue getPTXCmpMode(const CondCodeSDNode &CondCode);
SDValue selectPossiblyImm(SDValue V);
@@ -1331,6 +1334,24 @@ NVPTXMemCacheHintOperands NVPTXDAGToDAGISel::getMemCacheHintOperands(
return {getI32Imm(EvictionAndPrefetchHint, DL), PolicyReg};
}
+bool NVPTXDAGToDAGISel::SelectAtomicADDR(
+ SDNode *Root, SDValue Addr, SDValue &Base, SDValue &Offset,
+ SDValue &EvictionAndPrefetchHint, SDValue &CachePolicyReg) {
+ if (!SelectADDR(Addr, Base, Offset))
+ return false;
+
+ auto *MN = cast<MemSDNode>(Root);
+ unsigned EltWidth = MN->getMemoryVT().getFixedSizeInBits();
+ NVPTXMemCacheHintAccess Access{
+ NVPTXMemCacheHintInstruction::Atom, getAddrSpace(MN),
+ /*NumElts=*/1, EltWidth, MN->isVolatile()};
+ NVPTXMemCacheHintOperands CacheHintOperands =
+ getMemCacheHintOperands(MN, Access, SDLoc(Root));
+ EvictionAndPrefetchHint = CacheHintOperands.EvictionAndPrefetchHint;
+ CachePolicyReg = CacheHintOperands.CachePolicyReg;
+ return true;
+}
+
bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) {
MemSDNode *LD = cast<MemSDNode>(N);
assert(LD->readMem() && "Expected load");
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index f9a79d9a777d5..50d1de1566960 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2693,25 +2693,10 @@ class AtomicSDNodeXForm<SDNode atomic> {
SDNodeXForm GetAddSp = SDNodeXForm<atomic, [{
return getI32Imm(getAddrSpace(cast<MemSDNode>(N)), SDLoc(N));
}]>;
+}
- SDNodeXForm GetEvictionAndPrefetchHint = SDNodeXForm<atomic, [{
- auto *MN = cast<MemSDNode>(N);
- unsigned EltWidth = MN->getMemoryVT().getFixedSizeInBits();
- NVPTXMemCacheHintAccess Access{
- NVPTXMemCacheHintInstruction::Atom, getAddrSpace(MN),
- /*NumElts=*/1, EltWidth, MN->isVolatile()};
- return getMemCacheHintOperands(MN, Access, SDLoc(N))
- .EvictionAndPrefetchHint;
- }]>;
-
- SDNodeXForm GetCachePolicy = SDNodeXForm<atomic, [{
- auto *MN = cast<MemSDNode>(N);
- unsigned EltWidth = MN->getMemoryVT().getFixedSizeInBits();
- NVPTXMemCacheHintAccess Access{
- NVPTXMemCacheHintInstruction::Atom, getAddrSpace(MN),
- /*NumElts=*/1, EltWidth, MN->isVolatile()};
- return getMemCacheHintOperands(MN, Access, SDLoc(N)).CachePolicyReg;
- }]>;
+def atomic_addr : ComplexPattern<pAny, 4, "SelectAtomicADDR"> {
+ let WantsRoot = true;
}
multiclass F_ATOMIC_2<RegTyInfo t, SDPatternOperator op, string op_str,
@@ -2740,18 +2725,25 @@ multiclass F_ATOMIC_2<RegTyInfo t, SDPatternOperator op, string op_str,
defvar XForm = AtomicSDNodeXForm<atomic>;
- def : Pat<(op:$this addr:$addr, t.Ty:$b),
+ def : Pat<(op:$this
+ (atomic_addr ADDR:$addr,
+ EvictionAndPrefetchHint:$evictionAndPrefetchHint,
+ CachePolicy:$policy),
+ t.Ty:$b),
(!cast<Instruction>(NAME # _r) ADDR:$addr, t.Ty:$b, (XForm.GetSem $this),
(XForm.GetScope $this), (XForm.GetAddSp $this),
- (XForm.GetEvictionAndPrefetchHint $this),
- (XForm.GetCachePolicy $this))>;
+ EvictionAndPrefetchHint:$evictionAndPrefetchHint, CachePolicy:$policy)>;
if t.SupportsImm then
- def : Pat<(op:$this addr:$addr, (t.Ty t.ImmNode:$b)),
+ def : Pat<(op:$this
+ (atomic_addr ADDR:$addr,
+ EvictionAndPrefetchHint:$evictionAndPrefetchHint,
+ CachePolicy:$policy),
+ (t.Ty t.ImmNode:$b)),
(!cast<Instruction>(NAME # _i) ADDR:$addr, (t.Ty t.ImmNode:$b),
(XForm.GetSem $this), (XForm.GetScope $this), (XForm.GetAddSp $this),
- (XForm.GetEvictionAndPrefetchHint $this),
- (XForm.GetCachePolicy $this))>;
+ EvictionAndPrefetchHint:$evictionAndPrefetchHint,
+ CachePolicy:$policy)>;
}
multiclass F_ATOMIC_3<RegTyInfo t, string op_str, SDPatternOperator op, SDNode atomic> {
diff --git a/llvm/test/CodeGen/NVPTX/cache-hint-invalid.ll b/llvm/test/CodeGen/NVPTX/cache-hint-invalid.ll
index e17e4b0b496b8..adab9aede86e8 100644
--- a/llvm/test/CodeGen/NVPTX/cache-hint-invalid.ll
+++ b/llvm/test/CodeGen/NVPTX/cache-hint-invalid.ll
@@ -12,6 +12,7 @@
; RUN: llc < %t/bad-policy.ll -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 2>&1 | FileCheck %s --check-prefix=BAD-POLICY
; RUN: llc < %t/bad-policy.ll -mtriple=nvptx64 -mcpu=sm_60 -mattr=+ptx50 2>&1 | FileCheck %s --check-prefix=BAD-POLICY
; RUN: llc < %t/bad-policy-shared.ll -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx74 2>&1 | FileCheck %s --check-prefix=BAD-POLICY-SHARED
+; RUN: llc < %t/bad-policy-atomics.ll -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx88 2>&1 | FileCheck %s --check-prefix=BAD-POLICY-ATOMICS --implicit-check-not=warning:
;--- bad-empty.ll
@@ -133,6 +134,35 @@ define i32 @bad_cache_policy_value(ptr addrspace(1) %p) {
!0 = !{i32 0, !1}
!1 = !{!"nvvm.l2_cache_hint", !"not_an_integer"}
+;--- bad-policy-atomics.ll
+
+; Invalid cache policies on each supported kind of atomic memory operation
+; should produce exactly one warning per instruction.
+; BAD-POLICY-ATOMICS-COUNT-4: warning: invalid NVPTX !mem.cache_hint metadata: 'nvvm.l2_cache_hint' expects an integer value
+define i32 @bad_atomic_load_cache_policy(ptr addrspace(1) %p) {
+ %v = load atomic i32, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !0
+ ret i32 %v
+}
+
+define void @bad_atomic_store_cache_policy(ptr addrspace(1) %p, i32 %v) {
+ store atomic i32 %v, ptr addrspace(1) %p monotonic, align 4, !mem.cache_hint !2
+ ret void
+}
+
+define i32 @bad_atomicrmw_cache_policy(ptr addrspace(1) %p, i32 %v) {
+ %old = atomicrmw add ptr addrspace(1) %p, i32 %v monotonic, align 4, !mem.cache_hint !0
+ ret i32 %old
+}
+
+define i128 @bad_atomicrmw_i128_cache_policy(ptr addrspace(1) %p, i128 %v) {
+ %old = atomicrmw xchg ptr addrspace(1) %p, i128 %v monotonic, align 16, !mem.cache_hint !0
+ ret i128 %old
+}
+
+!0 = !{i32 0, !1}
+!1 = !{!"nvvm.l2_cache_hint", !"not_an_integer"}
+!2 = !{i32 1, !1}
+
;--- bad-policy-shared.ll
; nvvm.l2_cache_hint expects an integer cache-policy value even when the
>From dc8f6f22707955216f3df9786d0d47a2d9754428 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Mon, 24 Aug 2026 20:44:00 +0000
Subject: [PATCH 4/6] format
---
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 13 +++++++------
1 file changed, 7 insertions(+), 6 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 490059b2e53f9..0115d61fd8b0c 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -1334,17 +1334,18 @@ NVPTXMemCacheHintOperands NVPTXDAGToDAGISel::getMemCacheHintOperands(
return {getI32Imm(EvictionAndPrefetchHint, DL), PolicyReg};
}
-bool NVPTXDAGToDAGISel::SelectAtomicADDR(
- SDNode *Root, SDValue Addr, SDValue &Base, SDValue &Offset,
- SDValue &EvictionAndPrefetchHint, SDValue &CachePolicyReg) {
+bool NVPTXDAGToDAGISel::SelectAtomicADDR(SDNode *Root, SDValue Addr,
+ SDValue &Base, SDValue &Offset,
+ SDValue &EvictionAndPrefetchHint,
+ SDValue &CachePolicyReg) {
if (!SelectADDR(Addr, Base, Offset))
return false;
auto *MN = cast<MemSDNode>(Root);
unsigned EltWidth = MN->getMemoryVT().getFixedSizeInBits();
- NVPTXMemCacheHintAccess Access{
- NVPTXMemCacheHintInstruction::Atom, getAddrSpace(MN),
- /*NumElts=*/1, EltWidth, MN->isVolatile()};
+ NVPTXMemCacheHintAccess Access{NVPTXMemCacheHintInstruction::Atom,
+ getAddrSpace(MN),
+ /*NumElts=*/1, EltWidth, MN->isVolatile()};
NVPTXMemCacheHintOperands CacheHintOperands =
getMemCacheHintOperands(MN, Access, SDLoc(Root));
EvictionAndPrefetchHint = CacheHintOperands.EvictionAndPrefetchHint;
>From 37350f652e421b3a143a8083cec391219f709dab Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Tue, 25 Aug 2026 23:48:02 +0000
Subject: [PATCH 5/6] emitDiagnostics parameter
---
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 64 ++++++++-------------
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 43 ++++++++------
2 files changed, 52 insertions(+), 55 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 0115d61fd8b0c..801a2af35a576 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -148,9 +148,6 @@ class NVPTXDAGToDAGISel : public SelectionDAGISel {
NVPTX::Scope getAtomicScope(const MemSDNode *N) const;
bool SelectADDR(SDValue Addr, SDValue &Base, SDValue &Offset);
- bool SelectAtomicADDR(SDNode *Root, SDValue Addr, SDValue &Base,
- SDValue &Offset, SDValue &EvictionAndPrefetchHint,
- SDValue &CachePolicyReg);
SDValue getPTXCmpMode(const CondCodeSDNode &CondCode);
SDValue selectPossiblyImm(SDValue V);
@@ -161,7 +158,7 @@ class NVPTXDAGToDAGISel : public SelectionDAGISel {
// value. Otherwise returns NOREG for the policy operand.
NVPTXMemCacheHintOperands
getMemCacheHintOperands(const MemSDNode *N, NVPTXMemCacheHintAccess Access,
- const SDLoc &DL);
+ const SDLoc &DL, bool EmitDiagnostics = true);
// Returns the Memory Order and Scope that the PTX memory instruction should
// use, and inserts appropriate fence instruction before the memory
@@ -1190,16 +1187,19 @@ template <typename T>
static std::optional<T>
parseMemCacheHintStringValue(LLVMContext &Ctx, StringRef Key,
const Metadata *Value,
- std::optional<T> (*Parse)(StringRef)) {
+ std::optional<T> (*Parse)(StringRef),
+ bool EmitDiagnostics) {
const auto *Val = dyn_cast<MDString>(Value);
if (!Val) {
- emitInvalidMemCacheHint(Ctx, Twine("'") + Key + "' expects a string value");
+ if (EmitDiagnostics)
+ emitInvalidMemCacheHint(Ctx,
+ Twine("'") + Key + "' expects a string value");
return std::nullopt;
}
StringRef ValStr = Val->getString();
auto Parsed = Parse(ValStr);
- if (!Parsed)
+ if (!Parsed && EmitDiagnostics)
emitInvalidMemCacheHint(Ctx, Twine("unknown value '") + ValStr + "' for '" +
Key + "'");
return Parsed;
@@ -1263,14 +1263,16 @@ static bool isCachePolicySupported(const NVPTXSubtarget &Subtarget,
}
NVPTXMemCacheHintOperands NVPTXDAGToDAGISel::getMemCacheHintOperands(
- const MemSDNode *N, NVPTXMemCacheHintAccess Access, const SDLoc &DL) {
+ const MemSDNode *N, NVPTXMemCacheHintAccess Access, const SDLoc &DL,
+ bool EmitDiagnostics) {
LLVMContext &Ctx = *CurDAG->getContext();
const MDNode *Node = N->getMemCacheHint();
SDValue PolicyReg = CurDAG->getRegister(NVPTX::NoRegister, MVT::i64);
if (!Node)
return {getI32Imm(0, DL), PolicyReg};
if (Node->getNumOperands() == 0) {
- emitInvalidMemCacheHint(Ctx, "empty hint node");
+ if (EmitDiagnostics)
+ emitInvalidMemCacheHint(Ctx, "empty hint node");
return {getI32Imm(0, DL), PolicyReg};
}
@@ -1285,24 +1287,24 @@ NVPTXMemCacheHintOperands NVPTXDAGToDAGISel::getMemCacheHintOperands(
const Metadata *Value = Node->getOperand(I + 1).get();
if (KeyStr == "nvvm.l1_eviction") {
- auto ParsedL1 =
- parseMemCacheHintStringValue(Ctx, KeyStr, Value, parseL1Eviction);
+ auto ParsedL1 = parseMemCacheHintStringValue(
+ Ctx, KeyStr, Value, parseL1Eviction, EmitDiagnostics);
if (ParsedL1 && isL1EvictionSupported(*Subtarget, *ParsedL1, Access))
L1 = *ParsedL1;
continue;
}
if (KeyStr == "nvvm.l2_eviction") {
- auto ParsedL2 =
- parseMemCacheHintStringValue(Ctx, KeyStr, Value, parseL2Eviction);
+ auto ParsedL2 = parseMemCacheHintStringValue(
+ Ctx, KeyStr, Value, parseL2Eviction, EmitDiagnostics);
if (ParsedL2 && isL2EvictionSupported(*Subtarget, *ParsedL2, Access))
L2 = *ParsedL2;
continue;
}
if (KeyStr == "nvvm.l2_prefetch_size") {
- auto ParsedPrefetch =
- parseMemCacheHintStringValue(Ctx, KeyStr, Value, parseL2Prefetch);
+ auto ParsedPrefetch = parseMemCacheHintStringValue(
+ Ctx, KeyStr, Value, parseL2Prefetch, EmitDiagnostics);
if (ParsedPrefetch &&
isL2PrefetchSupported(*Subtarget, *ParsedPrefetch, Access))
Prefetch = *ParsedPrefetch;
@@ -1311,15 +1313,18 @@ NVPTXMemCacheHintOperands NVPTXDAGToDAGISel::getMemCacheHintOperands(
if (KeyStr == "nvvm.l2_cache_hint") {
const auto *ValCI = mdconst::dyn_extract<ConstantInt>(Value);
- if (!ValCI)
- emitInvalidMemCacheHint(
- Ctx, "'nvvm.l2_cache_hint' expects an integer value");
- else if (isCachePolicySupported(*Subtarget, Access))
+ if (!ValCI) {
+ if (EmitDiagnostics)
+ emitInvalidMemCacheHint(
+ Ctx, "'nvvm.l2_cache_hint' expects an integer value");
+ } else if (isCachePolicySupported(*Subtarget, Access)) {
CachePolicy = ValCI->getZExtValue();
+ }
continue;
}
- emitInvalidMemCacheHint(Ctx, Twine("unknown key '") + KeyStr + "'");
+ if (EmitDiagnostics)
+ emitInvalidMemCacheHint(Ctx, Twine("unknown key '") + KeyStr + "'");
}
unsigned EvictionAndPrefetchHint =
@@ -1334,25 +1339,6 @@ NVPTXMemCacheHintOperands NVPTXDAGToDAGISel::getMemCacheHintOperands(
return {getI32Imm(EvictionAndPrefetchHint, DL), PolicyReg};
}
-bool NVPTXDAGToDAGISel::SelectAtomicADDR(SDNode *Root, SDValue Addr,
- SDValue &Base, SDValue &Offset,
- SDValue &EvictionAndPrefetchHint,
- SDValue &CachePolicyReg) {
- if (!SelectADDR(Addr, Base, Offset))
- return false;
-
- auto *MN = cast<MemSDNode>(Root);
- unsigned EltWidth = MN->getMemoryVT().getFixedSizeInBits();
- NVPTXMemCacheHintAccess Access{NVPTXMemCacheHintInstruction::Atom,
- getAddrSpace(MN),
- /*NumElts=*/1, EltWidth, MN->isVolatile()};
- NVPTXMemCacheHintOperands CacheHintOperands =
- getMemCacheHintOperands(MN, Access, SDLoc(Root));
- EvictionAndPrefetchHint = CacheHintOperands.EvictionAndPrefetchHint;
- CachePolicyReg = CacheHintOperands.CachePolicyReg;
- return true;
-}
-
bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) {
MemSDNode *LD = cast<MemSDNode>(N);
assert(LD->readMem() && "Expected load");
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 50d1de1566960..520a567c600f9 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2693,10 +2693,28 @@ class AtomicSDNodeXForm<SDNode atomic> {
SDNodeXForm GetAddSp = SDNodeXForm<atomic, [{
return getI32Imm(getAddrSpace(cast<MemSDNode>(N)), SDLoc(N));
}]>;
-}
-def atomic_addr : ComplexPattern<pAny, 4, "SelectAtomicADDR"> {
- let WantsRoot = true;
+ SDNodeXForm GetEvictionAndPrefetchHint = SDNodeXForm<atomic, [{
+ auto *MN = cast<MemSDNode>(N);
+ unsigned EltWidth = MN->getMemoryVT().getFixedSizeInBits();
+ NVPTXMemCacheHintAccess Access{
+ NVPTXMemCacheHintInstruction::Atom, getAddrSpace(MN),
+ /*NumElts=*/1, EltWidth, MN->isVolatile()};
+ return getMemCacheHintOperands(MN, Access, SDLoc(N),
+ /*EmitDiagnostics=*/true)
+ .EvictionAndPrefetchHint;
+ }]>;
+
+ SDNodeXForm GetCachePolicy = SDNodeXForm<atomic, [{
+ auto *MN = cast<MemSDNode>(N);
+ unsigned EltWidth = MN->getMemoryVT().getFixedSizeInBits();
+ NVPTXMemCacheHintAccess Access{
+ NVPTXMemCacheHintInstruction::Atom, getAddrSpace(MN),
+ /*NumElts=*/1, EltWidth, MN->isVolatile()};
+ return getMemCacheHintOperands(MN, Access, SDLoc(N),
+ /*EmitDiagnostics=*/false)
+ .CachePolicyReg;
+ }]>;
}
multiclass F_ATOMIC_2<RegTyInfo t, SDPatternOperator op, string op_str,
@@ -2725,25 +2743,18 @@ multiclass F_ATOMIC_2<RegTyInfo t, SDPatternOperator op, string op_str,
defvar XForm = AtomicSDNodeXForm<atomic>;
- def : Pat<(op:$this
- (atomic_addr ADDR:$addr,
- EvictionAndPrefetchHint:$evictionAndPrefetchHint,
- CachePolicy:$policy),
- t.Ty:$b),
+ def : Pat<(op:$this addr:$addr, t.Ty:$b),
(!cast<Instruction>(NAME # _r) ADDR:$addr, t.Ty:$b, (XForm.GetSem $this),
(XForm.GetScope $this), (XForm.GetAddSp $this),
- EvictionAndPrefetchHint:$evictionAndPrefetchHint, CachePolicy:$policy)>;
+ (XForm.GetEvictionAndPrefetchHint $this),
+ (XForm.GetCachePolicy $this))>;
if t.SupportsImm then
- def : Pat<(op:$this
- (atomic_addr ADDR:$addr,
- EvictionAndPrefetchHint:$evictionAndPrefetchHint,
- CachePolicy:$policy),
- (t.Ty t.ImmNode:$b)),
+ def : Pat<(op:$this addr:$addr, (t.Ty t.ImmNode:$b)),
(!cast<Instruction>(NAME # _i) ADDR:$addr, (t.Ty t.ImmNode:$b),
(XForm.GetSem $this), (XForm.GetScope $this), (XForm.GetAddSp $this),
- EvictionAndPrefetchHint:$evictionAndPrefetchHint,
- CachePolicy:$policy)>;
+ (XForm.GetEvictionAndPrefetchHint $this),
+ (XForm.GetCachePolicy $this))>;
}
multiclass F_ATOMIC_3<RegTyInfo t, string op_str, SDPatternOperator op, SDNode atomic> {
>From e01f0992c11da69113958f456c1351055f4ac7b9 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Tue, 25 Aug 2026 23:48:35 +0000
Subject: [PATCH 6/6] remove SM90
---
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 520a567c600f9..1b9145cc074d7 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2804,7 +2804,7 @@ defm INT_PTX_ATOM_ADD_32 : F_ATOMIC_2<I32RT, atomic_load_add_i32, "add", "u32",
defm INT_PTX_ATOM_ADD_64 : F_ATOMIC_2<I64RT, atomic_load_add_i64, "add", "u64", atomic_load_add>;
defm INT_PTX_ATOM_ADD_F16 : F_ATOMIC_2<F16RT, atomic_load_fadd, "add.noftz", "f16", atomic_load_fadd, [SM70, PTX63]>;
-defm INT_PTX_ATOM_ADD_BF16 : F_ATOMIC_2<BF16RT, atomic_load_fadd, "add.noftz", "bf16", atomic_load_fadd, [SM90, SM90]>;
+defm INT_PTX_ATOM_ADD_BF16 : F_ATOMIC_2<BF16RT, atomic_load_fadd, "add.noftz", "bf16", atomic_load_fadd, [SM90]>;
defm INT_PTX_ATOM_ADD_F32 : F_ATOMIC_2<F32RT, atomic_load_fadd, "add", "f32", atomic_load_fadd>;
defm INT_PTX_ATOM_ADD_F64 : F_ATOMIC_2<F64RT, atomic_load_fadd, "add", "f64", atomic_load_fadd, [hasAtomAddF64]>;
More information about the llvm-commits
mailing list