[llvm] [AtomicExpand] Implement SUB → ADD(-x) and FSUB → FADD(-x) (PR #221425)
Yonah Goldberg via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 30 19:06:13 PDT 2026
https://github.com/YonahGoldberg updated https://github.com/llvm/llvm-project/pull/221425
>From 5a6ee2988e51bf2e1f6f3958f1a74af18dc058ee Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Sat, 5 Sep 2026 08:43:10 +0000
Subject: [PATCH 1/2] atomic expand
---
llvm/lib/CodeGen/AtomicExpandPass.cpp | 26 ++
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 43 ++-
llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll | 82 ++----
llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll | 222 +++------------
llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll | 257 ++++--------------
.../AtomicExpand/NVPTX/atomicrmw-sub.ll | 96 +++++++
6 files changed, 258 insertions(+), 468 deletions(-)
create mode 100644 llvm/test/Transforms/AtomicExpand/NVPTX/atomicrmw-sub.ll
diff --git a/llvm/lib/CodeGen/AtomicExpandPass.cpp b/llvm/lib/CodeGen/AtomicExpandPass.cpp
index f2ffa40030cc0..c595de6c380b5 100644
--- a/llvm/lib/CodeGen/AtomicExpandPass.cpp
+++ b/llvm/lib/CodeGen/AtomicExpandPass.cpp
@@ -105,6 +105,7 @@ class AtomicExpandImpl {
bool tryExpandAtomicStore(StoreInst *SI);
void expandAtomicStoreToXChg(StoreInst *SI);
bool tryExpandAtomicRMW(AtomicRMWInst *AI);
+ void expandAtomicSubToAdd(AtomicRMWInst *AI);
AtomicRMWInst *convertAtomicXchgToIntegerType(AtomicRMWInst *RMWI);
Value *
insertRMWLLSCLoop(IRBuilderBase &Builder, Type *ResultTy, Value *Addr,
@@ -771,6 +772,28 @@ static void createCmpXchgInstFun(IRBuilderBase &Builder, Value *Addr,
NewLoaded = Builder.CreateBitCast(NewLoaded, OrigTy);
}
+void AtomicExpandImpl::expandAtomicSubToAdd(AtomicRMWInst *AI) {
+ ReplacementIRBuilder Builder(AI, *DL);
+ AtomicRMWInst::BinOp NewOp;
+ Value *NewVal;
+
+ switch (AI->getOperation()) {
+ case AtomicRMWInst::Sub:
+ NewOp = AtomicRMWInst::Add;
+ NewVal = Builder.CreateNeg(AI->getValOperand(), "neg");
+ break;
+ case AtomicRMWInst::FSub:
+ NewOp = AtomicRMWInst::FAdd;
+ NewVal = Builder.CreateFNeg(AI->getValOperand(), "fneg");
+ break;
+ default:
+ llvm_unreachable("unsupported atomicrmw expansion");
+ }
+
+ AI->setOperation(NewOp);
+ AI->setOperand(1, NewVal);
+}
+
bool AtomicExpandImpl::tryExpandAtomicRMW(AtomicRMWInst *AI) {
LLVMContext &Ctx = AI->getModule()->getContext();
TargetLowering::AtomicExpansionKind Kind = TLI->shouldExpandAtomicRMWInIR(AI);
@@ -839,6 +862,9 @@ bool AtomicExpandImpl::tryExpandAtomicRMW(AtomicRMWInst *AI) {
TLI->emitCmpArithAtomicRMWIntrinsic(AI);
return true;
}
+ case TargetLoweringBase::AtomicExpansionKind::Expand:
+ expandAtomicSubToAdd(AI);
+ return true;
case TargetLoweringBase::AtomicExpansionKind::NotAtomic:
return lowerAtomicRMWInst(AI);
case TargetLoweringBase::AtomicExpansionKind::CustomExpand:
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 90c11f74ec5fa..a470ed5dca2fc 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1128,8 +1128,6 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
setOperationAction(ISD::ADDRSPACECAST, {MVT::i32, MVT::i64}, Custom);
- setOperationAction(ISD::ATOMIC_LOAD_SUB, {MVT::i32, MVT::i64}, Expand);
-
// atom.b128 is legal in PTX but since we don't represent i128 as a legal
// type, we need to custom lower it.
setOperationAction({ISD::ATOMIC_CMP_SWAP, ISD::ATOMIC_SWAP}, MVT::i128,
@@ -7636,8 +7634,9 @@ NVPTXTargetLowering::AtomicExpansionKind
NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
Type *Ty = AI->getValOperand()->getType();
- // Try to lower LLVM atomicrmw fadd to PTX atomic.add. This is complicated
- // by the weird FTZ behavior PTX atom.add has:
+ // Try to lower LLVM atomicrmw fadd/fsub to PTX atomic.add. Fsub is first
+ // expanded to an fadd with a negated operand. This is complicated by the
+ // weird FTZ behavior PTX atom.add has:
// - atom.add.f32 on global memory flushes denormals
// - atom.add.f32 on shared memory does not flush denormals
// - atom.add.f16 and atomic.add.bf16 never flush denormals
@@ -7647,8 +7646,13 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
// atomic.add.bf16; even though it never flushes denormals, we never flush
// bf16 denormals when doing regular arithmetic, even when FTZ is enabled.
if (AI->isFloatingPointOperation() &&
- AI->getOperation() == AtomicRMWInst::BinOp::FAdd) {
+ (AI->getOperation() == AtomicRMWInst::BinOp::FAdd ||
+ AI->getOperation() == AtomicRMWInst::BinOp::FSub)) {
const Function *F = AI->getFunction();
+ AtomicExpansionKind ExpansionKind =
+ AI->getOperation() == AtomicRMWInst::BinOp::FSub
+ ? AtomicExpansionKind::Expand
+ : AtomicExpansionKind::None;
// AllowFTZAtomics forces atom.add regardless of the FTZ mismatch.
if (Ty->isFloatTy()) {
@@ -7665,7 +7669,7 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
break;
}
if (UseNative)
- return AtomicExpansionKind::None;
+ return ExpansionKind;
}
if (Ty->isHalfTy()) {
@@ -7675,14 +7679,14 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
DenormalMode::PreserveSign;
if ((!FTZ || AllowFTZAtomics) && STI.hasFeature(NVPTX::SM70) &&
STI.hasFeature(NVPTX::PTX63))
- return AtomicExpansionKind::None;
+ return ExpansionKind;
}
if (Ty->isBFloatTy() && STI.hasFeature(NVPTX::SM90))
- return AtomicExpansionKind::None;
+ return ExpansionKind;
if (Ty->isDoubleTy() && STI.hasAtomAddF64())
- return AtomicExpansionKind::None;
+ return ExpansionKind;
}
// PTX's only atomic fp op is `add`; all other ops expand to a CAS loop.
@@ -7725,22 +7729,27 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
case AtomicRMWInst::BinOp::Max:
case AtomicRMWInst::BinOp::Min:
case AtomicRMWInst::BinOp::UMax:
- case AtomicRMWInst::BinOp::UMin:
+ case AtomicRMWInst::BinOp::UMin: {
+ AtomicExpansionKind ExpansionKind =
+ AI->getOperation() == AtomicRMWInst::BinOp::Sub
+ ? AtomicExpansionKind::Expand
+ : AtomicExpansionKind::None;
switch (BitWidth) {
case 8:
case 16:
return AtomicExpansionKind::CmpXChg;
case 32:
- return AtomicExpansionKind::None;
+ return ExpansionKind;
case 64:
if (STI.hasAtomMinMax64())
- return AtomicExpansionKind::None;
+ return ExpansionKind;
return AtomicExpansionKind::CmpXChg;
case 128:
return AtomicExpansionKind::CmpXChg;
default:
llvm_unreachable("unsupported width encountered");
}
+ }
case AtomicRMWInst::BinOp::UIncWrap:
case AtomicRMWInst::BinOp::UDecWrap:
switch (BitWidth) {
@@ -7812,9 +7821,13 @@ AtomicOrdering NVPTXTargetLowering::atomicOperationOrderAfterFenceSplit(
STI.getMinCmpXchgSizeInBits())
return AtomicOrdering::Acquire;
else if (auto *RI = dyn_cast<AtomicRMWInst>(I);
- RI && RI->getOrdering() == AtomicOrdering::SequentiallyConsistent &&
- shouldExpandAtomicRMWInIR(RI) == AtomicExpansionKind::None)
- return AtomicOrdering::Acquire;
+ RI &&
+ RI->getOrdering() == AtomicOrdering::SequentiallyConsistent) {
+ AtomicExpansionKind ExpansionKind = shouldExpandAtomicRMWInIR(RI);
+ if (ExpansionKind == AtomicExpansionKind::None ||
+ ExpansionKind == AtomicExpansionKind::Expand)
+ return AtomicOrdering::Acquire;
+ }
return AtomicOrdering::Monotonic;
}
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
index 5a43d78501ae0..7ea96bdba553b 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
@@ -1835,71 +1835,41 @@ define float @fsub_acq_rel_float_global_cta(ptr addrspace(1) %addr, float %val)
;
; SM60-NOFTZ-ALLOW-LABEL: fsub_acq_rel_float_global_cta(
; SM60-NOFTZ-ALLOW: {
-; SM60-NOFTZ-ALLOW-NEXT: .reg .pred %p<2>;
-; SM60-NOFTZ-ALLOW-NEXT: .reg .b32 %r<5>;
+; SM60-NOFTZ-ALLOW-NEXT: .reg .b32 %r<4>;
; SM60-NOFTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
; SM60-NOFTZ-ALLOW-EMPTY:
; SM60-NOFTZ-ALLOW-NEXT: // %bb.0:
-; SM60-NOFTZ-ALLOW-NEXT: ld.param.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
; SM60-NOFTZ-ALLOW-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
-; SM60-NOFTZ-ALLOW-NEXT: membar.cta;
-; SM60-NOFTZ-ALLOW-NEXT: ld.volatile.global.b32 %r4, [%rd1];
-; SM60-NOFTZ-ALLOW-NEXT: $L__BB61_1: // %atomicrmw.start
-; SM60-NOFTZ-ALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM60-NOFTZ-ALLOW-NEXT: sub.rn.f32 %r3, %r4, %r2;
-; SM60-NOFTZ-ALLOW-NEXT: atom.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM60-NOFTZ-ALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; SM60-NOFTZ-ALLOW-NEXT: mov.b32 %r4, %r1;
-; SM60-NOFTZ-ALLOW-NEXT: @%p1 bra $L__BB61_1;
-; SM60-NOFTZ-ALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM60-NOFTZ-ALLOW-NEXT: membar.cta;
-; SM60-NOFTZ-ALLOW-NEXT: st.param.b32 [func_retval0], %r1;
+; SM60-NOFTZ-ALLOW-NEXT: ld.param.b32 %r1, [fsub_acq_rel_float_global_cta_param_1];
+; SM60-NOFTZ-ALLOW-NEXT: neg.f32 %r2, %r1;
+; SM60-NOFTZ-ALLOW-NEXT: atom.cta.global.add.f32 %r3, [%rd1], %r2;
+; SM60-NOFTZ-ALLOW-NEXT: st.param.b32 [func_retval0], %r3;
; SM60-NOFTZ-ALLOW-NEXT: ret;
;
; SM60-FTZ-DISALLOW-LABEL: fsub_acq_rel_float_global_cta(
; SM60-FTZ-DISALLOW: {
-; SM60-FTZ-DISALLOW-NEXT: .reg .pred %p<2>;
-; SM60-FTZ-DISALLOW-NEXT: .reg .b32 %r<5>;
+; SM60-FTZ-DISALLOW-NEXT: .reg .b32 %r<4>;
; SM60-FTZ-DISALLOW-NEXT: .reg .b64 %rd<2>;
; SM60-FTZ-DISALLOW-EMPTY:
; SM60-FTZ-DISALLOW-NEXT: // %bb.0:
-; SM60-FTZ-DISALLOW-NEXT: ld.param.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
; SM60-FTZ-DISALLOW-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
-; SM60-FTZ-DISALLOW-NEXT: membar.cta;
-; SM60-FTZ-DISALLOW-NEXT: ld.volatile.global.b32 %r4, [%rd1];
-; SM60-FTZ-DISALLOW-NEXT: $L__BB61_1: // %atomicrmw.start
-; SM60-FTZ-DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM60-FTZ-DISALLOW-NEXT: sub.rn.ftz.f32 %r3, %r4, %r2;
-; SM60-FTZ-DISALLOW-NEXT: atom.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM60-FTZ-DISALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; SM60-FTZ-DISALLOW-NEXT: mov.b32 %r4, %r1;
-; SM60-FTZ-DISALLOW-NEXT: @%p1 bra $L__BB61_1;
-; SM60-FTZ-DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM60-FTZ-DISALLOW-NEXT: membar.cta;
-; SM60-FTZ-DISALLOW-NEXT: st.param.b32 [func_retval0], %r1;
+; SM60-FTZ-DISALLOW-NEXT: ld.param.b32 %r1, [fsub_acq_rel_float_global_cta_param_1];
+; SM60-FTZ-DISALLOW-NEXT: neg.ftz.f32 %r2, %r1;
+; SM60-FTZ-DISALLOW-NEXT: atom.cta.global.add.f32 %r3, [%rd1], %r2;
+; SM60-FTZ-DISALLOW-NEXT: st.param.b32 [func_retval0], %r3;
; SM60-FTZ-DISALLOW-NEXT: ret;
;
; SM60-FTZ-ALLOW-LABEL: fsub_acq_rel_float_global_cta(
; SM60-FTZ-ALLOW: {
-; SM60-FTZ-ALLOW-NEXT: .reg .pred %p<2>;
-; SM60-FTZ-ALLOW-NEXT: .reg .b32 %r<5>;
+; SM60-FTZ-ALLOW-NEXT: .reg .b32 %r<4>;
; SM60-FTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
; SM60-FTZ-ALLOW-EMPTY:
; SM60-FTZ-ALLOW-NEXT: // %bb.0:
-; SM60-FTZ-ALLOW-NEXT: ld.param.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
; SM60-FTZ-ALLOW-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
-; SM60-FTZ-ALLOW-NEXT: membar.cta;
-; SM60-FTZ-ALLOW-NEXT: ld.volatile.global.b32 %r4, [%rd1];
-; SM60-FTZ-ALLOW-NEXT: $L__BB61_1: // %atomicrmw.start
-; SM60-FTZ-ALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM60-FTZ-ALLOW-NEXT: sub.rn.ftz.f32 %r3, %r4, %r2;
-; SM60-FTZ-ALLOW-NEXT: atom.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM60-FTZ-ALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; SM60-FTZ-ALLOW-NEXT: mov.b32 %r4, %r1;
-; SM60-FTZ-ALLOW-NEXT: @%p1 bra $L__BB61_1;
-; SM60-FTZ-ALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM60-FTZ-ALLOW-NEXT: membar.cta;
-; SM60-FTZ-ALLOW-NEXT: st.param.b32 [func_retval0], %r1;
+; SM60-FTZ-ALLOW-NEXT: ld.param.b32 %r1, [fsub_acq_rel_float_global_cta_param_1];
+; SM60-FTZ-ALLOW-NEXT: neg.ftz.f32 %r2, %r1;
+; SM60-FTZ-ALLOW-NEXT: atom.cta.global.add.f32 %r3, [%rd1], %r2;
+; SM60-FTZ-ALLOW-NEXT: st.param.b32 [func_retval0], %r3;
; SM60-FTZ-ALLOW-NEXT: ret;
%retval = atomicrmw fsub ptr addrspace(1) %addr, float %val syncscope("block") acq_rel
ret float %retval
@@ -2355,24 +2325,14 @@ define double @fadd_acq_rel_double_global_cta(ptr addrspace(1) %addr, double %va
define double @fsub_acq_rel_double_global_cta(ptr addrspace(1) %addr, double %val) {
; SM60-LABEL: fsub_acq_rel_double_global_cta(
; SM60: {
-; SM60-NEXT: .reg .pred %p<2>;
-; SM60-NEXT: .reg .b64 %rd<6>;
+; SM60-NEXT: .reg .b64 %rd<5>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
-; SM60-NEXT: ld.param.b64 %rd3, [fsub_acq_rel_double_global_cta_param_1];
-; SM60-NEXT: ld.param.b64 %rd2, [fsub_acq_rel_double_global_cta_param_0];
-; SM60-NEXT: membar.cta;
-; SM60-NEXT: ld.volatile.global.b64 %rd5, [%rd2];
-; SM60-NEXT: $L__BB67_1: // %atomicrmw.start
-; SM60-NEXT: // =>This Inner Loop Header: Depth=1
-; SM60-NEXT: sub.rn.f64 %rd4, %rd5, %rd3;
-; SM60-NEXT: atom.cta.global.cas.b64 %rd1, [%rd2], %rd5, %rd4;
-; SM60-NEXT: setp.ne.b64 %p1, %rd1, %rd5;
-; SM60-NEXT: mov.b64 %rd5, %rd1;
-; SM60-NEXT: @%p1 bra $L__BB67_1;
-; SM60-NEXT: // %bb.2: // %atomicrmw.end
-; SM60-NEXT: membar.cta;
-; SM60-NEXT: st.param.b64 [func_retval0], %rd1;
+; SM60-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_double_global_cta_param_0];
+; SM60-NEXT: ld.param.b64 %rd2, [fsub_acq_rel_double_global_cta_param_1];
+; SM60-NEXT: neg.f64 %rd3, %rd2;
+; SM60-NEXT: atom.cta.global.add.f64 %rd4, [%rd1], %rd3;
+; SM60-NEXT: st.param.b64 [func_retval0], %rd4;
; SM60-NEXT: ret;
%retval = atomicrmw fsub ptr addrspace(1) %addr, double %val syncscope("block") acq_rel
ret double %retval
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
index 5b9f1a68bceab..00fbbf4db9ff2 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
@@ -1835,71 +1835,41 @@ define float @fsub_acq_rel_float_global_cta(ptr addrspace(1) %addr, float %val)
;
; SM70-NOFTZ-ALLOW-LABEL: fsub_acq_rel_float_global_cta(
; SM70-NOFTZ-ALLOW: {
-; SM70-NOFTZ-ALLOW-NEXT: .reg .pred %p<2>;
-; SM70-NOFTZ-ALLOW-NEXT: .reg .b32 %r<5>;
+; SM70-NOFTZ-ALLOW-NEXT: .reg .b32 %r<4>;
; SM70-NOFTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
; SM70-NOFTZ-ALLOW-EMPTY:
; SM70-NOFTZ-ALLOW-NEXT: // %bb.0:
-; SM70-NOFTZ-ALLOW-NEXT: ld.param.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
; SM70-NOFTZ-ALLOW-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
-; SM70-NOFTZ-ALLOW-NEXT: fence.acq_rel.cta;
-; SM70-NOFTZ-ALLOW-NEXT: ld.relaxed.cta.global.b32 %r4, [%rd1];
-; SM70-NOFTZ-ALLOW-NEXT: $L__BB61_1: // %atomicrmw.start
-; SM70-NOFTZ-ALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM70-NOFTZ-ALLOW-NEXT: sub.rn.f32 %r3, %r4, %r2;
-; SM70-NOFTZ-ALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM70-NOFTZ-ALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; SM70-NOFTZ-ALLOW-NEXT: mov.b32 %r4, %r1;
-; SM70-NOFTZ-ALLOW-NEXT: @%p1 bra $L__BB61_1;
-; SM70-NOFTZ-ALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM70-NOFTZ-ALLOW-NEXT: fence.acq_rel.cta;
-; SM70-NOFTZ-ALLOW-NEXT: st.param.b32 [func_retval0], %r1;
+; SM70-NOFTZ-ALLOW-NEXT: ld.param.b32 %r1, [fsub_acq_rel_float_global_cta_param_1];
+; SM70-NOFTZ-ALLOW-NEXT: neg.f32 %r2, %r1;
+; SM70-NOFTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.f32 %r3, [%rd1], %r2;
+; SM70-NOFTZ-ALLOW-NEXT: st.param.b32 [func_retval0], %r3;
; SM70-NOFTZ-ALLOW-NEXT: ret;
;
; SM70-FTZ-DISALLOW-LABEL: fsub_acq_rel_float_global_cta(
; SM70-FTZ-DISALLOW: {
-; SM70-FTZ-DISALLOW-NEXT: .reg .pred %p<2>;
-; SM70-FTZ-DISALLOW-NEXT: .reg .b32 %r<5>;
+; SM70-FTZ-DISALLOW-NEXT: .reg .b32 %r<4>;
; SM70-FTZ-DISALLOW-NEXT: .reg .b64 %rd<2>;
; SM70-FTZ-DISALLOW-EMPTY:
; SM70-FTZ-DISALLOW-NEXT: // %bb.0:
-; SM70-FTZ-DISALLOW-NEXT: ld.param.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
; SM70-FTZ-DISALLOW-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
-; SM70-FTZ-DISALLOW-NEXT: fence.acq_rel.cta;
-; SM70-FTZ-DISALLOW-NEXT: ld.relaxed.cta.global.b32 %r4, [%rd1];
-; SM70-FTZ-DISALLOW-NEXT: $L__BB61_1: // %atomicrmw.start
-; SM70-FTZ-DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM70-FTZ-DISALLOW-NEXT: sub.rn.ftz.f32 %r3, %r4, %r2;
-; SM70-FTZ-DISALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM70-FTZ-DISALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; SM70-FTZ-DISALLOW-NEXT: mov.b32 %r4, %r1;
-; SM70-FTZ-DISALLOW-NEXT: @%p1 bra $L__BB61_1;
-; SM70-FTZ-DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM70-FTZ-DISALLOW-NEXT: fence.acq_rel.cta;
-; SM70-FTZ-DISALLOW-NEXT: st.param.b32 [func_retval0], %r1;
+; SM70-FTZ-DISALLOW-NEXT: ld.param.b32 %r1, [fsub_acq_rel_float_global_cta_param_1];
+; SM70-FTZ-DISALLOW-NEXT: neg.ftz.f32 %r2, %r1;
+; SM70-FTZ-DISALLOW-NEXT: atom.acq_rel.cta.global.add.f32 %r3, [%rd1], %r2;
+; SM70-FTZ-DISALLOW-NEXT: st.param.b32 [func_retval0], %r3;
; SM70-FTZ-DISALLOW-NEXT: ret;
;
; SM70-FTZ-ALLOW-LABEL: fsub_acq_rel_float_global_cta(
; SM70-FTZ-ALLOW: {
-; SM70-FTZ-ALLOW-NEXT: .reg .pred %p<2>;
-; SM70-FTZ-ALLOW-NEXT: .reg .b32 %r<5>;
+; SM70-FTZ-ALLOW-NEXT: .reg .b32 %r<4>;
; SM70-FTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
; SM70-FTZ-ALLOW-EMPTY:
; SM70-FTZ-ALLOW-NEXT: // %bb.0:
-; SM70-FTZ-ALLOW-NEXT: ld.param.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
; SM70-FTZ-ALLOW-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
-; SM70-FTZ-ALLOW-NEXT: fence.acq_rel.cta;
-; SM70-FTZ-ALLOW-NEXT: ld.relaxed.cta.global.b32 %r4, [%rd1];
-; SM70-FTZ-ALLOW-NEXT: $L__BB61_1: // %atomicrmw.start
-; SM70-FTZ-ALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM70-FTZ-ALLOW-NEXT: sub.rn.ftz.f32 %r3, %r4, %r2;
-; SM70-FTZ-ALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM70-FTZ-ALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; SM70-FTZ-ALLOW-NEXT: mov.b32 %r4, %r1;
-; SM70-FTZ-ALLOW-NEXT: @%p1 bra $L__BB61_1;
-; SM70-FTZ-ALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM70-FTZ-ALLOW-NEXT: fence.acq_rel.cta;
-; SM70-FTZ-ALLOW-NEXT: st.param.b32 [func_retval0], %r1;
+; SM70-FTZ-ALLOW-NEXT: ld.param.b32 %r1, [fsub_acq_rel_float_global_cta_param_1];
+; SM70-FTZ-ALLOW-NEXT: neg.ftz.f32 %r2, %r1;
+; SM70-FTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.f32 %r3, [%rd1], %r2;
+; SM70-FTZ-ALLOW-NEXT: st.param.b32 [func_retval0], %r3;
; SM70-FTZ-ALLOW-NEXT: ret;
%retval = atomicrmw fsub ptr addrspace(1) %addr, float %val syncscope("block") acq_rel
ret float %retval
@@ -2355,24 +2325,14 @@ define double @fadd_acq_rel_double_global_cta(ptr addrspace(1) %addr, double %va
define double @fsub_acq_rel_double_global_cta(ptr addrspace(1) %addr, double %val) {
; SM70-LABEL: fsub_acq_rel_double_global_cta(
; SM70: {
-; SM70-NEXT: .reg .pred %p<2>;
-; SM70-NEXT: .reg .b64 %rd<6>;
+; SM70-NEXT: .reg .b64 %rd<5>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
-; SM70-NEXT: ld.param.b64 %rd3, [fsub_acq_rel_double_global_cta_param_1];
-; SM70-NEXT: ld.param.b64 %rd2, [fsub_acq_rel_double_global_cta_param_0];
-; SM70-NEXT: fence.acq_rel.cta;
-; SM70-NEXT: ld.relaxed.cta.global.b64 %rd5, [%rd2];
-; SM70-NEXT: $L__BB67_1: // %atomicrmw.start
-; SM70-NEXT: // =>This Inner Loop Header: Depth=1
-; SM70-NEXT: sub.rn.f64 %rd4, %rd5, %rd3;
-; SM70-NEXT: atom.relaxed.cta.global.cas.b64 %rd1, [%rd2], %rd5, %rd4;
-; SM70-NEXT: setp.ne.b64 %p1, %rd1, %rd5;
-; SM70-NEXT: mov.b64 %rd5, %rd1;
-; SM70-NEXT: @%p1 bra $L__BB67_1;
-; SM70-NEXT: // %bb.2: // %atomicrmw.end
-; SM70-NEXT: fence.acq_rel.cta;
-; SM70-NEXT: st.param.b64 [func_retval0], %rd1;
+; SM70-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_double_global_cta_param_0];
+; SM70-NEXT: ld.param.b64 %rd2, [fsub_acq_rel_double_global_cta_param_1];
+; SM70-NEXT: neg.f64 %rd3, %rd2;
+; SM70-NEXT: atom.acq_rel.cta.global.add.f64 %rd4, [%rd1], %rd3;
+; SM70-NEXT: st.param.b64 [func_retval0], %rd4;
; SM70-NEXT: ret;
%retval = atomicrmw fsub ptr addrspace(1) %addr, double %val syncscope("block") acq_rel
ret double %retval
@@ -2513,154 +2473,54 @@ define half @fadd_acq_rel_half_global_cta(ptr addrspace(1) %addr, half %val) {
define half @fsub_acq_rel_half_global_cta(ptr addrspace(1) %addr, half %val) {
; SM70-NOFTZ-DISALLOW-LABEL: fsub_acq_rel_half_global_cta(
; SM70-NOFTZ-DISALLOW: {
-; SM70-NOFTZ-DISALLOW-NEXT: .reg .pred %p<2>;
; SM70-NOFTZ-DISALLOW-NEXT: .reg .b16 %rs<4>;
-; SM70-NOFTZ-DISALLOW-NEXT: .reg .b32 %r<15>;
-; SM70-NOFTZ-DISALLOW-NEXT: .reg .b64 %rd<3>;
+; SM70-NOFTZ-DISALLOW-NEXT: .reg .b64 %rd<2>;
; SM70-NOFTZ-DISALLOW-EMPTY:
; SM70-NOFTZ-DISALLOW-NEXT: // %bb.0:
+; SM70-NOFTZ-DISALLOW-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_half_global_cta_param_0];
; SM70-NOFTZ-DISALLOW-NEXT: ld.param.b16 %rs1, [fsub_acq_rel_half_global_cta_param_1];
-; SM70-NOFTZ-DISALLOW-NEXT: ld.param.b64 %rd2, [fsub_acq_rel_half_global_cta_param_0];
-; SM70-NOFTZ-DISALLOW-NEXT: fence.acq_rel.cta;
-; SM70-NOFTZ-DISALLOW-NEXT: and.b64 %rd1, %rd2, -4;
-; SM70-NOFTZ-DISALLOW-NEXT: cvt.u32.u64 %r4, %rd2;
-; SM70-NOFTZ-DISALLOW-NEXT: and.b32 %r5, %r4, 3;
-; SM70-NOFTZ-DISALLOW-NEXT: shl.b32 %r1, %r5, 3;
-; SM70-NOFTZ-DISALLOW-NEXT: mov.b32 %r6, 65535;
-; SM70-NOFTZ-DISALLOW-NEXT: shl.b32 %r7, %r6, %r1;
-; SM70-NOFTZ-DISALLOW-NEXT: not.b32 %r2, %r7;
-; SM70-NOFTZ-DISALLOW-NEXT: ld.relaxed.cta.global.b32 %r14, [%rd1];
-; SM70-NOFTZ-DISALLOW-NEXT: $L__BB73_1: // %atomicrmw.start
-; SM70-NOFTZ-DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM70-NOFTZ-DISALLOW-NEXT: shr.u32 %r8, %r14, %r1;
-; SM70-NOFTZ-DISALLOW-NEXT: cvt.u16.u32 %rs2, %r8;
-; SM70-NOFTZ-DISALLOW-NEXT: sub.rn.f16 %rs3, %rs2, %rs1;
-; SM70-NOFTZ-DISALLOW-NEXT: cvt.u32.u16 %r9, %rs3;
-; SM70-NOFTZ-DISALLOW-NEXT: shl.b32 %r10, %r9, %r1;
-; SM70-NOFTZ-DISALLOW-NEXT: and.b32 %r11, %r14, %r2;
-; SM70-NOFTZ-DISALLOW-NEXT: or.b32 %r12, %r11, %r10;
-; SM70-NOFTZ-DISALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r3, [%rd1], %r14, %r12;
-; SM70-NOFTZ-DISALLOW-NEXT: setp.ne.b32 %p1, %r3, %r14;
-; SM70-NOFTZ-DISALLOW-NEXT: mov.b32 %r14, %r3;
-; SM70-NOFTZ-DISALLOW-NEXT: @%p1 bra $L__BB73_1;
-; SM70-NOFTZ-DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM70-NOFTZ-DISALLOW-NEXT: shr.u32 %r13, %r3, %r1;
-; SM70-NOFTZ-DISALLOW-NEXT: fence.acq_rel.cta;
-; SM70-NOFTZ-DISALLOW-NEXT: st.param.b16 [func_retval0], %r13;
+; SM70-NOFTZ-DISALLOW-NEXT: neg.f16 %rs2, %rs1;
+; SM70-NOFTZ-DISALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs3, [%rd1], %rs2;
+; SM70-NOFTZ-DISALLOW-NEXT: st.param.b16 [func_retval0], %rs3;
; SM70-NOFTZ-DISALLOW-NEXT: ret;
;
; SM70-NOFTZ-ALLOW-LABEL: fsub_acq_rel_half_global_cta(
; SM70-NOFTZ-ALLOW: {
-; SM70-NOFTZ-ALLOW-NEXT: .reg .pred %p<2>;
; SM70-NOFTZ-ALLOW-NEXT: .reg .b16 %rs<4>;
-; SM70-NOFTZ-ALLOW-NEXT: .reg .b32 %r<15>;
-; SM70-NOFTZ-ALLOW-NEXT: .reg .b64 %rd<3>;
+; SM70-NOFTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
; SM70-NOFTZ-ALLOW-EMPTY:
; SM70-NOFTZ-ALLOW-NEXT: // %bb.0:
+; SM70-NOFTZ-ALLOW-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_half_global_cta_param_0];
; SM70-NOFTZ-ALLOW-NEXT: ld.param.b16 %rs1, [fsub_acq_rel_half_global_cta_param_1];
-; SM70-NOFTZ-ALLOW-NEXT: ld.param.b64 %rd2, [fsub_acq_rel_half_global_cta_param_0];
-; SM70-NOFTZ-ALLOW-NEXT: fence.acq_rel.cta;
-; SM70-NOFTZ-ALLOW-NEXT: and.b64 %rd1, %rd2, -4;
-; SM70-NOFTZ-ALLOW-NEXT: cvt.u32.u64 %r4, %rd2;
-; SM70-NOFTZ-ALLOW-NEXT: and.b32 %r5, %r4, 3;
-; SM70-NOFTZ-ALLOW-NEXT: shl.b32 %r1, %r5, 3;
-; SM70-NOFTZ-ALLOW-NEXT: mov.b32 %r6, 65535;
-; SM70-NOFTZ-ALLOW-NEXT: shl.b32 %r7, %r6, %r1;
-; SM70-NOFTZ-ALLOW-NEXT: not.b32 %r2, %r7;
-; SM70-NOFTZ-ALLOW-NEXT: ld.relaxed.cta.global.b32 %r14, [%rd1];
-; SM70-NOFTZ-ALLOW-NEXT: $L__BB73_1: // %atomicrmw.start
-; SM70-NOFTZ-ALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM70-NOFTZ-ALLOW-NEXT: shr.u32 %r8, %r14, %r1;
-; SM70-NOFTZ-ALLOW-NEXT: cvt.u16.u32 %rs2, %r8;
-; SM70-NOFTZ-ALLOW-NEXT: sub.rn.f16 %rs3, %rs2, %rs1;
-; SM70-NOFTZ-ALLOW-NEXT: cvt.u32.u16 %r9, %rs3;
-; SM70-NOFTZ-ALLOW-NEXT: shl.b32 %r10, %r9, %r1;
-; SM70-NOFTZ-ALLOW-NEXT: and.b32 %r11, %r14, %r2;
-; SM70-NOFTZ-ALLOW-NEXT: or.b32 %r12, %r11, %r10;
-; SM70-NOFTZ-ALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r3, [%rd1], %r14, %r12;
-; SM70-NOFTZ-ALLOW-NEXT: setp.ne.b32 %p1, %r3, %r14;
-; SM70-NOFTZ-ALLOW-NEXT: mov.b32 %r14, %r3;
-; SM70-NOFTZ-ALLOW-NEXT: @%p1 bra $L__BB73_1;
-; SM70-NOFTZ-ALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM70-NOFTZ-ALLOW-NEXT: shr.u32 %r13, %r3, %r1;
-; SM70-NOFTZ-ALLOW-NEXT: fence.acq_rel.cta;
-; SM70-NOFTZ-ALLOW-NEXT: st.param.b16 [func_retval0], %r13;
+; SM70-NOFTZ-ALLOW-NEXT: neg.f16 %rs2, %rs1;
+; SM70-NOFTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs3, [%rd1], %rs2;
+; SM70-NOFTZ-ALLOW-NEXT: st.param.b16 [func_retval0], %rs3;
; SM70-NOFTZ-ALLOW-NEXT: ret;
;
; SM70-FTZ-DISALLOW-LABEL: fsub_acq_rel_half_global_cta(
; SM70-FTZ-DISALLOW: {
-; SM70-FTZ-DISALLOW-NEXT: .reg .pred %p<2>;
; SM70-FTZ-DISALLOW-NEXT: .reg .b16 %rs<4>;
-; SM70-FTZ-DISALLOW-NEXT: .reg .b32 %r<15>;
-; SM70-FTZ-DISALLOW-NEXT: .reg .b64 %rd<3>;
+; SM70-FTZ-DISALLOW-NEXT: .reg .b64 %rd<2>;
; SM70-FTZ-DISALLOW-EMPTY:
; SM70-FTZ-DISALLOW-NEXT: // %bb.0:
+; SM70-FTZ-DISALLOW-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_half_global_cta_param_0];
; SM70-FTZ-DISALLOW-NEXT: ld.param.b16 %rs1, [fsub_acq_rel_half_global_cta_param_1];
-; SM70-FTZ-DISALLOW-NEXT: ld.param.b64 %rd2, [fsub_acq_rel_half_global_cta_param_0];
-; SM70-FTZ-DISALLOW-NEXT: fence.acq_rel.cta;
-; SM70-FTZ-DISALLOW-NEXT: and.b64 %rd1, %rd2, -4;
-; SM70-FTZ-DISALLOW-NEXT: cvt.u32.u64 %r4, %rd2;
-; SM70-FTZ-DISALLOW-NEXT: and.b32 %r5, %r4, 3;
-; SM70-FTZ-DISALLOW-NEXT: shl.b32 %r1, %r5, 3;
-; SM70-FTZ-DISALLOW-NEXT: mov.b32 %r6, 65535;
-; SM70-FTZ-DISALLOW-NEXT: shl.b32 %r7, %r6, %r1;
-; SM70-FTZ-DISALLOW-NEXT: not.b32 %r2, %r7;
-; SM70-FTZ-DISALLOW-NEXT: ld.relaxed.cta.global.b32 %r14, [%rd1];
-; SM70-FTZ-DISALLOW-NEXT: $L__BB73_1: // %atomicrmw.start
-; SM70-FTZ-DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM70-FTZ-DISALLOW-NEXT: shr.u32 %r8, %r14, %r1;
-; SM70-FTZ-DISALLOW-NEXT: cvt.u16.u32 %rs2, %r8;
-; SM70-FTZ-DISALLOW-NEXT: sub.rn.ftz.f16 %rs3, %rs2, %rs1;
-; SM70-FTZ-DISALLOW-NEXT: cvt.u32.u16 %r9, %rs3;
-; SM70-FTZ-DISALLOW-NEXT: shl.b32 %r10, %r9, %r1;
-; SM70-FTZ-DISALLOW-NEXT: and.b32 %r11, %r14, %r2;
-; SM70-FTZ-DISALLOW-NEXT: or.b32 %r12, %r11, %r10;
-; SM70-FTZ-DISALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r3, [%rd1], %r14, %r12;
-; SM70-FTZ-DISALLOW-NEXT: setp.ne.b32 %p1, %r3, %r14;
-; SM70-FTZ-DISALLOW-NEXT: mov.b32 %r14, %r3;
-; SM70-FTZ-DISALLOW-NEXT: @%p1 bra $L__BB73_1;
-; SM70-FTZ-DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM70-FTZ-DISALLOW-NEXT: shr.u32 %r13, %r3, %r1;
-; SM70-FTZ-DISALLOW-NEXT: fence.acq_rel.cta;
-; SM70-FTZ-DISALLOW-NEXT: st.param.b16 [func_retval0], %r13;
+; SM70-FTZ-DISALLOW-NEXT: neg.ftz.f16 %rs2, %rs1;
+; SM70-FTZ-DISALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs3, [%rd1], %rs2;
+; SM70-FTZ-DISALLOW-NEXT: st.param.b16 [func_retval0], %rs3;
; SM70-FTZ-DISALLOW-NEXT: ret;
;
; SM70-FTZ-ALLOW-LABEL: fsub_acq_rel_half_global_cta(
; SM70-FTZ-ALLOW: {
-; SM70-FTZ-ALLOW-NEXT: .reg .pred %p<2>;
; SM70-FTZ-ALLOW-NEXT: .reg .b16 %rs<4>;
-; SM70-FTZ-ALLOW-NEXT: .reg .b32 %r<15>;
-; SM70-FTZ-ALLOW-NEXT: .reg .b64 %rd<3>;
+; SM70-FTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
; SM70-FTZ-ALLOW-EMPTY:
; SM70-FTZ-ALLOW-NEXT: // %bb.0:
+; SM70-FTZ-ALLOW-NEXT: ld.param.b64 %rd1, [fsub_acq_rel_half_global_cta_param_0];
; SM70-FTZ-ALLOW-NEXT: ld.param.b16 %rs1, [fsub_acq_rel_half_global_cta_param_1];
-; SM70-FTZ-ALLOW-NEXT: ld.param.b64 %rd2, [fsub_acq_rel_half_global_cta_param_0];
-; SM70-FTZ-ALLOW-NEXT: fence.acq_rel.cta;
-; SM70-FTZ-ALLOW-NEXT: and.b64 %rd1, %rd2, -4;
-; SM70-FTZ-ALLOW-NEXT: cvt.u32.u64 %r4, %rd2;
-; SM70-FTZ-ALLOW-NEXT: and.b32 %r5, %r4, 3;
-; SM70-FTZ-ALLOW-NEXT: shl.b32 %r1, %r5, 3;
-; SM70-FTZ-ALLOW-NEXT: mov.b32 %r6, 65535;
-; SM70-FTZ-ALLOW-NEXT: shl.b32 %r7, %r6, %r1;
-; SM70-FTZ-ALLOW-NEXT: not.b32 %r2, %r7;
-; SM70-FTZ-ALLOW-NEXT: ld.relaxed.cta.global.b32 %r14, [%rd1];
-; SM70-FTZ-ALLOW-NEXT: $L__BB73_1: // %atomicrmw.start
-; SM70-FTZ-ALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM70-FTZ-ALLOW-NEXT: shr.u32 %r8, %r14, %r1;
-; SM70-FTZ-ALLOW-NEXT: cvt.u16.u32 %rs2, %r8;
-; SM70-FTZ-ALLOW-NEXT: sub.rn.ftz.f16 %rs3, %rs2, %rs1;
-; SM70-FTZ-ALLOW-NEXT: cvt.u32.u16 %r9, %rs3;
-; SM70-FTZ-ALLOW-NEXT: shl.b32 %r10, %r9, %r1;
-; SM70-FTZ-ALLOW-NEXT: and.b32 %r11, %r14, %r2;
-; SM70-FTZ-ALLOW-NEXT: or.b32 %r12, %r11, %r10;
-; SM70-FTZ-ALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r3, [%rd1], %r14, %r12;
-; SM70-FTZ-ALLOW-NEXT: setp.ne.b32 %p1, %r3, %r14;
-; SM70-FTZ-ALLOW-NEXT: mov.b32 %r14, %r3;
-; SM70-FTZ-ALLOW-NEXT: @%p1 bra $L__BB73_1;
-; SM70-FTZ-ALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM70-FTZ-ALLOW-NEXT: shr.u32 %r13, %r3, %r1;
-; SM70-FTZ-ALLOW-NEXT: fence.acq_rel.cta;
-; SM70-FTZ-ALLOW-NEXT: st.param.b16 [func_retval0], %r13;
+; SM70-FTZ-ALLOW-NEXT: neg.ftz.f16 %rs2, %rs1;
+; SM70-FTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs3, [%rd1], %rs2;
+; SM70-FTZ-ALLOW-NEXT: st.param.b16 [func_retval0], %rs3;
; SM70-FTZ-ALLOW-NEXT: ret;
%retval = atomicrmw fsub ptr addrspace(1) %addr, half %val syncscope("block") acq_rel
ret half %retval
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
index a1470453fa427..6c55435e9b0cf 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
@@ -1835,71 +1835,41 @@ define float @fsub_acq_rel_float_global_cta(ptr addrspace(1) %addr, float %val)
;
; SM90-NOFTZ-ALLOW-LABEL: fsub_acq_rel_float_global_cta(
; SM90-NOFTZ-ALLOW: {
-; SM90-NOFTZ-ALLOW-NEXT: .reg .pred %p<2>;
-; SM90-NOFTZ-ALLOW-NEXT: .reg .b32 %r<5>;
+; SM90-NOFTZ-ALLOW-NEXT: .reg .b32 %r<4>;
; SM90-NOFTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
; SM90-NOFTZ-ALLOW-EMPTY:
; SM90-NOFTZ-ALLOW-NEXT: // %bb.0:
-; SM90-NOFTZ-ALLOW-NEXT: ld.param::func.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
; SM90-NOFTZ-ALLOW-NEXT: ld.param::func.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
-; SM90-NOFTZ-ALLOW-NEXT: fence.release.cta;
-; SM90-NOFTZ-ALLOW-NEXT: ld.relaxed.cta.global.b32 %r4, [%rd1];
-; SM90-NOFTZ-ALLOW-NEXT: $L__BB61_1: // %atomicrmw.start
-; SM90-NOFTZ-ALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM90-NOFTZ-ALLOW-NEXT: sub.rn.f32 %r3, %r4, %r2;
-; SM90-NOFTZ-ALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM90-NOFTZ-ALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; SM90-NOFTZ-ALLOW-NEXT: mov.b32 %r4, %r1;
-; SM90-NOFTZ-ALLOW-NEXT: @%p1 bra $L__BB61_1;
-; SM90-NOFTZ-ALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM90-NOFTZ-ALLOW-NEXT: fence.acquire.cta;
-; SM90-NOFTZ-ALLOW-NEXT: st.param::func.b32 [func_retval0], %r1;
+; SM90-NOFTZ-ALLOW-NEXT: ld.param::func.b32 %r1, [fsub_acq_rel_float_global_cta_param_1];
+; SM90-NOFTZ-ALLOW-NEXT: neg.f32 %r2, %r1;
+; SM90-NOFTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.f32 %r3, [%rd1], %r2;
+; SM90-NOFTZ-ALLOW-NEXT: st.param::func.b32 [func_retval0], %r3;
; SM90-NOFTZ-ALLOW-NEXT: ret;
;
; SM90-FTZ-DISALLOW-LABEL: fsub_acq_rel_float_global_cta(
; SM90-FTZ-DISALLOW: {
-; SM90-FTZ-DISALLOW-NEXT: .reg .pred %p<2>;
-; SM90-FTZ-DISALLOW-NEXT: .reg .b32 %r<5>;
+; SM90-FTZ-DISALLOW-NEXT: .reg .b32 %r<4>;
; SM90-FTZ-DISALLOW-NEXT: .reg .b64 %rd<2>;
; SM90-FTZ-DISALLOW-EMPTY:
; SM90-FTZ-DISALLOW-NEXT: // %bb.0:
-; SM90-FTZ-DISALLOW-NEXT: ld.param::func.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
; SM90-FTZ-DISALLOW-NEXT: ld.param::func.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
-; SM90-FTZ-DISALLOW-NEXT: fence.release.cta;
-; SM90-FTZ-DISALLOW-NEXT: ld.relaxed.cta.global.b32 %r4, [%rd1];
-; SM90-FTZ-DISALLOW-NEXT: $L__BB61_1: // %atomicrmw.start
-; SM90-FTZ-DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM90-FTZ-DISALLOW-NEXT: sub.rn.ftz.f32 %r3, %r4, %r2;
-; SM90-FTZ-DISALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM90-FTZ-DISALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; SM90-FTZ-DISALLOW-NEXT: mov.b32 %r4, %r1;
-; SM90-FTZ-DISALLOW-NEXT: @%p1 bra $L__BB61_1;
-; SM90-FTZ-DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM90-FTZ-DISALLOW-NEXT: fence.acquire.cta;
-; SM90-FTZ-DISALLOW-NEXT: st.param::func.b32 [func_retval0], %r1;
+; SM90-FTZ-DISALLOW-NEXT: ld.param::func.b32 %r1, [fsub_acq_rel_float_global_cta_param_1];
+; SM90-FTZ-DISALLOW-NEXT: neg.ftz.f32 %r2, %r1;
+; SM90-FTZ-DISALLOW-NEXT: atom.acq_rel.cta.global.add.f32 %r3, [%rd1], %r2;
+; SM90-FTZ-DISALLOW-NEXT: st.param::func.b32 [func_retval0], %r3;
; SM90-FTZ-DISALLOW-NEXT: ret;
;
; SM90-FTZ-ALLOW-LABEL: fsub_acq_rel_float_global_cta(
; SM90-FTZ-ALLOW: {
-; SM90-FTZ-ALLOW-NEXT: .reg .pred %p<2>;
-; SM90-FTZ-ALLOW-NEXT: .reg .b32 %r<5>;
+; SM90-FTZ-ALLOW-NEXT: .reg .b32 %r<4>;
; SM90-FTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
; SM90-FTZ-ALLOW-EMPTY:
; SM90-FTZ-ALLOW-NEXT: // %bb.0:
-; SM90-FTZ-ALLOW-NEXT: ld.param::func.b32 %r2, [fsub_acq_rel_float_global_cta_param_1];
; SM90-FTZ-ALLOW-NEXT: ld.param::func.b64 %rd1, [fsub_acq_rel_float_global_cta_param_0];
-; SM90-FTZ-ALLOW-NEXT: fence.release.cta;
-; SM90-FTZ-ALLOW-NEXT: ld.relaxed.cta.global.b32 %r4, [%rd1];
-; SM90-FTZ-ALLOW-NEXT: $L__BB61_1: // %atomicrmw.start
-; SM90-FTZ-ALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM90-FTZ-ALLOW-NEXT: sub.rn.ftz.f32 %r3, %r4, %r2;
-; SM90-FTZ-ALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; SM90-FTZ-ALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; SM90-FTZ-ALLOW-NEXT: mov.b32 %r4, %r1;
-; SM90-FTZ-ALLOW-NEXT: @%p1 bra $L__BB61_1;
-; SM90-FTZ-ALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM90-FTZ-ALLOW-NEXT: fence.acquire.cta;
-; SM90-FTZ-ALLOW-NEXT: st.param::func.b32 [func_retval0], %r1;
+; SM90-FTZ-ALLOW-NEXT: ld.param::func.b32 %r1, [fsub_acq_rel_float_global_cta_param_1];
+; SM90-FTZ-ALLOW-NEXT: neg.ftz.f32 %r2, %r1;
+; SM90-FTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.f32 %r3, [%rd1], %r2;
+; SM90-FTZ-ALLOW-NEXT: st.param::func.b32 [func_retval0], %r3;
; SM90-FTZ-ALLOW-NEXT: ret;
%retval = atomicrmw fsub ptr addrspace(1) %addr, float %val syncscope("block") acq_rel
ret float %retval
@@ -2307,24 +2277,14 @@ define double @fadd_acq_rel_double_global_cta(ptr addrspace(1) %addr, double %va
define double @fsub_acq_rel_double_global_cta(ptr addrspace(1) %addr, double %val) {
; SM90-LABEL: fsub_acq_rel_double_global_cta(
; SM90: {
-; SM90-NEXT: .reg .pred %p<2>;
-; SM90-NEXT: .reg .b64 %rd<6>;
+; SM90-NEXT: .reg .b64 %rd<5>;
; SM90-EMPTY:
; SM90-NEXT: // %bb.0:
-; SM90-NEXT: ld.param::func.b64 %rd3, [fsub_acq_rel_double_global_cta_param_1];
-; SM90-NEXT: ld.param::func.b64 %rd2, [fsub_acq_rel_double_global_cta_param_0];
-; SM90-NEXT: fence.release.cta;
-; SM90-NEXT: ld.relaxed.cta.global.b64 %rd5, [%rd2];
-; SM90-NEXT: $L__BB67_1: // %atomicrmw.start
-; SM90-NEXT: // =>This Inner Loop Header: Depth=1
-; SM90-NEXT: sub.rn.f64 %rd4, %rd5, %rd3;
-; SM90-NEXT: atom.relaxed.cta.global.cas.b64 %rd1, [%rd2], %rd5, %rd4;
-; SM90-NEXT: setp.ne.b64 %p1, %rd1, %rd5;
-; SM90-NEXT: mov.b64 %rd5, %rd1;
-; SM90-NEXT: @%p1 bra $L__BB67_1;
-; SM90-NEXT: // %bb.2: // %atomicrmw.end
-; SM90-NEXT: fence.acquire.cta;
-; SM90-NEXT: st.param::func.b64 [func_retval0], %rd1;
+; SM90-NEXT: ld.param::func.b64 %rd1, [fsub_acq_rel_double_global_cta_param_0];
+; SM90-NEXT: ld.param::func.b64 %rd2, [fsub_acq_rel_double_global_cta_param_1];
+; SM90-NEXT: neg.f64 %rd3, %rd2;
+; SM90-NEXT: atom.acq_rel.cta.global.add.f64 %rd4, [%rd1], %rd3;
+; SM90-NEXT: st.param::func.b64 [func_retval0], %rd4;
; SM90-NEXT: ret;
%retval = atomicrmw fsub ptr addrspace(1) %addr, double %val syncscope("block") acq_rel
ret double %retval
@@ -2465,154 +2425,54 @@ define half @fadd_acq_rel_half_global_cta(ptr addrspace(1) %addr, half %val) {
define half @fsub_acq_rel_half_global_cta(ptr addrspace(1) %addr, half %val) {
; SM90-NOFTZ-DISALLOW-LABEL: fsub_acq_rel_half_global_cta(
; SM90-NOFTZ-DISALLOW: {
-; SM90-NOFTZ-DISALLOW-NEXT: .reg .pred %p<2>;
; SM90-NOFTZ-DISALLOW-NEXT: .reg .b16 %rs<4>;
-; SM90-NOFTZ-DISALLOW-NEXT: .reg .b32 %r<15>;
-; SM90-NOFTZ-DISALLOW-NEXT: .reg .b64 %rd<3>;
+; SM90-NOFTZ-DISALLOW-NEXT: .reg .b64 %rd<2>;
; SM90-NOFTZ-DISALLOW-EMPTY:
; SM90-NOFTZ-DISALLOW-NEXT: // %bb.0:
+; SM90-NOFTZ-DISALLOW-NEXT: ld.param::func.b64 %rd1, [fsub_acq_rel_half_global_cta_param_0];
; SM90-NOFTZ-DISALLOW-NEXT: ld.param::func.b16 %rs1, [fsub_acq_rel_half_global_cta_param_1];
-; SM90-NOFTZ-DISALLOW-NEXT: ld.param::func.b64 %rd2, [fsub_acq_rel_half_global_cta_param_0];
-; SM90-NOFTZ-DISALLOW-NEXT: fence.release.cta;
-; SM90-NOFTZ-DISALLOW-NEXT: and.b64 %rd1, %rd2, -4;
-; SM90-NOFTZ-DISALLOW-NEXT: cvt.u32.u64 %r4, %rd2;
-; SM90-NOFTZ-DISALLOW-NEXT: and.b32 %r5, %r4, 3;
-; SM90-NOFTZ-DISALLOW-NEXT: shl.b32 %r1, %r5, 3;
-; SM90-NOFTZ-DISALLOW-NEXT: mov.b32 %r6, 65535;
-; SM90-NOFTZ-DISALLOW-NEXT: shl.b32 %r7, %r6, %r1;
-; SM90-NOFTZ-DISALLOW-NEXT: not.b32 %r2, %r7;
-; SM90-NOFTZ-DISALLOW-NEXT: ld.relaxed.cta.global.b32 %r14, [%rd1];
-; SM90-NOFTZ-DISALLOW-NEXT: $L__BB73_1: // %atomicrmw.start
-; SM90-NOFTZ-DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM90-NOFTZ-DISALLOW-NEXT: shr.u32 %r8, %r14, %r1;
-; SM90-NOFTZ-DISALLOW-NEXT: cvt.u16.u32 %rs2, %r8;
-; SM90-NOFTZ-DISALLOW-NEXT: sub.rn.f16 %rs3, %rs2, %rs1;
-; SM90-NOFTZ-DISALLOW-NEXT: cvt.u32.u16 %r9, %rs3;
-; SM90-NOFTZ-DISALLOW-NEXT: shl.b32 %r10, %r9, %r1;
-; SM90-NOFTZ-DISALLOW-NEXT: and.b32 %r11, %r14, %r2;
-; SM90-NOFTZ-DISALLOW-NEXT: or.b32 %r12, %r11, %r10;
-; SM90-NOFTZ-DISALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r3, [%rd1], %r14, %r12;
-; SM90-NOFTZ-DISALLOW-NEXT: setp.ne.b32 %p1, %r3, %r14;
-; SM90-NOFTZ-DISALLOW-NEXT: mov.b32 %r14, %r3;
-; SM90-NOFTZ-DISALLOW-NEXT: @%p1 bra $L__BB73_1;
-; SM90-NOFTZ-DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM90-NOFTZ-DISALLOW-NEXT: shr.u32 %r13, %r3, %r1;
-; SM90-NOFTZ-DISALLOW-NEXT: fence.acquire.cta;
-; SM90-NOFTZ-DISALLOW-NEXT: st.param::func.b16 [func_retval0], %r13;
+; SM90-NOFTZ-DISALLOW-NEXT: neg.f16 %rs2, %rs1;
+; SM90-NOFTZ-DISALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs3, [%rd1], %rs2;
+; SM90-NOFTZ-DISALLOW-NEXT: st.param::func.b16 [func_retval0], %rs3;
; SM90-NOFTZ-DISALLOW-NEXT: ret;
;
; SM90-NOFTZ-ALLOW-LABEL: fsub_acq_rel_half_global_cta(
; SM90-NOFTZ-ALLOW: {
-; SM90-NOFTZ-ALLOW-NEXT: .reg .pred %p<2>;
; SM90-NOFTZ-ALLOW-NEXT: .reg .b16 %rs<4>;
-; SM90-NOFTZ-ALLOW-NEXT: .reg .b32 %r<15>;
-; SM90-NOFTZ-ALLOW-NEXT: .reg .b64 %rd<3>;
+; SM90-NOFTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
; SM90-NOFTZ-ALLOW-EMPTY:
; SM90-NOFTZ-ALLOW-NEXT: // %bb.0:
+; SM90-NOFTZ-ALLOW-NEXT: ld.param::func.b64 %rd1, [fsub_acq_rel_half_global_cta_param_0];
; SM90-NOFTZ-ALLOW-NEXT: ld.param::func.b16 %rs1, [fsub_acq_rel_half_global_cta_param_1];
-; SM90-NOFTZ-ALLOW-NEXT: ld.param::func.b64 %rd2, [fsub_acq_rel_half_global_cta_param_0];
-; SM90-NOFTZ-ALLOW-NEXT: fence.release.cta;
-; SM90-NOFTZ-ALLOW-NEXT: and.b64 %rd1, %rd2, -4;
-; SM90-NOFTZ-ALLOW-NEXT: cvt.u32.u64 %r4, %rd2;
-; SM90-NOFTZ-ALLOW-NEXT: and.b32 %r5, %r4, 3;
-; SM90-NOFTZ-ALLOW-NEXT: shl.b32 %r1, %r5, 3;
-; SM90-NOFTZ-ALLOW-NEXT: mov.b32 %r6, 65535;
-; SM90-NOFTZ-ALLOW-NEXT: shl.b32 %r7, %r6, %r1;
-; SM90-NOFTZ-ALLOW-NEXT: not.b32 %r2, %r7;
-; SM90-NOFTZ-ALLOW-NEXT: ld.relaxed.cta.global.b32 %r14, [%rd1];
-; SM90-NOFTZ-ALLOW-NEXT: $L__BB73_1: // %atomicrmw.start
-; SM90-NOFTZ-ALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM90-NOFTZ-ALLOW-NEXT: shr.u32 %r8, %r14, %r1;
-; SM90-NOFTZ-ALLOW-NEXT: cvt.u16.u32 %rs2, %r8;
-; SM90-NOFTZ-ALLOW-NEXT: sub.rn.f16 %rs3, %rs2, %rs1;
-; SM90-NOFTZ-ALLOW-NEXT: cvt.u32.u16 %r9, %rs3;
-; SM90-NOFTZ-ALLOW-NEXT: shl.b32 %r10, %r9, %r1;
-; SM90-NOFTZ-ALLOW-NEXT: and.b32 %r11, %r14, %r2;
-; SM90-NOFTZ-ALLOW-NEXT: or.b32 %r12, %r11, %r10;
-; SM90-NOFTZ-ALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r3, [%rd1], %r14, %r12;
-; SM90-NOFTZ-ALLOW-NEXT: setp.ne.b32 %p1, %r3, %r14;
-; SM90-NOFTZ-ALLOW-NEXT: mov.b32 %r14, %r3;
-; SM90-NOFTZ-ALLOW-NEXT: @%p1 bra $L__BB73_1;
-; SM90-NOFTZ-ALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM90-NOFTZ-ALLOW-NEXT: shr.u32 %r13, %r3, %r1;
-; SM90-NOFTZ-ALLOW-NEXT: fence.acquire.cta;
-; SM90-NOFTZ-ALLOW-NEXT: st.param::func.b16 [func_retval0], %r13;
+; SM90-NOFTZ-ALLOW-NEXT: neg.f16 %rs2, %rs1;
+; SM90-NOFTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs3, [%rd1], %rs2;
+; SM90-NOFTZ-ALLOW-NEXT: st.param::func.b16 [func_retval0], %rs3;
; SM90-NOFTZ-ALLOW-NEXT: ret;
;
; SM90-FTZ-DISALLOW-LABEL: fsub_acq_rel_half_global_cta(
; SM90-FTZ-DISALLOW: {
-; SM90-FTZ-DISALLOW-NEXT: .reg .pred %p<2>;
; SM90-FTZ-DISALLOW-NEXT: .reg .b16 %rs<4>;
-; SM90-FTZ-DISALLOW-NEXT: .reg .b32 %r<15>;
-; SM90-FTZ-DISALLOW-NEXT: .reg .b64 %rd<3>;
+; SM90-FTZ-DISALLOW-NEXT: .reg .b64 %rd<2>;
; SM90-FTZ-DISALLOW-EMPTY:
; SM90-FTZ-DISALLOW-NEXT: // %bb.0:
+; SM90-FTZ-DISALLOW-NEXT: ld.param::func.b64 %rd1, [fsub_acq_rel_half_global_cta_param_0];
; SM90-FTZ-DISALLOW-NEXT: ld.param::func.b16 %rs1, [fsub_acq_rel_half_global_cta_param_1];
-; SM90-FTZ-DISALLOW-NEXT: ld.param::func.b64 %rd2, [fsub_acq_rel_half_global_cta_param_0];
-; SM90-FTZ-DISALLOW-NEXT: fence.release.cta;
-; SM90-FTZ-DISALLOW-NEXT: and.b64 %rd1, %rd2, -4;
-; SM90-FTZ-DISALLOW-NEXT: cvt.u32.u64 %r4, %rd2;
-; SM90-FTZ-DISALLOW-NEXT: and.b32 %r5, %r4, 3;
-; SM90-FTZ-DISALLOW-NEXT: shl.b32 %r1, %r5, 3;
-; SM90-FTZ-DISALLOW-NEXT: mov.b32 %r6, 65535;
-; SM90-FTZ-DISALLOW-NEXT: shl.b32 %r7, %r6, %r1;
-; SM90-FTZ-DISALLOW-NEXT: not.b32 %r2, %r7;
-; SM90-FTZ-DISALLOW-NEXT: ld.relaxed.cta.global.b32 %r14, [%rd1];
-; SM90-FTZ-DISALLOW-NEXT: $L__BB73_1: // %atomicrmw.start
-; SM90-FTZ-DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM90-FTZ-DISALLOW-NEXT: shr.u32 %r8, %r14, %r1;
-; SM90-FTZ-DISALLOW-NEXT: cvt.u16.u32 %rs2, %r8;
-; SM90-FTZ-DISALLOW-NEXT: sub.rn.ftz.f16 %rs3, %rs2, %rs1;
-; SM90-FTZ-DISALLOW-NEXT: cvt.u32.u16 %r9, %rs3;
-; SM90-FTZ-DISALLOW-NEXT: shl.b32 %r10, %r9, %r1;
-; SM90-FTZ-DISALLOW-NEXT: and.b32 %r11, %r14, %r2;
-; SM90-FTZ-DISALLOW-NEXT: or.b32 %r12, %r11, %r10;
-; SM90-FTZ-DISALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r3, [%rd1], %r14, %r12;
-; SM90-FTZ-DISALLOW-NEXT: setp.ne.b32 %p1, %r3, %r14;
-; SM90-FTZ-DISALLOW-NEXT: mov.b32 %r14, %r3;
-; SM90-FTZ-DISALLOW-NEXT: @%p1 bra $L__BB73_1;
-; SM90-FTZ-DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM90-FTZ-DISALLOW-NEXT: shr.u32 %r13, %r3, %r1;
-; SM90-FTZ-DISALLOW-NEXT: fence.acquire.cta;
-; SM90-FTZ-DISALLOW-NEXT: st.param::func.b16 [func_retval0], %r13;
+; SM90-FTZ-DISALLOW-NEXT: neg.ftz.f16 %rs2, %rs1;
+; SM90-FTZ-DISALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs3, [%rd1], %rs2;
+; SM90-FTZ-DISALLOW-NEXT: st.param::func.b16 [func_retval0], %rs3;
; SM90-FTZ-DISALLOW-NEXT: ret;
;
; SM90-FTZ-ALLOW-LABEL: fsub_acq_rel_half_global_cta(
; SM90-FTZ-ALLOW: {
-; SM90-FTZ-ALLOW-NEXT: .reg .pred %p<2>;
; SM90-FTZ-ALLOW-NEXT: .reg .b16 %rs<4>;
-; SM90-FTZ-ALLOW-NEXT: .reg .b32 %r<15>;
-; SM90-FTZ-ALLOW-NEXT: .reg .b64 %rd<3>;
+; SM90-FTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
; SM90-FTZ-ALLOW-EMPTY:
; SM90-FTZ-ALLOW-NEXT: // %bb.0:
+; SM90-FTZ-ALLOW-NEXT: ld.param::func.b64 %rd1, [fsub_acq_rel_half_global_cta_param_0];
; SM90-FTZ-ALLOW-NEXT: ld.param::func.b16 %rs1, [fsub_acq_rel_half_global_cta_param_1];
-; SM90-FTZ-ALLOW-NEXT: ld.param::func.b64 %rd2, [fsub_acq_rel_half_global_cta_param_0];
-; SM90-FTZ-ALLOW-NEXT: fence.release.cta;
-; SM90-FTZ-ALLOW-NEXT: and.b64 %rd1, %rd2, -4;
-; SM90-FTZ-ALLOW-NEXT: cvt.u32.u64 %r4, %rd2;
-; SM90-FTZ-ALLOW-NEXT: and.b32 %r5, %r4, 3;
-; SM90-FTZ-ALLOW-NEXT: shl.b32 %r1, %r5, 3;
-; SM90-FTZ-ALLOW-NEXT: mov.b32 %r6, 65535;
-; SM90-FTZ-ALLOW-NEXT: shl.b32 %r7, %r6, %r1;
-; SM90-FTZ-ALLOW-NEXT: not.b32 %r2, %r7;
-; SM90-FTZ-ALLOW-NEXT: ld.relaxed.cta.global.b32 %r14, [%rd1];
-; SM90-FTZ-ALLOW-NEXT: $L__BB73_1: // %atomicrmw.start
-; SM90-FTZ-ALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM90-FTZ-ALLOW-NEXT: shr.u32 %r8, %r14, %r1;
-; SM90-FTZ-ALLOW-NEXT: cvt.u16.u32 %rs2, %r8;
-; SM90-FTZ-ALLOW-NEXT: sub.rn.ftz.f16 %rs3, %rs2, %rs1;
-; SM90-FTZ-ALLOW-NEXT: cvt.u32.u16 %r9, %rs3;
-; SM90-FTZ-ALLOW-NEXT: shl.b32 %r10, %r9, %r1;
-; SM90-FTZ-ALLOW-NEXT: and.b32 %r11, %r14, %r2;
-; SM90-FTZ-ALLOW-NEXT: or.b32 %r12, %r11, %r10;
-; SM90-FTZ-ALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r3, [%rd1], %r14, %r12;
-; SM90-FTZ-ALLOW-NEXT: setp.ne.b32 %p1, %r3, %r14;
-; SM90-FTZ-ALLOW-NEXT: mov.b32 %r14, %r3;
-; SM90-FTZ-ALLOW-NEXT: @%p1 bra $L__BB73_1;
-; SM90-FTZ-ALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM90-FTZ-ALLOW-NEXT: shr.u32 %r13, %r3, %r1;
-; SM90-FTZ-ALLOW-NEXT: fence.acquire.cta;
-; SM90-FTZ-ALLOW-NEXT: st.param::func.b16 [func_retval0], %r13;
+; SM90-FTZ-ALLOW-NEXT: neg.ftz.f16 %rs2, %rs1;
+; SM90-FTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs3, [%rd1], %rs2;
+; SM90-FTZ-ALLOW-NEXT: st.param::func.b16 [func_retval0], %rs3;
; SM90-FTZ-ALLOW-NEXT: ret;
%retval = atomicrmw fsub ptr addrspace(1) %addr, half %val syncscope("block") acq_rel
ret half %retval
@@ -3261,40 +3121,15 @@ define bfloat @fadd_acq_rel_bfloat_global_cta(ptr addrspace(1) %addr, bfloat %va
define bfloat @fsub_acq_rel_bfloat_global_cta(ptr addrspace(1) %addr, bfloat %val) {
; SM90-LABEL: fsub_acq_rel_bfloat_global_cta(
; SM90: {
-; SM90-NEXT: .reg .pred %p<2>;
; SM90-NEXT: .reg .b16 %rs<4>;
-; SM90-NEXT: .reg .b32 %r<15>;
-; SM90-NEXT: .reg .b64 %rd<3>;
+; SM90-NEXT: .reg .b64 %rd<2>;
; SM90-EMPTY:
; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param::func.b64 %rd1, [fsub_acq_rel_bfloat_global_cta_param_0];
; SM90-NEXT: ld.param::func.b16 %rs1, [fsub_acq_rel_bfloat_global_cta_param_1];
-; SM90-NEXT: ld.param::func.b64 %rd2, [fsub_acq_rel_bfloat_global_cta_param_0];
-; SM90-NEXT: fence.release.cta;
-; SM90-NEXT: and.b64 %rd1, %rd2, -4;
-; SM90-NEXT: cvt.u32.u64 %r4, %rd2;
-; SM90-NEXT: and.b32 %r5, %r4, 3;
-; SM90-NEXT: shl.b32 %r1, %r5, 3;
-; SM90-NEXT: mov.b32 %r6, 65535;
-; SM90-NEXT: shl.b32 %r7, %r6, %r1;
-; SM90-NEXT: not.b32 %r2, %r7;
-; SM90-NEXT: ld.relaxed.cta.global.b32 %r14, [%rd1];
-; SM90-NEXT: $L__BB79_1: // %atomicrmw.start
-; SM90-NEXT: // =>This Inner Loop Header: Depth=1
-; SM90-NEXT: shr.u32 %r8, %r14, %r1;
-; SM90-NEXT: cvt.u16.u32 %rs2, %r8;
-; SM90-NEXT: sub.rn.bf16 %rs3, %rs2, %rs1;
-; SM90-NEXT: cvt.u32.u16 %r9, %rs3;
-; SM90-NEXT: shl.b32 %r10, %r9, %r1;
-; SM90-NEXT: and.b32 %r11, %r14, %r2;
-; SM90-NEXT: or.b32 %r12, %r11, %r10;
-; SM90-NEXT: atom.relaxed.cta.global.cas.b32 %r3, [%rd1], %r14, %r12;
-; SM90-NEXT: setp.ne.b32 %p1, %r3, %r14;
-; SM90-NEXT: mov.b32 %r14, %r3;
-; SM90-NEXT: @%p1 bra $L__BB79_1;
-; SM90-NEXT: // %bb.2: // %atomicrmw.end
-; SM90-NEXT: shr.u32 %r13, %r3, %r1;
-; SM90-NEXT: fence.acquire.cta;
-; SM90-NEXT: st.param::func.b16 [func_retval0], %r13;
+; SM90-NEXT: neg.bf16 %rs2, %rs1;
+; SM90-NEXT: atom.acq_rel.cta.global.add.noftz.bf16 %rs3, [%rd1], %rs2;
+; SM90-NEXT: st.param::func.b16 [func_retval0], %rs3;
; SM90-NEXT: ret;
%retval = atomicrmw fsub ptr addrspace(1) %addr, bfloat %val syncscope("block") acq_rel
ret bfloat %retval
diff --git a/llvm/test/Transforms/AtomicExpand/NVPTX/atomicrmw-sub.ll b/llvm/test/Transforms/AtomicExpand/NVPTX/atomicrmw-sub.ll
new file mode 100644
index 0000000000000..b23845d6ab1a1
--- /dev/null
+++ b/llvm/test/Transforms/AtomicExpand/NVPTX/atomicrmw-sub.ll
@@ -0,0 +1,96 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
+; RUN: opt -S %s -passes='require<libcall-lowering-info>,atomic-expand' -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 -mattr=+ptx83 -nvptx-allow-ftz-atomics=true | FileCheck %s
+
+target triple = "nvptx64-nvidia-cuda"
+
+define i32 @sub_i32(ptr addrspace(1) %addr, i32 %val) {
+; CHECK-LABEL: @sub_i32(
+; CHECK-NEXT: [[NEG:%.*]] = sub i32 0, [[VAL:%.*]]
+; CHECK-NEXT: [[OLD:%.*]] = atomicrmw add ptr addrspace(1) [[ADDR:%.*]], i32 [[NEG]] acq_rel, align 4
+; CHECK-NEXT: ret i32 [[OLD]]
+;
+ %old = atomicrmw sub ptr addrspace(1) %addr, i32 %val acq_rel, align 4
+ ret i32 %old
+}
+
+define i32 @sub_i32_constant(ptr addrspace(1) %addr) {
+; CHECK-LABEL: @sub_i32_constant(
+; CHECK-NEXT: [[OLD:%.*]] = atomicrmw add ptr addrspace(1) [[ADDR:%.*]], i32 -7 acq_rel, align 4
+; CHECK-NEXT: ret i32 [[OLD]]
+;
+ %old = atomicrmw sub ptr addrspace(1) %addr, i32 7 acq_rel, align 4
+ ret i32 %old
+}
+
+define i64 @sub_i64_seq_cst(ptr addrspace(1) %addr, i64 %val) {
+; CHECK-LABEL: @sub_i64_seq_cst(
+; CHECK-NEXT: fence seq_cst
+; CHECK-NEXT: [[NEG:%.*]] = sub i64 0, [[VAL:%.*]]
+; CHECK-NEXT: [[OLD:%.*]] = atomicrmw add ptr addrspace(1) [[ADDR:%.*]], i64 [[NEG]] acquire, align 8
+; CHECK-NEXT: ret i64 [[OLD]]
+;
+ %old = atomicrmw sub ptr addrspace(1) %addr, i64 %val seq_cst, align 8
+ ret i64 %old
+}
+
+define float @fsub_float(ptr addrspace(1) %addr, float %val) {
+; CHECK-LABEL: @fsub_float(
+; CHECK-NEXT: [[FNEG:%.*]] = fneg float [[VAL:%.*]]
+; CHECK-NEXT: [[OLD:%.*]] = atomicrmw fadd ptr addrspace(1) [[ADDR:%.*]], float [[FNEG]] acq_rel, align 4
+; CHECK-NEXT: ret float [[OLD]]
+;
+ %old = atomicrmw fsub ptr addrspace(1) %addr, float %val acq_rel, align 4
+ ret float %old
+}
+
+define float @fsub_float_constant(ptr addrspace(1) %addr) {
+; CHECK-LABEL: @fsub_float_constant(
+; CHECK-NEXT: [[OLD:%.*]] = atomicrmw fadd ptr addrspace(1) [[ADDR:%.*]], float -1.000000e+00 acq_rel, align 4
+; CHECK-NEXT: ret float [[OLD]]
+;
+ %old = atomicrmw fsub ptr addrspace(1) %addr, float 1.000000e+00 acq_rel, align 4
+ ret float %old
+}
+
+define double @fsub_double_seq_cst(ptr addrspace(1) %addr, double %val) {
+; CHECK-LABEL: @fsub_double_seq_cst(
+; CHECK-NEXT: fence seq_cst
+; CHECK-NEXT: [[FNEG:%.*]] = fneg double [[VAL:%.*]]
+; CHECK-NEXT: [[OLD:%.*]] = atomicrmw fadd ptr addrspace(1) [[ADDR:%.*]], double [[FNEG]] acquire, align 8
+; CHECK-NEXT: ret double [[OLD]]
+;
+ %old = atomicrmw fsub ptr addrspace(1) %addr, double %val seq_cst, align 8
+ ret double %old
+}
+
+define i16 @sub_i16(ptr addrspace(1) %addr, i16 %val) {
+; CHECK-LABEL: @sub_i16(
+; CHECK-NEXT: [[ALIGNEDADDR:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[ADDR:%.*]], i64 -4)
+; CHECK-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[ADDR]] to i64
+; CHECK-NEXT: [[PTRLSB:%.*]] = and i64 [[TMP1]], 3
+; CHECK-NEXT: [[TMP2:%.*]] = shl i64 [[PTRLSB]], 3
+; CHECK-NEXT: [[SHIFTAMT:%.*]] = trunc i64 [[TMP2]] to i32
+; CHECK-NEXT: [[MASK:%.*]] = shl i32 65535, [[SHIFTAMT]]
+; CHECK-NEXT: [[INV_MASK:%.*]] = xor i32 [[MASK]], -1
+; CHECK-NEXT: [[TMP3:%.*]] = zext i16 [[VAL:%.*]] to i32
+; CHECK-NEXT: [[VALOPERAND_SHIFTED:%.*]] = shl i32 [[TMP3]], [[SHIFTAMT]]
+; CHECK-NEXT: [[TMP4:%.*]] = load atomic i32, ptr addrspace(1) [[ALIGNEDADDR]] monotonic, align 4
+; CHECK-NEXT: br label [[ATOMICRMW_START:%.*]]
+; CHECK: atomicrmw.start:
+; CHECK-NEXT: [[LOADED:%.*]] = phi i32 [ [[TMP4]], [[TMP0:%.*]] ], [ [[NEWLOADED:%.*]], [[ATOMICRMW_START]] ]
+; CHECK-NEXT: [[NEW:%.*]] = sub i32 [[LOADED]], [[VALOPERAND_SHIFTED]]
+; CHECK-NEXT: [[TMP5:%.*]] = and i32 [[NEW]], [[MASK]]
+; CHECK-NEXT: [[TMP6:%.*]] = and i32 [[LOADED]], [[INV_MASK]]
+; CHECK-NEXT: [[TMP7:%.*]] = or i32 [[TMP6]], [[TMP5]]
+; CHECK-NEXT: [[TMP8:%.*]] = cmpxchg ptr addrspace(1) [[ALIGNEDADDR]], i32 [[LOADED]], i32 [[TMP7]] monotonic monotonic, align 4
+; CHECK-NEXT: [[SUCCESS:%.*]] = extractvalue { i32, i1 } [[TMP8]], 1
+; CHECK-NEXT: [[NEWLOADED]] = extractvalue { i32, i1 } [[TMP8]], 0
+; CHECK-NEXT: br i1 [[SUCCESS]], label [[ATOMICRMW_END:%.*]], label [[ATOMICRMW_START]]
+; CHECK: atomicrmw.end:
+; CHECK-NEXT: [[SHIFTED:%.*]] = lshr i32 [[NEWLOADED]], [[SHIFTAMT]]
+; CHECK-NEXT: [[EXTRACTED:%.*]] = trunc i32 [[SHIFTED]] to i16
+; CHECK-NEXT: ret i16 [[EXTRACTED]]
+;
+ %old = atomicrmw sub ptr addrspace(1) %addr, i16 %val monotonic, align 2
+ ret i16 %old
+}
>From 6ee404420322fa121160c711f294d7197207f952 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Sat, 5 Sep 2026 09:14:57 +0000
Subject: [PATCH 2/2] format
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index a470ed5dca2fc..7479a30885ae6 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -7821,8 +7821,7 @@ AtomicOrdering NVPTXTargetLowering::atomicOperationOrderAfterFenceSplit(
STI.getMinCmpXchgSizeInBits())
return AtomicOrdering::Acquire;
else if (auto *RI = dyn_cast<AtomicRMWInst>(I);
- RI &&
- RI->getOrdering() == AtomicOrdering::SequentiallyConsistent) {
+ RI && RI->getOrdering() == AtomicOrdering::SequentiallyConsistent) {
AtomicExpansionKind ExpansionKind = shouldExpandAtomicRMWInIR(RI);
if (ExpansionKind == AtomicExpansionKind::None ||
ExpansionKind == AtomicExpansionKind::Expand)
More information about the llvm-commits
mailing list