[llvm] b08cb38 - [AtomicExpand] Implement SUB → ADD(-x) and FSUB → FADD(-x) (#221425)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Oct 1 17:53:22 PDT 2026
Author: Yonah Goldberg
Date: 2026-10-02T00:53:14Z
New Revision: b08cb38bd7bcacb1a740f2dd566eef19c2b9ae40
URL: https://github.com/llvm/llvm-project/commit/b08cb38bd7bcacb1a740f2dd566eef19c2b9ae40
DIFF: https://github.com/llvm/llvm-project/commit/b08cb38bd7bcacb1a740f2dd566eef19c2b9ae40.diff
LOG: [AtomicExpand] Implement SUB → ADD(-x) and FSUB → FADD(-x) (#221425)
Implement ATOMIC_SUB → ATOMIC_ADD(-x) and ATOMIC_FSUB → ATOMIC_FADD(-x).
Many targets have atomic adds, but I don't think any have atomic subs.
This transform prevents cmpxchg expansion of atomic SUB. Enable these
transformations in NVPTX.
The integer variant exists in many different places currently. For
example, NVPTX currently expands it in DAG legalization.
begin AI generated
- SelectionDAG generic legalization:
[LegalizeDAG.cpp:3410](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp#L3410)
- GlobalISel generic legalization:
[LegalizerHelper.cpp:5117](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/CodeGen/GlobalISel/LegalizerHelper.cpp#L5117)
- GlobalISel outlined-atomic libcall path, mapping SUB to negated
`LDADD`:
[LegalizerHelper.cpp:919](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/CodeGen/GlobalISel/LegalizerHelper.cpp#L919)
and
[LegalizerHelper.cpp:973](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/CodeGen/GlobalISel/LegalizerHelper.cpp#L973)
- X86 SelectionDAG, when the result is used:
[X86ISelLowering.cpp:33966](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/X86/X86ISelLowering.cpp#L33966)
- SystemZ i32/i64:
[SystemZISelLowering.cpp:5281](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/SystemZ/SystemZISelLowering.cpp#L5281)
- SystemZ i8/i16, constant RHS only:
[SystemZISelLowering.cpp:5238](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/SystemZ/SystemZISelLowering.cpp#L5238)
- BPF i32 and i64 instruction-selection patterns:
[BPFInstrInfo.td:1040](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/BPF/BPFInstrInfo.td#L1040)
and
[BPFInstrInfo.td:1062](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/BPF/BPFInstrInfo.td#L1062)
- LoongArch i8/i16 patterns:
[LoongArchInstrInfo.td:2314](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/LoongArch/LoongArchInstrInfo.td#L2314)
- LoongArch i32/i64 patterns:
[LoongArchInstrInfo.td:2371](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/LoongArch/LoongArchInstrInfo.td#L2371)
- AArch64 GlobalISel/LSE pattern:
[AArch64InstrAtomics.td:546](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/AArch64/AArch64InstrAtomics.td#L546),
with the negating pattern machinery in
[AArch64InstrFormats.td:12625](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/AArch64/AArch64InstrFormats.td#L12625)
- AMDGPU changes only `atomic sub 0` to `atomic add 0`:
[SIISelLowering.cpp:21345](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/AMDGPU/SIISelLowering.cpp#L21345)
end AI generated
It looks like the only place that does the float transformation is
SPIRV:
begin AI generated
- SPIR-V converts `G_ATOMICRMW_FSUB` into `OpAtomicFAddEXT` with an
`OpFNegate`:
[SPIRVInstructionSelector.cpp:1413](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp#L1413),
with operand-negation machinery at
[SPIRVInstructionSelector.cpp:2614](https://github.com/llvm/llvm-project/blob/4977a0c815c39fd90bb67c270174f2356ee9b5a7/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp#L2614)
end AI generated
I guess less targets have a native atomic fadd, but it looks like
Aarch64 and AMDGPU do, so this could be useful there.
AI assisted.
Added:
llvm/test/Transforms/AtomicExpand/NVPTX/atomicrmw-sub.ll
Modified:
llvm/lib/CodeGen/AtomicExpandPass.cpp
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
llvm/test/Transforms/AtomicExpand/NVPTX/atomicrmw-i64-sm20.ll
Removed:
################################################################################
diff --git a/llvm/lib/CodeGen/AtomicExpandPass.cpp b/llvm/lib/CodeGen/AtomicExpandPass.cpp
index b3d186fdf5f01..a07fb82a88fd1 100644
--- a/llvm/lib/CodeGen/AtomicExpandPass.cpp
+++ b/llvm/lib/CodeGen/AtomicExpandPass.cpp
@@ -106,6 +106,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,
@@ -819,6 +820,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);
@@ -887,6 +910,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 33ff7f8180d65..5d78c3edbf154 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1134,8 +1134,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,
@@ -7711,8 +7709,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
@@ -7722,8 +7721,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()) {
@@ -7740,7 +7744,7 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
break;
}
if (UseNative)
- return AtomicExpansionKind::None;
+ return ExpansionKind;
}
if (Ty->isHalfTy()) {
@@ -7750,14 +7754,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.
@@ -7778,19 +7782,24 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
return AtomicExpansionKind::None;
[[fallthrough]];
case AtomicRMWInst::BinOp::Add:
- case AtomicRMWInst::BinOp::Sub:
+ case AtomicRMWInst::BinOp::Sub: {
+ AtomicExpansionKind ExpansionKind =
+ AI->getOperation() == AtomicRMWInst::BinOp::Sub
+ ? AtomicExpansionKind::Expand
+ : AtomicExpansionKind::None;
switch (BitWidth) {
case 8:
case 16:
return AtomicExpansionKind::CmpXChg;
case 32:
case 64:
- return AtomicExpansionKind::None;
+ return ExpansionKind;
case 128:
return AtomicExpansionKind::CmpXChg;
default:
llvm_unreachable("unsupported width encountered");
}
+ }
case AtomicRMWInst::BinOp::And:
case AtomicRMWInst::BinOp::Or:
case AtomicRMWInst::BinOp::Xor:
@@ -7891,9 +7900,12 @@ 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 9f3f24902100e..f67752f73eb2a 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
@@ -1885,71 +1885,47 @@ 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: 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: membar.cta;
-; SM60-NOFTZ-ALLOW-NEXT: st.param.b32 [func_retval0], %r1;
+; 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: 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: membar.cta;
-; SM60-FTZ-DISALLOW-NEXT: st.param.b32 [func_retval0], %r1;
+; 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: 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: membar.cta;
-; SM60-FTZ-ALLOW-NEXT: st.param.b32 [func_retval0], %r1;
+; 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
@@ -2407,24 +2383,16 @@ 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: ld.param.b64 %rd1, [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: 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: membar.cta;
-; SM60-NEXT: st.param.b64 [func_retval0], %rd1;
+; 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-i64-sm20.ll b/llvm/test/Transforms/AtomicExpand/NVPTX/atomicrmw-i64-sm20.ll
index 203cfbf89d1d3..2485d080efd8e 100644
--- a/llvm/test/Transforms/AtomicExpand/NVPTX/atomicrmw-i64-sm20.ll
+++ b/llvm/test/Transforms/AtomicExpand/NVPTX/atomicrmw-i64-sm20.ll
@@ -8,14 +8,16 @@ define i64 @add_sub_xchg(ptr addrspace(1) %addr, i64 %val) {
; SM20-LABEL: define i64 @add_sub_xchg(
; SM20-SAME: ptr addrspace(1) [[ADDR:%.*]], i64 [[VAL:%.*]]) #[[ATTR0:[0-9]+]] {
; SM20-NEXT: [[ADD:%.*]] = atomicrmw add ptr addrspace(1) [[ADDR]], i64 [[VAL]] monotonic, align 8
-; SM20-NEXT: [[SUB:%.*]] = atomicrmw sub ptr addrspace(1) [[ADDR]], i64 [[VAL]] monotonic, align 8
+; SM20-NEXT: [[NEG:%.*]] = sub i64 0, [[VAL]]
+; SM20-NEXT: [[SUB:%.*]] = atomicrmw add ptr addrspace(1) [[ADDR]], i64 [[NEG]] monotonic, align 8
; SM20-NEXT: [[XCHG:%.*]] = atomicrmw xchg ptr addrspace(1) [[ADDR]], i64 [[VAL]] monotonic, align 8
; SM20-NEXT: ret i64 [[XCHG]]
;
; SM32-LABEL: define i64 @add_sub_xchg(
; SM32-SAME: ptr addrspace(1) [[ADDR:%.*]], i64 [[VAL:%.*]]) #[[ATTR0:[0-9]+]] {
; SM32-NEXT: [[ADD:%.*]] = atomicrmw add ptr addrspace(1) [[ADDR]], i64 [[VAL]] monotonic, align 8
-; SM32-NEXT: [[SUB:%.*]] = atomicrmw sub ptr addrspace(1) [[ADDR]], i64 [[VAL]] monotonic, align 8
+; SM32-NEXT: [[NEG:%.*]] = sub i64 0, [[VAL]]
+; SM32-NEXT: [[SUB:%.*]] = atomicrmw add ptr addrspace(1) [[ADDR]], i64 [[NEG]] monotonic, align 8
; SM32-NEXT: [[XCHG:%.*]] = atomicrmw xchg ptr addrspace(1) [[ADDR]], i64 [[VAL]] monotonic, align 8
; SM32-NEXT: ret i64 [[XCHG]]
;
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
+}
More information about the llvm-commits
mailing list