Skip to content

[NVPTX] Fixup some issues introduced by 128-bit atomics #155921

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Aug 28, 2025
Merged

[NVPTX] Fixup some issues introduced by 128-bit atomics #155921

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2325,13 +2325,15 @@ void NVPTXDAGToDAGISel::selectAtomicSwap128(SDNode *N) {
MemSDNode *AN = cast<MemSDNode>(N);
SDLoc dl(N);

const SDValue Chain = N->getOperand(0);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

How does this change relate to the ToT CI failure?

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It doesn't directly, but is also an issue introduced in #154852.

const auto [Base, Offset] = selectADDR(N->getOperand(1), CurDAG);
SmallVector<SDValue, 5> Ops{Base, Offset};
Ops.append(N->op_begin() + 2, N->op_end());
Ops.append({
getI32Imm(getMemOrder(AN), dl),
getI32Imm(getAtomicScope(AN), dl),
getI32Imm(getAddrSpace(AN), dl),
Chain,
});

assert(N->getOpcode() == NVPTXISD::ATOMIC_CMP_SWAP_B128 ||
Expand Down
30 changes: 30 additions & 0 deletions llvm/test/CodeGen/NVPTX/atomics-b128.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -458,6 +458,7 @@ define i128 @test_cmpxchg_monotonic_seq_cst(ptr %addr, i128 %cmp, i128 %new) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_monotonic_seq_cst_param_0];
; CHECK-NEXT: fence.sc.sys;
; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_monotonic_seq_cst_param_1];
; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_monotonic_seq_cst_param_2];
; CHECK-NEXT: {
Expand Down Expand Up @@ -524,6 +525,7 @@ define i128 @test_cmpxchg_acquire_seq_cst(ptr %addr, i128 %cmp, i128 %new) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_acquire_seq_cst_param_0];
; CHECK-NEXT: fence.sc.sys;
; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_acquire_seq_cst_param_1];
; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_acquire_seq_cst_param_2];
; CHECK-NEXT: {
Expand Down Expand Up @@ -590,6 +592,7 @@ define i128 @test_cmpxchg_release_seq_cst(ptr %addr, i128 %cmp, i128 %new) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_release_seq_cst_param_0];
; CHECK-NEXT: fence.sc.sys;
; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_release_seq_cst_param_1];
; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_release_seq_cst_param_2];
; CHECK-NEXT: {
Expand Down Expand Up @@ -656,6 +659,7 @@ define i128 @test_cmpxchg_acq_rel_seq_cst(ptr %addr, i128 %cmp, i128 %new) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_acq_rel_seq_cst_param_0];
; CHECK-NEXT: fence.sc.sys;
; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_acq_rel_seq_cst_param_1];
; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_acq_rel_seq_cst_param_2];
; CHECK-NEXT: {
Expand All @@ -678,6 +682,7 @@ define i128 @test_cmpxchg_seq_cst_monotonic(ptr %addr, i128 %cmp, i128 %new) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_seq_cst_monotonic_param_0];
; CHECK-NEXT: fence.sc.sys;
; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_seq_cst_monotonic_param_1];
; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_seq_cst_monotonic_param_2];
; CHECK-NEXT: {
Expand All @@ -700,6 +705,7 @@ define i128 @test_cmpxchg_seq_cst_acquire(ptr %addr, i128 %cmp, i128 %new) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_seq_cst_acquire_param_0];
; CHECK-NEXT: fence.sc.sys;
; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_seq_cst_acquire_param_1];
; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_seq_cst_acquire_param_2];
; CHECK-NEXT: {
Expand All @@ -722,6 +728,7 @@ define i128 @test_cmpxchg_seq_cst_seq_cst(ptr %addr, i128 %cmp, i128 %new) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_seq_cst_seq_cst_param_0];
; CHECK-NEXT: fence.sc.sys;
; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_seq_cst_seq_cst_param_1];
; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_seq_cst_seq_cst_param_2];
; CHECK-NEXT: {
Expand Down Expand Up @@ -1001,3 +1008,26 @@ define i128 @test_atomicrmw_umax(ptr %ptr, i128 %val) {
%ret = atomicrmw umax ptr %ptr, i128 %val monotonic
ret i128 %ret
}


@si128 = internal addrspace(3) global i128 0, align 16

define void @test_atomicrmw_xchg_const() {
; CHECK-LABEL: test_atomicrmw_xchg_const(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-NEXT: // demoted variable
; CHECK-NEXT: .shared .align 16 .b8 si128[16];
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: mov.b64 %rd1, 0;
; CHECK-NEXT: mov.b64 %rd2, 23;
; CHECK-NEXT: {
; CHECK-NEXT: .reg .b128 amt, dst;
; CHECK-NEXT: mov.b128 amt, {%rd2, %rd1};
; CHECK-NEXT: atom.seq_cst.sys.shared.exch.b128 dst, [si128], amt;
; CHECK-NEXT: mov.b128 {%rd3, %rd4}, dst;
; CHECK-NEXT: }
; CHECK-NEXT: ret;
%res = atomicrmw xchg ptr addrspace(3) @si128, i128 23 seq_cst
ret void
}
10 changes: 0 additions & 10 deletions llvm/test/CodeGen/NVPTX/load-store-atomic.err.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,16 +2,6 @@

; CHECK: error: unsupported atomic store
; CHECK: error: unsupported atomic load
; CHECK: error: unsupported atomic store
; CHECK: error: unsupported atomic load

;; TODO: we could actually support this but we don't currently support b128
;; load lowering.
define void @test_i128_generic_atomic(ptr %a, ptr %b) {
%a.load = load atomic i128, ptr %a seq_cst, align 16
store atomic i128 %a.load, ptr %b seq_cst, align 16
ret void
}

define void @test_i256_global_atomic(ptr addrspace(1) %a, ptr addrspace(1) %b) {
%a.load = load atomic i256, ptr addrspace(1) %a seq_cst, align 32
Expand Down