[llvm] [NVPTX] Error on atomics with system scope pre-sm60 (PR #222458)
Yonah Goldberg via llvm-commits
llvm-commits at lists.llvm.org
Thu Sep 10 20:05:19 PDT 2026
https://github.com/YonahGoldberg updated https://github.com/llvm/llvm-project/pull/222458
>From d6c08ab70390f7eb67c601fde4a4abaa66c57c49 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Wed, 9 Sep 2026 21:47:10 +0000
Subject: [PATCH 1/2] fix bug
---
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 9 +-
.../NVPTX/atomics-system-scope-sm50.err.ll | 50 +
llvm/test/CodeGen/NVPTX/cmpxchg.ll | 1063 +++++++++--------
3 files changed, 598 insertions(+), 524 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/atomics-system-scope-sm50.err.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 9a58a7108a85f..68855f9e92819 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -676,9 +676,14 @@ static NVPTX::Scope resolveScope(NVPTX::Scope S, const NVPTXSubtarget *T) {
}
NVPTX::Scope NVPTXDAGToDAGISel::getAtomicScope(const MemSDNode *N) const {
- if (!Subtarget->hasAtomScope())
+ NVPTX::Scope Scope = resolveScope(Scopes[N->getSyncScopeID()], Subtarget);
+ if (!Subtarget->hasAtomScope()) {
+ if (Scope == NVPTX::Scope::System)
+ reportFatalUsageError(
+ "NVPTX system scope atomics require sm_60 or later");
return NVPTX::Scope::DefaultDevice;
- return resolveScope(Scopes[N->getSyncScopeID()], Subtarget);
+ }
+ return Scope;
}
namespace {
diff --git a/llvm/test/CodeGen/NVPTX/atomics-system-scope-sm50.err.ll b/llvm/test/CodeGen/NVPTX/atomics-system-scope-sm50.err.ll
new file mode 100644
index 0000000000000..dad55c89f6119
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/atomics-system-scope-sm50.err.ll
@@ -0,0 +1,50 @@
+; RUN: split-file %s %t
+; RUN: not llc -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx40 -filetype=null %t/cmpxchg.ll 2>&1 | FileCheck %s
+; RUN: not llc -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx40 -filetype=null %t/cmpxchg-partword.ll 2>&1 | FileCheck %s
+; RUN: not llc -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx40 -filetype=null %t/atomicrmw.ll 2>&1 | FileCheck %s
+; RUN: not llc -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx40 -filetype=null %t/atomicrmw-expand.ll 2>&1 | FileCheck %s
+; RUN: llc -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx40 %t/device.ll -o - | FileCheck %s --check-prefix=DEVICE
+
+; CHECK: LLVM ERROR: NVPTX system scope atomics require sm_60 or later
+
+;--- cmpxchg.ll
+define i32 @cmpxchg_system(ptr %addr, i32 %cmp, i32 %new) {
+ %result = cmpxchg ptr %addr, i32 %cmp, i32 %new monotonic monotonic
+ %value = extractvalue { i32, i1 } %result, 0
+ ret i32 %value
+}
+
+;--- cmpxchg-partword.ll
+define i8 @cmpxchg_system_i8(ptr %addr, i8 %cmp, i8 %new) {
+ %result = cmpxchg ptr %addr, i8 %cmp, i8 %new monotonic monotonic
+ %value = extractvalue { i8, i1 } %result, 0
+ ret i8 %value
+}
+
+;--- atomicrmw.ll
+define i32 @atomicrmw_system(ptr %addr, i32 %value) {
+ %result = atomicrmw add ptr %addr, i32 %value monotonic
+ ret i32 %result
+}
+
+;--- atomicrmw-expand.ll
+define i32 @atomicrmw_system_nand(ptr %addr, i32 %value) {
+ %result = atomicrmw nand ptr %addr, i32 %value monotonic
+ ret i32 %result
+}
+
+;--- device.ll
+; DEVICE-LABEL: cmpxchg_device(
+; DEVICE: atom.cas.b32
+define i32 @cmpxchg_device(ptr %addr, i32 %cmp, i32 %new) {
+ %result = cmpxchg ptr %addr, i32 %cmp, i32 %new syncscope("device") monotonic monotonic
+ %value = extractvalue { i32, i1 } %result, 0
+ ret i32 %value
+}
+
+; DEVICE-LABEL: atomicrmw_device(
+; DEVICE: atom.add.u32
+define i32 @atomicrmw_device(ptr %addr, i32 %value) {
+ %result = atomicrmw add ptr %addr, i32 %value syncscope("device") monotonic
+ ret i32 %result
+}
diff --git a/llvm/test/CodeGen/NVPTX/cmpxchg.ll b/llvm/test/CodeGen/NVPTX/cmpxchg.ll
index 89884a51fce1b..d0b802a81b918 100644
--- a/llvm/test/CodeGen/NVPTX/cmpxchg.ll
+++ b/llvm/test/CodeGen/NVPTX/cmpxchg.ll
@@ -1,54 +1,54 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --default-march nvptx64 --version 5
-; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_32 | FileCheck %s --check-prefixes=SM30,CHECK
-; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_32 | %ptxas-verify %}
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_60 -mattr=+ptx50 | FileCheck %s --check-prefixes=SM60,CHECK
+; RUN: %if ptxas-sm_60 && ptxas-isa-5.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_60 -mattr=+ptx50 | %ptxas-verify -arch=sm_60 %}
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | FileCheck %s --check-prefixes=SM70,CHECK
; RUN: %if ptxas-sm_70 && ptxas-isa-6.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | %ptxas-verify -arch=sm_70 %}
-; TODO: these are system scope, but are compiled to gpu scope..
; TODO: these are seq_cst, but are compiled to relaxed..
; CHECK-LABEL: relaxed_sys_i8
define i8 @relaxed_sys_i8(ptr %addr, i8 %cmp, i8 %new) {
-; SM30-LABEL: relaxed_sys_i8(
-; SM30: {
-; SM30-NEXT: .reg .pred %p<3>;
-; SM30-NEXT: .reg .b16 %rs<2>;
-; SM30-NEXT: .reg .b32 %r<17>;
-; SM30-NEXT: .reg .b64 %rd<3>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b8 %rs1, [relaxed_sys_i8_param_2];
-; SM30-NEXT: ld.param.b64 %rd2, [relaxed_sys_i8_param_0];
-; SM30-NEXT: and.b64 %rd1, %rd2, -4;
-; SM30-NEXT: ld.param.b8 %r7, [relaxed_sys_i8_param_1];
-; SM30-NEXT: cvt.u32.u64 %r8, %rd2;
-; SM30-NEXT: and.b32 %r9, %r8, 3;
-; SM30-NEXT: shl.b32 %r1, %r9, 3;
-; SM30-NEXT: mov.b32 %r10, 255;
-; SM30-NEXT: shl.b32 %r11, %r10, %r1;
-; SM30-NEXT: not.b32 %r2, %r11;
-; SM30-NEXT: cvt.u32.u16 %r12, %rs1;
-; SM30-NEXT: shl.b32 %r3, %r12, %r1;
-; SM30-NEXT: shl.b32 %r4, %r7, %r1;
-; SM30-NEXT: ld.volatile.b32 %r13, [%rd1];
-; SM30-NEXT: and.b32 %r16, %r13, %r2;
-; SM30-NEXT: $L__BB0_1: // %partword.cmpxchg.loop
-; SM30-NEXT: // =>This Inner Loop Header: Depth=1
-; SM30-NEXT: or.b32 %r14, %r16, %r3;
-; SM30-NEXT: or.b32 %r15, %r16, %r4;
-; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;
-; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;
-; SM30-NEXT: @%p1 bra $L__BB0_3;
-; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
-; SM30-NEXT: // in Loop: Header=BB0_1 Depth=1
-; SM30-NEXT: and.b32 %r6, %r5, %r2;
-; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;
-; SM30-NEXT: mov.b32 %r16, %r6;
-; SM30-NEXT: @%p2 bra $L__BB0_1;
-; SM30-NEXT: $L__BB0_3: // %partword.cmpxchg.end
-; SM30-NEXT: st.param.b32 [func_retval0], %r12;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: relaxed_sys_i8(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<3>;
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<17>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b8 %rs1, [relaxed_sys_i8_param_2];
+; SM60-NEXT: ld.param.b64 %rd2, [relaxed_sys_i8_param_0];
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: ld.param.b8 %r7, [relaxed_sys_i8_param_1];
+; SM60-NEXT: cvt.u32.u64 %r8, %rd2;
+; SM60-NEXT: and.b32 %r9, %r8, 3;
+; SM60-NEXT: shl.b32 %r1, %r9, 3;
+; SM60-NEXT: mov.b32 %r10, 255;
+; SM60-NEXT: shl.b32 %r11, %r10, %r1;
+; SM60-NEXT: not.b32 %r2, %r11;
+; SM60-NEXT: cvt.u32.u16 %r12, %rs1;
+; SM60-NEXT: shl.b32 %r3, %r12, %r1;
+; SM60-NEXT: shl.b32 %r4, %r7, %r1;
+; SM60-NEXT: ld.volatile.b32 %r13, [%rd1];
+; SM60-NEXT: and.b32 %r16, %r13, %r2;
+; SM60-NEXT: $L__BB0_1: // %partword.cmpxchg.loop
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: or.b32 %r14, %r16, %r3;
+; SM60-NEXT: or.b32 %r15, %r16, %r4;
+; SM60-NEXT: atom.sys.cas.b32 %r5, [%rd1], %r15, %r14;
+; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;
+; SM60-NEXT: @%p1 bra $L__BB0_3;
+; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure
+; SM60-NEXT: // in Loop: Header=BB0_1 Depth=1
+; SM60-NEXT: and.b32 %r6, %r5, %r2;
+; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;
+; SM60-NEXT: mov.b32 %r16, %r6;
+; SM60-NEXT: @%p2 bra $L__BB0_1;
+; SM60-NEXT: $L__BB0_3: // %partword.cmpxchg.end
+; SM60-NEXT: st.param.b32 [func_retval0], %r12;
+; SM60-NEXT: ret;
;
; SM70-LABEL: relaxed_sys_i8(
; SM70: {
@@ -134,46 +134,47 @@ define i8 @relaxed_sys_i8(ptr %addr, i8 %cmp, i8 %new) {
}
define i8 @acquire_sys_i8(ptr %addr, i8 %cmp, i8 %new) {
-; SM30-LABEL: acquire_sys_i8(
-; SM30: {
-; SM30-NEXT: .reg .pred %p<3>;
-; SM30-NEXT: .reg .b16 %rs<2>;
-; SM30-NEXT: .reg .b32 %r<17>;
-; SM30-NEXT: .reg .b64 %rd<3>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b8 %rs1, [acquire_sys_i8_param_2];
-; SM30-NEXT: ld.param.b64 %rd2, [acquire_sys_i8_param_0];
-; SM30-NEXT: and.b64 %rd1, %rd2, -4;
-; SM30-NEXT: ld.param.b8 %r7, [acquire_sys_i8_param_1];
-; SM30-NEXT: cvt.u32.u64 %r8, %rd2;
-; SM30-NEXT: and.b32 %r9, %r8, 3;
-; SM30-NEXT: shl.b32 %r1, %r9, 3;
-; SM30-NEXT: mov.b32 %r10, 255;
-; SM30-NEXT: shl.b32 %r11, %r10, %r1;
-; SM30-NEXT: not.b32 %r2, %r11;
-; SM30-NEXT: cvt.u32.u16 %r12, %rs1;
-; SM30-NEXT: shl.b32 %r3, %r12, %r1;
-; SM30-NEXT: shl.b32 %r4, %r7, %r1;
-; SM30-NEXT: ld.volatile.b32 %r13, [%rd1];
-; SM30-NEXT: and.b32 %r16, %r13, %r2;
-; SM30-NEXT: $L__BB1_1: // %partword.cmpxchg.loop
-; SM30-NEXT: // =>This Inner Loop Header: Depth=1
-; SM30-NEXT: or.b32 %r14, %r16, %r3;
-; SM30-NEXT: or.b32 %r15, %r16, %r4;
-; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;
-; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;
-; SM30-NEXT: @%p1 bra $L__BB1_3;
-; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
-; SM30-NEXT: // in Loop: Header=BB1_1 Depth=1
-; SM30-NEXT: and.b32 %r6, %r5, %r2;
-; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;
-; SM30-NEXT: mov.b32 %r16, %r6;
-; SM30-NEXT: @%p2 bra $L__BB1_1;
-; SM30-NEXT: $L__BB1_3: // %partword.cmpxchg.end
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: st.param.b32 [func_retval0], %r12;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: acquire_sys_i8(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<3>;
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<17>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b8 %rs1, [acquire_sys_i8_param_2];
+; SM60-NEXT: ld.param.b64 %rd2, [acquire_sys_i8_param_0];
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: ld.param.b8 %r7, [acquire_sys_i8_param_1];
+; SM60-NEXT: cvt.u32.u64 %r8, %rd2;
+; SM60-NEXT: and.b32 %r9, %r8, 3;
+; SM60-NEXT: shl.b32 %r1, %r9, 3;
+; SM60-NEXT: mov.b32 %r10, 255;
+; SM60-NEXT: shl.b32 %r11, %r10, %r1;
+; SM60-NEXT: not.b32 %r2, %r11;
+; SM60-NEXT: cvt.u32.u16 %r12, %rs1;
+; SM60-NEXT: shl.b32 %r3, %r12, %r1;
+; SM60-NEXT: shl.b32 %r4, %r7, %r1;
+; SM60-NEXT: ld.volatile.b32 %r13, [%rd1];
+; SM60-NEXT: and.b32 %r16, %r13, %r2;
+; SM60-NEXT: $L__BB1_1: // %partword.cmpxchg.loop
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: or.b32 %r14, %r16, %r3;
+; SM60-NEXT: or.b32 %r15, %r16, %r4;
+; SM60-NEXT: atom.sys.cas.b32 %r5, [%rd1], %r15, %r14;
+; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;
+; SM60-NEXT: @%p1 bra $L__BB1_3;
+; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure
+; SM60-NEXT: // in Loop: Header=BB1_1 Depth=1
+; SM60-NEXT: and.b32 %r6, %r5, %r2;
+; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;
+; SM60-NEXT: mov.b32 %r16, %r6;
+; SM60-NEXT: @%p2 bra $L__BB1_1;
+; SM60-NEXT: $L__BB1_3: // %partword.cmpxchg.end
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: st.param.b32 [func_retval0], %r12;
+; SM60-NEXT: ret;
;
; SM70-LABEL: acquire_sys_i8(
; SM70: {
@@ -261,46 +262,47 @@ define i8 @acquire_sys_i8(ptr %addr, i8 %cmp, i8 %new) {
}
define i8 @release_sys_i8(ptr %addr, i8 %cmp, i8 %new) {
-; SM30-LABEL: release_sys_i8(
-; SM30: {
-; SM30-NEXT: .reg .pred %p<3>;
-; SM30-NEXT: .reg .b16 %rs<2>;
-; SM30-NEXT: .reg .b32 %r<17>;
-; SM30-NEXT: .reg .b64 %rd<3>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b8 %rs1, [release_sys_i8_param_2];
-; SM30-NEXT: ld.param.b64 %rd2, [release_sys_i8_param_0];
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: ld.param.b8 %r7, [release_sys_i8_param_1];
-; SM30-NEXT: and.b64 %rd1, %rd2, -4;
-; SM30-NEXT: cvt.u32.u64 %r8, %rd2;
-; SM30-NEXT: and.b32 %r9, %r8, 3;
-; SM30-NEXT: shl.b32 %r1, %r9, 3;
-; SM30-NEXT: mov.b32 %r10, 255;
-; SM30-NEXT: shl.b32 %r11, %r10, %r1;
-; SM30-NEXT: not.b32 %r2, %r11;
-; SM30-NEXT: cvt.u32.u16 %r12, %rs1;
-; SM30-NEXT: shl.b32 %r3, %r12, %r1;
-; SM30-NEXT: shl.b32 %r4, %r7, %r1;
-; SM30-NEXT: ld.volatile.b32 %r13, [%rd1];
-; SM30-NEXT: and.b32 %r16, %r13, %r2;
-; SM30-NEXT: $L__BB2_1: // %partword.cmpxchg.loop
-; SM30-NEXT: // =>This Inner Loop Header: Depth=1
-; SM30-NEXT: or.b32 %r14, %r16, %r3;
-; SM30-NEXT: or.b32 %r15, %r16, %r4;
-; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;
-; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;
-; SM30-NEXT: @%p1 bra $L__BB2_3;
-; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
-; SM30-NEXT: // in Loop: Header=BB2_1 Depth=1
-; SM30-NEXT: and.b32 %r6, %r5, %r2;
-; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;
-; SM30-NEXT: mov.b32 %r16, %r6;
-; SM30-NEXT: @%p2 bra $L__BB2_1;
-; SM30-NEXT: $L__BB2_3: // %partword.cmpxchg.end
-; SM30-NEXT: st.param.b32 [func_retval0], %r12;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: release_sys_i8(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<3>;
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<17>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b8 %rs1, [release_sys_i8_param_2];
+; SM60-NEXT: ld.param.b64 %rd2, [release_sys_i8_param_0];
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: ld.param.b8 %r7, [release_sys_i8_param_1];
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: cvt.u32.u64 %r8, %rd2;
+; SM60-NEXT: and.b32 %r9, %r8, 3;
+; SM60-NEXT: shl.b32 %r1, %r9, 3;
+; SM60-NEXT: mov.b32 %r10, 255;
+; SM60-NEXT: shl.b32 %r11, %r10, %r1;
+; SM60-NEXT: not.b32 %r2, %r11;
+; SM60-NEXT: cvt.u32.u16 %r12, %rs1;
+; SM60-NEXT: shl.b32 %r3, %r12, %r1;
+; SM60-NEXT: shl.b32 %r4, %r7, %r1;
+; SM60-NEXT: ld.volatile.b32 %r13, [%rd1];
+; SM60-NEXT: and.b32 %r16, %r13, %r2;
+; SM60-NEXT: $L__BB2_1: // %partword.cmpxchg.loop
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: or.b32 %r14, %r16, %r3;
+; SM60-NEXT: or.b32 %r15, %r16, %r4;
+; SM60-NEXT: atom.sys.cas.b32 %r5, [%rd1], %r15, %r14;
+; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;
+; SM60-NEXT: @%p1 bra $L__BB2_3;
+; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure
+; SM60-NEXT: // in Loop: Header=BB2_1 Depth=1
+; SM60-NEXT: and.b32 %r6, %r5, %r2;
+; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;
+; SM60-NEXT: mov.b32 %r16, %r6;
+; SM60-NEXT: @%p2 bra $L__BB2_1;
+; SM60-NEXT: $L__BB2_3: // %partword.cmpxchg.end
+; SM60-NEXT: st.param.b32 [func_retval0], %r12;
+; SM60-NEXT: ret;
;
; SM70-LABEL: release_sys_i8(
; SM70: {
@@ -388,47 +390,48 @@ define i8 @release_sys_i8(ptr %addr, i8 %cmp, i8 %new) {
}
define i8 @acq_rel_sys_i8(ptr %addr, i8 %cmp, i8 %new) {
-; SM30-LABEL: acq_rel_sys_i8(
-; SM30: {
-; SM30-NEXT: .reg .pred %p<3>;
-; SM30-NEXT: .reg .b16 %rs<2>;
-; SM30-NEXT: .reg .b32 %r<17>;
-; SM30-NEXT: .reg .b64 %rd<3>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b8 %rs1, [acq_rel_sys_i8_param_2];
-; SM30-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i8_param_0];
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: ld.param.b8 %r7, [acq_rel_sys_i8_param_1];
-; SM30-NEXT: and.b64 %rd1, %rd2, -4;
-; SM30-NEXT: cvt.u32.u64 %r8, %rd2;
-; SM30-NEXT: and.b32 %r9, %r8, 3;
-; SM30-NEXT: shl.b32 %r1, %r9, 3;
-; SM30-NEXT: mov.b32 %r10, 255;
-; SM30-NEXT: shl.b32 %r11, %r10, %r1;
-; SM30-NEXT: not.b32 %r2, %r11;
-; SM30-NEXT: cvt.u32.u16 %r12, %rs1;
-; SM30-NEXT: shl.b32 %r3, %r12, %r1;
-; SM30-NEXT: shl.b32 %r4, %r7, %r1;
-; SM30-NEXT: ld.volatile.b32 %r13, [%rd1];
-; SM30-NEXT: and.b32 %r16, %r13, %r2;
-; SM30-NEXT: $L__BB3_1: // %partword.cmpxchg.loop
-; SM30-NEXT: // =>This Inner Loop Header: Depth=1
-; SM30-NEXT: or.b32 %r14, %r16, %r3;
-; SM30-NEXT: or.b32 %r15, %r16, %r4;
-; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;
-; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;
-; SM30-NEXT: @%p1 bra $L__BB3_3;
-; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
-; SM30-NEXT: // in Loop: Header=BB3_1 Depth=1
-; SM30-NEXT: and.b32 %r6, %r5, %r2;
-; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;
-; SM30-NEXT: mov.b32 %r16, %r6;
-; SM30-NEXT: @%p2 bra $L__BB3_1;
-; SM30-NEXT: $L__BB3_3: // %partword.cmpxchg.end
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: st.param.b32 [func_retval0], %r12;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: acq_rel_sys_i8(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<3>;
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<17>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b8 %rs1, [acq_rel_sys_i8_param_2];
+; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i8_param_0];
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: ld.param.b8 %r7, [acq_rel_sys_i8_param_1];
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: cvt.u32.u64 %r8, %rd2;
+; SM60-NEXT: and.b32 %r9, %r8, 3;
+; SM60-NEXT: shl.b32 %r1, %r9, 3;
+; SM60-NEXT: mov.b32 %r10, 255;
+; SM60-NEXT: shl.b32 %r11, %r10, %r1;
+; SM60-NEXT: not.b32 %r2, %r11;
+; SM60-NEXT: cvt.u32.u16 %r12, %rs1;
+; SM60-NEXT: shl.b32 %r3, %r12, %r1;
+; SM60-NEXT: shl.b32 %r4, %r7, %r1;
+; SM60-NEXT: ld.volatile.b32 %r13, [%rd1];
+; SM60-NEXT: and.b32 %r16, %r13, %r2;
+; SM60-NEXT: $L__BB3_1: // %partword.cmpxchg.loop
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: or.b32 %r14, %r16, %r3;
+; SM60-NEXT: or.b32 %r15, %r16, %r4;
+; SM60-NEXT: atom.sys.cas.b32 %r5, [%rd1], %r15, %r14;
+; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;
+; SM60-NEXT: @%p1 bra $L__BB3_3;
+; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure
+; SM60-NEXT: // in Loop: Header=BB3_1 Depth=1
+; SM60-NEXT: and.b32 %r6, %r5, %r2;
+; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;
+; SM60-NEXT: mov.b32 %r16, %r6;
+; SM60-NEXT: @%p2 bra $L__BB3_1;
+; SM60-NEXT: $L__BB3_3: // %partword.cmpxchg.end
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: st.param.b32 [func_retval0], %r12;
+; SM60-NEXT: ret;
;
; SM70-LABEL: acq_rel_sys_i8(
; SM70: {
@@ -518,47 +521,48 @@ define i8 @acq_rel_sys_i8(ptr %addr, i8 %cmp, i8 %new) {
}
define i8 @seq_cst_sys_i8(ptr %addr, i8 %cmp, i8 %new) {
-; SM30-LABEL: seq_cst_sys_i8(
-; SM30: {
-; SM30-NEXT: .reg .pred %p<3>;
-; SM30-NEXT: .reg .b16 %rs<2>;
-; SM30-NEXT: .reg .b32 %r<17>;
-; SM30-NEXT: .reg .b64 %rd<3>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b8 %rs1, [seq_cst_sys_i8_param_2];
-; SM30-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i8_param_0];
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: ld.param.b8 %r7, [seq_cst_sys_i8_param_1];
-; SM30-NEXT: and.b64 %rd1, %rd2, -4;
-; SM30-NEXT: cvt.u32.u64 %r8, %rd2;
-; SM30-NEXT: and.b32 %r9, %r8, 3;
-; SM30-NEXT: shl.b32 %r1, %r9, 3;
-; SM30-NEXT: mov.b32 %r10, 255;
-; SM30-NEXT: shl.b32 %r11, %r10, %r1;
-; SM30-NEXT: not.b32 %r2, %r11;
-; SM30-NEXT: cvt.u32.u16 %r12, %rs1;
-; SM30-NEXT: shl.b32 %r3, %r12, %r1;
-; SM30-NEXT: shl.b32 %r4, %r7, %r1;
-; SM30-NEXT: ld.volatile.b32 %r13, [%rd1];
-; SM30-NEXT: and.b32 %r16, %r13, %r2;
-; SM30-NEXT: $L__BB4_1: // %partword.cmpxchg.loop
-; SM30-NEXT: // =>This Inner Loop Header: Depth=1
-; SM30-NEXT: or.b32 %r14, %r16, %r3;
-; SM30-NEXT: or.b32 %r15, %r16, %r4;
-; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;
-; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;
-; SM30-NEXT: @%p1 bra $L__BB4_3;
-; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
-; SM30-NEXT: // in Loop: Header=BB4_1 Depth=1
-; SM30-NEXT: and.b32 %r6, %r5, %r2;
-; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;
-; SM30-NEXT: mov.b32 %r16, %r6;
-; SM30-NEXT: @%p2 bra $L__BB4_1;
-; SM30-NEXT: $L__BB4_3: // %partword.cmpxchg.end
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: st.param.b32 [func_retval0], %r12;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: seq_cst_sys_i8(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<3>;
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<17>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b8 %rs1, [seq_cst_sys_i8_param_2];
+; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i8_param_0];
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: ld.param.b8 %r7, [seq_cst_sys_i8_param_1];
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: cvt.u32.u64 %r8, %rd2;
+; SM60-NEXT: and.b32 %r9, %r8, 3;
+; SM60-NEXT: shl.b32 %r1, %r9, 3;
+; SM60-NEXT: mov.b32 %r10, 255;
+; SM60-NEXT: shl.b32 %r11, %r10, %r1;
+; SM60-NEXT: not.b32 %r2, %r11;
+; SM60-NEXT: cvt.u32.u16 %r12, %rs1;
+; SM60-NEXT: shl.b32 %r3, %r12, %r1;
+; SM60-NEXT: shl.b32 %r4, %r7, %r1;
+; SM60-NEXT: ld.volatile.b32 %r13, [%rd1];
+; SM60-NEXT: and.b32 %r16, %r13, %r2;
+; SM60-NEXT: $L__BB4_1: // %partword.cmpxchg.loop
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: or.b32 %r14, %r16, %r3;
+; SM60-NEXT: or.b32 %r15, %r16, %r4;
+; SM60-NEXT: atom.sys.cas.b32 %r5, [%rd1], %r15, %r14;
+; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;
+; SM60-NEXT: @%p1 bra $L__BB4_3;
+; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure
+; SM60-NEXT: // in Loop: Header=BB4_1 Depth=1
+; SM60-NEXT: and.b32 %r6, %r5, %r2;
+; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;
+; SM60-NEXT: mov.b32 %r16, %r6;
+; SM60-NEXT: @%p2 bra $L__BB4_1;
+; SM60-NEXT: $L__BB4_3: // %partword.cmpxchg.end
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: st.param.b32 [func_retval0], %r12;
+; SM60-NEXT: ret;
;
; SM70-LABEL: seq_cst_sys_i8(
; SM70: {
@@ -649,45 +653,46 @@ define i8 @seq_cst_sys_i8(ptr %addr, i8 %cmp, i8 %new) {
; CHECK-LABEL: relaxed_sys_i16
define i16 @relaxed_sys_i16(ptr %addr, i16 %cmp, i16 %new) {
-; SM30-LABEL: relaxed_sys_i16(
-; SM30: {
-; SM30-NEXT: .reg .pred %p<3>;
-; SM30-NEXT: .reg .b16 %rs<2>;
-; SM30-NEXT: .reg .b32 %r<17>;
-; SM30-NEXT: .reg .b64 %rd<3>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b16 %rs1, [relaxed_sys_i16_param_2];
-; SM30-NEXT: ld.param.b64 %rd2, [relaxed_sys_i16_param_0];
-; SM30-NEXT: and.b64 %rd1, %rd2, -4;
-; SM30-NEXT: ld.param.b16 %r7, [relaxed_sys_i16_param_1];
-; SM30-NEXT: cvt.u32.u64 %r8, %rd2;
-; SM30-NEXT: and.b32 %r9, %r8, 3;
-; SM30-NEXT: shl.b32 %r1, %r9, 3;
-; SM30-NEXT: mov.b32 %r10, 65535;
-; SM30-NEXT: shl.b32 %r11, %r10, %r1;
-; SM30-NEXT: not.b32 %r2, %r11;
-; SM30-NEXT: cvt.u32.u16 %r12, %rs1;
-; SM30-NEXT: shl.b32 %r3, %r12, %r1;
-; SM30-NEXT: shl.b32 %r4, %r7, %r1;
-; SM30-NEXT: ld.volatile.b32 %r13, [%rd1];
-; SM30-NEXT: and.b32 %r16, %r13, %r2;
-; SM30-NEXT: $L__BB5_1: // %partword.cmpxchg.loop
-; SM30-NEXT: // =>This Inner Loop Header: Depth=1
-; SM30-NEXT: or.b32 %r14, %r16, %r3;
-; SM30-NEXT: or.b32 %r15, %r16, %r4;
-; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;
-; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;
-; SM30-NEXT: @%p1 bra $L__BB5_3;
-; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
-; SM30-NEXT: // in Loop: Header=BB5_1 Depth=1
-; SM30-NEXT: and.b32 %r6, %r5, %r2;
-; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;
-; SM30-NEXT: mov.b32 %r16, %r6;
-; SM30-NEXT: @%p2 bra $L__BB5_1;
-; SM30-NEXT: $L__BB5_3: // %partword.cmpxchg.end
-; SM30-NEXT: st.param.b32 [func_retval0], %r12;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: relaxed_sys_i16(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<3>;
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<17>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b16 %rs1, [relaxed_sys_i16_param_2];
+; SM60-NEXT: ld.param.b64 %rd2, [relaxed_sys_i16_param_0];
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: ld.param.b16 %r7, [relaxed_sys_i16_param_1];
+; SM60-NEXT: cvt.u32.u64 %r8, %rd2;
+; SM60-NEXT: and.b32 %r9, %r8, 3;
+; SM60-NEXT: shl.b32 %r1, %r9, 3;
+; SM60-NEXT: mov.b32 %r10, 65535;
+; SM60-NEXT: shl.b32 %r11, %r10, %r1;
+; SM60-NEXT: not.b32 %r2, %r11;
+; SM60-NEXT: cvt.u32.u16 %r12, %rs1;
+; SM60-NEXT: shl.b32 %r3, %r12, %r1;
+; SM60-NEXT: shl.b32 %r4, %r7, %r1;
+; SM60-NEXT: ld.volatile.b32 %r13, [%rd1];
+; SM60-NEXT: and.b32 %r16, %r13, %r2;
+; SM60-NEXT: $L__BB5_1: // %partword.cmpxchg.loop
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: or.b32 %r14, %r16, %r3;
+; SM60-NEXT: or.b32 %r15, %r16, %r4;
+; SM60-NEXT: atom.sys.cas.b32 %r5, [%rd1], %r15, %r14;
+; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;
+; SM60-NEXT: @%p1 bra $L__BB5_3;
+; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure
+; SM60-NEXT: // in Loop: Header=BB5_1 Depth=1
+; SM60-NEXT: and.b32 %r6, %r5, %r2;
+; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;
+; SM60-NEXT: mov.b32 %r16, %r6;
+; SM60-NEXT: @%p2 bra $L__BB5_1;
+; SM60-NEXT: $L__BB5_3: // %partword.cmpxchg.end
+; SM60-NEXT: st.param.b32 [func_retval0], %r12;
+; SM60-NEXT: ret;
;
; SM70-LABEL: relaxed_sys_i16(
; SM70: {
@@ -772,46 +777,47 @@ define i16 @relaxed_sys_i16(ptr %addr, i16 %cmp, i16 %new) {
}
define i16 @acquire_sys_i16(ptr %addr, i16 %cmp, i16 %new) {
-; SM30-LABEL: acquire_sys_i16(
-; SM30: {
-; SM30-NEXT: .reg .pred %p<3>;
-; SM30-NEXT: .reg .b16 %rs<2>;
-; SM30-NEXT: .reg .b32 %r<17>;
-; SM30-NEXT: .reg .b64 %rd<3>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b16 %rs1, [acquire_sys_i16_param_2];
-; SM30-NEXT: ld.param.b64 %rd2, [acquire_sys_i16_param_0];
-; SM30-NEXT: and.b64 %rd1, %rd2, -4;
-; SM30-NEXT: ld.param.b16 %r7, [acquire_sys_i16_param_1];
-; SM30-NEXT: cvt.u32.u64 %r8, %rd2;
-; SM30-NEXT: and.b32 %r9, %r8, 3;
-; SM30-NEXT: shl.b32 %r1, %r9, 3;
-; SM30-NEXT: mov.b32 %r10, 65535;
-; SM30-NEXT: shl.b32 %r11, %r10, %r1;
-; SM30-NEXT: not.b32 %r2, %r11;
-; SM30-NEXT: cvt.u32.u16 %r12, %rs1;
-; SM30-NEXT: shl.b32 %r3, %r12, %r1;
-; SM30-NEXT: shl.b32 %r4, %r7, %r1;
-; SM30-NEXT: ld.volatile.b32 %r13, [%rd1];
-; SM30-NEXT: and.b32 %r16, %r13, %r2;
-; SM30-NEXT: $L__BB6_1: // %partword.cmpxchg.loop
-; SM30-NEXT: // =>This Inner Loop Header: Depth=1
-; SM30-NEXT: or.b32 %r14, %r16, %r3;
-; SM30-NEXT: or.b32 %r15, %r16, %r4;
-; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;
-; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;
-; SM30-NEXT: @%p1 bra $L__BB6_3;
-; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
-; SM30-NEXT: // in Loop: Header=BB6_1 Depth=1
-; SM30-NEXT: and.b32 %r6, %r5, %r2;
-; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;
-; SM30-NEXT: mov.b32 %r16, %r6;
-; SM30-NEXT: @%p2 bra $L__BB6_1;
-; SM30-NEXT: $L__BB6_3: // %partword.cmpxchg.end
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: st.param.b32 [func_retval0], %r12;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: acquire_sys_i16(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<3>;
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<17>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b16 %rs1, [acquire_sys_i16_param_2];
+; SM60-NEXT: ld.param.b64 %rd2, [acquire_sys_i16_param_0];
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: ld.param.b16 %r7, [acquire_sys_i16_param_1];
+; SM60-NEXT: cvt.u32.u64 %r8, %rd2;
+; SM60-NEXT: and.b32 %r9, %r8, 3;
+; SM60-NEXT: shl.b32 %r1, %r9, 3;
+; SM60-NEXT: mov.b32 %r10, 65535;
+; SM60-NEXT: shl.b32 %r11, %r10, %r1;
+; SM60-NEXT: not.b32 %r2, %r11;
+; SM60-NEXT: cvt.u32.u16 %r12, %rs1;
+; SM60-NEXT: shl.b32 %r3, %r12, %r1;
+; SM60-NEXT: shl.b32 %r4, %r7, %r1;
+; SM60-NEXT: ld.volatile.b32 %r13, [%rd1];
+; SM60-NEXT: and.b32 %r16, %r13, %r2;
+; SM60-NEXT: $L__BB6_1: // %partword.cmpxchg.loop
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: or.b32 %r14, %r16, %r3;
+; SM60-NEXT: or.b32 %r15, %r16, %r4;
+; SM60-NEXT: atom.sys.cas.b32 %r5, [%rd1], %r15, %r14;
+; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;
+; SM60-NEXT: @%p1 bra $L__BB6_3;
+; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure
+; SM60-NEXT: // in Loop: Header=BB6_1 Depth=1
+; SM60-NEXT: and.b32 %r6, %r5, %r2;
+; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;
+; SM60-NEXT: mov.b32 %r16, %r6;
+; SM60-NEXT: @%p2 bra $L__BB6_1;
+; SM60-NEXT: $L__BB6_3: // %partword.cmpxchg.end
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: st.param.b32 [func_retval0], %r12;
+; SM60-NEXT: ret;
;
; SM70-LABEL: acquire_sys_i16(
; SM70: {
@@ -898,46 +904,47 @@ define i16 @acquire_sys_i16(ptr %addr, i16 %cmp, i16 %new) {
}
define i16 @release_sys_i16(ptr %addr, i16 %cmp, i16 %new) {
-; SM30-LABEL: release_sys_i16(
-; SM30: {
-; SM30-NEXT: .reg .pred %p<3>;
-; SM30-NEXT: .reg .b16 %rs<2>;
-; SM30-NEXT: .reg .b32 %r<17>;
-; SM30-NEXT: .reg .b64 %rd<3>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b16 %rs1, [release_sys_i16_param_2];
-; SM30-NEXT: ld.param.b64 %rd2, [release_sys_i16_param_0];
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: ld.param.b16 %r7, [release_sys_i16_param_1];
-; SM30-NEXT: and.b64 %rd1, %rd2, -4;
-; SM30-NEXT: cvt.u32.u64 %r8, %rd2;
-; SM30-NEXT: and.b32 %r9, %r8, 3;
-; SM30-NEXT: shl.b32 %r1, %r9, 3;
-; SM30-NEXT: mov.b32 %r10, 65535;
-; SM30-NEXT: shl.b32 %r11, %r10, %r1;
-; SM30-NEXT: not.b32 %r2, %r11;
-; SM30-NEXT: cvt.u32.u16 %r12, %rs1;
-; SM30-NEXT: shl.b32 %r3, %r12, %r1;
-; SM30-NEXT: shl.b32 %r4, %r7, %r1;
-; SM30-NEXT: ld.volatile.b32 %r13, [%rd1];
-; SM30-NEXT: and.b32 %r16, %r13, %r2;
-; SM30-NEXT: $L__BB7_1: // %partword.cmpxchg.loop
-; SM30-NEXT: // =>This Inner Loop Header: Depth=1
-; SM30-NEXT: or.b32 %r14, %r16, %r3;
-; SM30-NEXT: or.b32 %r15, %r16, %r4;
-; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;
-; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;
-; SM30-NEXT: @%p1 bra $L__BB7_3;
-; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
-; SM30-NEXT: // in Loop: Header=BB7_1 Depth=1
-; SM30-NEXT: and.b32 %r6, %r5, %r2;
-; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;
-; SM30-NEXT: mov.b32 %r16, %r6;
-; SM30-NEXT: @%p2 bra $L__BB7_1;
-; SM30-NEXT: $L__BB7_3: // %partword.cmpxchg.end
-; SM30-NEXT: st.param.b32 [func_retval0], %r12;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: release_sys_i16(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<3>;
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<17>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b16 %rs1, [release_sys_i16_param_2];
+; SM60-NEXT: ld.param.b64 %rd2, [release_sys_i16_param_0];
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: ld.param.b16 %r7, [release_sys_i16_param_1];
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: cvt.u32.u64 %r8, %rd2;
+; SM60-NEXT: and.b32 %r9, %r8, 3;
+; SM60-NEXT: shl.b32 %r1, %r9, 3;
+; SM60-NEXT: mov.b32 %r10, 65535;
+; SM60-NEXT: shl.b32 %r11, %r10, %r1;
+; SM60-NEXT: not.b32 %r2, %r11;
+; SM60-NEXT: cvt.u32.u16 %r12, %rs1;
+; SM60-NEXT: shl.b32 %r3, %r12, %r1;
+; SM60-NEXT: shl.b32 %r4, %r7, %r1;
+; SM60-NEXT: ld.volatile.b32 %r13, [%rd1];
+; SM60-NEXT: and.b32 %r16, %r13, %r2;
+; SM60-NEXT: $L__BB7_1: // %partword.cmpxchg.loop
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: or.b32 %r14, %r16, %r3;
+; SM60-NEXT: or.b32 %r15, %r16, %r4;
+; SM60-NEXT: atom.sys.cas.b32 %r5, [%rd1], %r15, %r14;
+; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;
+; SM60-NEXT: @%p1 bra $L__BB7_3;
+; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure
+; SM60-NEXT: // in Loop: Header=BB7_1 Depth=1
+; SM60-NEXT: and.b32 %r6, %r5, %r2;
+; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;
+; SM60-NEXT: mov.b32 %r16, %r6;
+; SM60-NEXT: @%p2 bra $L__BB7_1;
+; SM60-NEXT: $L__BB7_3: // %partword.cmpxchg.end
+; SM60-NEXT: st.param.b32 [func_retval0], %r12;
+; SM60-NEXT: ret;
;
; SM70-LABEL: release_sys_i16(
; SM70: {
@@ -1024,47 +1031,48 @@ define i16 @release_sys_i16(ptr %addr, i16 %cmp, i16 %new) {
}
define i16 @acq_rel_sys_i16(ptr %addr, i16 %cmp, i16 %new) {
-; SM30-LABEL: acq_rel_sys_i16(
-; SM30: {
-; SM30-NEXT: .reg .pred %p<3>;
-; SM30-NEXT: .reg .b16 %rs<2>;
-; SM30-NEXT: .reg .b32 %r<17>;
-; SM30-NEXT: .reg .b64 %rd<3>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b16 %rs1, [acq_rel_sys_i16_param_2];
-; SM30-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i16_param_0];
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: ld.param.b16 %r7, [acq_rel_sys_i16_param_1];
-; SM30-NEXT: and.b64 %rd1, %rd2, -4;
-; SM30-NEXT: cvt.u32.u64 %r8, %rd2;
-; SM30-NEXT: and.b32 %r9, %r8, 3;
-; SM30-NEXT: shl.b32 %r1, %r9, 3;
-; SM30-NEXT: mov.b32 %r10, 65535;
-; SM30-NEXT: shl.b32 %r11, %r10, %r1;
-; SM30-NEXT: not.b32 %r2, %r11;
-; SM30-NEXT: cvt.u32.u16 %r12, %rs1;
-; SM30-NEXT: shl.b32 %r3, %r12, %r1;
-; SM30-NEXT: shl.b32 %r4, %r7, %r1;
-; SM30-NEXT: ld.volatile.b32 %r13, [%rd1];
-; SM30-NEXT: and.b32 %r16, %r13, %r2;
-; SM30-NEXT: $L__BB8_1: // %partword.cmpxchg.loop
-; SM30-NEXT: // =>This Inner Loop Header: Depth=1
-; SM30-NEXT: or.b32 %r14, %r16, %r3;
-; SM30-NEXT: or.b32 %r15, %r16, %r4;
-; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;
-; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;
-; SM30-NEXT: @%p1 bra $L__BB8_3;
-; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
-; SM30-NEXT: // in Loop: Header=BB8_1 Depth=1
-; SM30-NEXT: and.b32 %r6, %r5, %r2;
-; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;
-; SM30-NEXT: mov.b32 %r16, %r6;
-; SM30-NEXT: @%p2 bra $L__BB8_1;
-; SM30-NEXT: $L__BB8_3: // %partword.cmpxchg.end
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: st.param.b32 [func_retval0], %r12;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: acq_rel_sys_i16(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<3>;
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<17>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b16 %rs1, [acq_rel_sys_i16_param_2];
+; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i16_param_0];
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: ld.param.b16 %r7, [acq_rel_sys_i16_param_1];
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: cvt.u32.u64 %r8, %rd2;
+; SM60-NEXT: and.b32 %r9, %r8, 3;
+; SM60-NEXT: shl.b32 %r1, %r9, 3;
+; SM60-NEXT: mov.b32 %r10, 65535;
+; SM60-NEXT: shl.b32 %r11, %r10, %r1;
+; SM60-NEXT: not.b32 %r2, %r11;
+; SM60-NEXT: cvt.u32.u16 %r12, %rs1;
+; SM60-NEXT: shl.b32 %r3, %r12, %r1;
+; SM60-NEXT: shl.b32 %r4, %r7, %r1;
+; SM60-NEXT: ld.volatile.b32 %r13, [%rd1];
+; SM60-NEXT: and.b32 %r16, %r13, %r2;
+; SM60-NEXT: $L__BB8_1: // %partword.cmpxchg.loop
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: or.b32 %r14, %r16, %r3;
+; SM60-NEXT: or.b32 %r15, %r16, %r4;
+; SM60-NEXT: atom.sys.cas.b32 %r5, [%rd1], %r15, %r14;
+; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;
+; SM60-NEXT: @%p1 bra $L__BB8_3;
+; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure
+; SM60-NEXT: // in Loop: Header=BB8_1 Depth=1
+; SM60-NEXT: and.b32 %r6, %r5, %r2;
+; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;
+; SM60-NEXT: mov.b32 %r16, %r6;
+; SM60-NEXT: @%p2 bra $L__BB8_1;
+; SM60-NEXT: $L__BB8_3: // %partword.cmpxchg.end
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: st.param.b32 [func_retval0], %r12;
+; SM60-NEXT: ret;
;
; SM70-LABEL: acq_rel_sys_i16(
; SM70: {
@@ -1154,47 +1162,48 @@ define i16 @acq_rel_sys_i16(ptr %addr, i16 %cmp, i16 %new) {
; CHECK-LABEL: seq_cst_sys_i16
define i16 @seq_cst_sys_i16(ptr %addr, i16 %cmp, i16 %new) {
-; SM30-LABEL: seq_cst_sys_i16(
-; SM30: {
-; SM30-NEXT: .reg .pred %p<3>;
-; SM30-NEXT: .reg .b16 %rs<2>;
-; SM30-NEXT: .reg .b32 %r<17>;
-; SM30-NEXT: .reg .b64 %rd<3>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b16 %rs1, [seq_cst_sys_i16_param_2];
-; SM30-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i16_param_0];
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: ld.param.b16 %r7, [seq_cst_sys_i16_param_1];
-; SM30-NEXT: and.b64 %rd1, %rd2, -4;
-; SM30-NEXT: cvt.u32.u64 %r8, %rd2;
-; SM30-NEXT: and.b32 %r9, %r8, 3;
-; SM30-NEXT: shl.b32 %r1, %r9, 3;
-; SM30-NEXT: mov.b32 %r10, 65535;
-; SM30-NEXT: shl.b32 %r11, %r10, %r1;
-; SM30-NEXT: not.b32 %r2, %r11;
-; SM30-NEXT: cvt.u32.u16 %r12, %rs1;
-; SM30-NEXT: shl.b32 %r3, %r12, %r1;
-; SM30-NEXT: shl.b32 %r4, %r7, %r1;
-; SM30-NEXT: ld.volatile.b32 %r13, [%rd1];
-; SM30-NEXT: and.b32 %r16, %r13, %r2;
-; SM30-NEXT: $L__BB9_1: // %partword.cmpxchg.loop
-; SM30-NEXT: // =>This Inner Loop Header: Depth=1
-; SM30-NEXT: or.b32 %r14, %r16, %r3;
-; SM30-NEXT: or.b32 %r15, %r16, %r4;
-; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;
-; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;
-; SM30-NEXT: @%p1 bra $L__BB9_3;
-; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
-; SM30-NEXT: // in Loop: Header=BB9_1 Depth=1
-; SM30-NEXT: and.b32 %r6, %r5, %r2;
-; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;
-; SM30-NEXT: mov.b32 %r16, %r6;
-; SM30-NEXT: @%p2 bra $L__BB9_1;
-; SM30-NEXT: $L__BB9_3: // %partword.cmpxchg.end
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: st.param.b32 [func_retval0], %r12;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: seq_cst_sys_i16(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<3>;
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<17>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b16 %rs1, [seq_cst_sys_i16_param_2];
+; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i16_param_0];
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: ld.param.b16 %r7, [seq_cst_sys_i16_param_1];
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: cvt.u32.u64 %r8, %rd2;
+; SM60-NEXT: and.b32 %r9, %r8, 3;
+; SM60-NEXT: shl.b32 %r1, %r9, 3;
+; SM60-NEXT: mov.b32 %r10, 65535;
+; SM60-NEXT: shl.b32 %r11, %r10, %r1;
+; SM60-NEXT: not.b32 %r2, %r11;
+; SM60-NEXT: cvt.u32.u16 %r12, %rs1;
+; SM60-NEXT: shl.b32 %r3, %r12, %r1;
+; SM60-NEXT: shl.b32 %r4, %r7, %r1;
+; SM60-NEXT: ld.volatile.b32 %r13, [%rd1];
+; SM60-NEXT: and.b32 %r16, %r13, %r2;
+; SM60-NEXT: $L__BB9_1: // %partword.cmpxchg.loop
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: or.b32 %r14, %r16, %r3;
+; SM60-NEXT: or.b32 %r15, %r16, %r4;
+; SM60-NEXT: atom.sys.cas.b32 %r5, [%rd1], %r15, %r14;
+; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;
+; SM60-NEXT: @%p1 bra $L__BB9_3;
+; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure
+; SM60-NEXT: // in Loop: Header=BB9_1 Depth=1
+; SM60-NEXT: and.b32 %r6, %r5, %r2;
+; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;
+; SM60-NEXT: mov.b32 %r16, %r6;
+; SM60-NEXT: @%p2 bra $L__BB9_1;
+; SM60-NEXT: $L__BB9_3: // %partword.cmpxchg.end
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: st.param.b32 [func_retval0], %r12;
+; SM60-NEXT: ret;
;
; SM70-LABEL: seq_cst_sys_i16(
; SM70: {
@@ -1284,18 +1293,19 @@ define i16 @seq_cst_sys_i16(ptr %addr, i16 %cmp, i16 %new) {
; CHECK-LABEL: relaxed_sys_i32
define i32 @relaxed_sys_i32(ptr %addr, i32 %cmp, i32 %new) {
-; SM30-LABEL: relaxed_sys_i32(
-; SM30: {
-; SM30-NEXT: .reg .b32 %r<4>;
-; SM30-NEXT: .reg .b64 %rd<2>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b64 %rd1, [relaxed_sys_i32_param_0];
-; SM30-NEXT: ld.param.b32 %r1, [relaxed_sys_i32_param_1];
-; SM30-NEXT: ld.param.b32 %r2, [relaxed_sys_i32_param_2];
-; SM30-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;
-; SM30-NEXT: st.param.b32 [func_retval0], %r2;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: relaxed_sys_i32(
+; SM60: {
+; SM60-NEXT: .reg .b32 %r<4>;
+; SM60-NEXT: .reg .b64 %rd<2>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [relaxed_sys_i32_param_0];
+; SM60-NEXT: ld.param.b32 %r1, [relaxed_sys_i32_param_1];
+; SM60-NEXT: ld.param.b32 %r2, [relaxed_sys_i32_param_2];
+; SM60-NEXT: atom.sys.cas.b32 %r3, [%rd1], %r1, %r2;
+; SM60-NEXT: st.param.b32 [func_retval0], %r2;
+; SM60-NEXT: ret;
;
; SM70-LABEL: relaxed_sys_i32(
; SM70: {
@@ -1326,18 +1336,19 @@ define i32 @relaxed_sys_i32(ptr %addr, i32 %cmp, i32 %new) {
}
define i32 @acq_rel_sys_i32(ptr %addr, i32 %cmp, i32 %new) {
-; SM30-LABEL: acq_rel_sys_i32(
-; SM30: {
-; SM30-NEXT: .reg .b32 %r<4>;
-; SM30-NEXT: .reg .b64 %rd<2>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i32_param_0];
-; SM30-NEXT: ld.param.b32 %r1, [acq_rel_sys_i32_param_1];
-; SM30-NEXT: ld.param.b32 %r2, [acq_rel_sys_i32_param_2];
-; SM30-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;
-; SM30-NEXT: st.param.b32 [func_retval0], %r2;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: acq_rel_sys_i32(
+; SM60: {
+; SM60-NEXT: .reg .b32 %r<4>;
+; SM60-NEXT: .reg .b64 %rd<2>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i32_param_0];
+; SM60-NEXT: ld.param.b32 %r1, [acq_rel_sys_i32_param_1];
+; SM60-NEXT: ld.param.b32 %r2, [acq_rel_sys_i32_param_2];
+; SM60-NEXT: atom.sys.cas.b32 %r3, [%rd1], %r1, %r2;
+; SM60-NEXT: st.param.b32 [func_retval0], %r2;
+; SM60-NEXT: ret;
;
; SM70-LABEL: acq_rel_sys_i32(
; SM70: {
@@ -1368,18 +1379,19 @@ define i32 @acq_rel_sys_i32(ptr %addr, i32 %cmp, i32 %new) {
}
define i32 @acquire_sys_i32(ptr %addr, i32 %cmp, i32 %new) {
-; SM30-LABEL: acquire_sys_i32(
-; SM30: {
-; SM30-NEXT: .reg .b32 %r<4>;
-; SM30-NEXT: .reg .b64 %rd<2>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b64 %rd1, [acquire_sys_i32_param_0];
-; SM30-NEXT: ld.param.b32 %r1, [acquire_sys_i32_param_1];
-; SM30-NEXT: ld.param.b32 %r2, [acquire_sys_i32_param_2];
-; SM30-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;
-; SM30-NEXT: st.param.b32 [func_retval0], %r2;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: acquire_sys_i32(
+; SM60: {
+; SM60-NEXT: .reg .b32 %r<4>;
+; SM60-NEXT: .reg .b64 %rd<2>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [acquire_sys_i32_param_0];
+; SM60-NEXT: ld.param.b32 %r1, [acquire_sys_i32_param_1];
+; SM60-NEXT: ld.param.b32 %r2, [acquire_sys_i32_param_2];
+; SM60-NEXT: atom.sys.cas.b32 %r3, [%rd1], %r1, %r2;
+; SM60-NEXT: st.param.b32 [func_retval0], %r2;
+; SM60-NEXT: ret;
;
; SM70-LABEL: acquire_sys_i32(
; SM70: {
@@ -1410,18 +1422,19 @@ define i32 @acquire_sys_i32(ptr %addr, i32 %cmp, i32 %new) {
}
define i32 @release_sys_i32(ptr %addr, i32 %cmp, i32 %new) {
-; SM30-LABEL: release_sys_i32(
-; SM30: {
-; SM30-NEXT: .reg .b32 %r<4>;
-; SM30-NEXT: .reg .b64 %rd<2>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b64 %rd1, [release_sys_i32_param_0];
-; SM30-NEXT: ld.param.b32 %r1, [release_sys_i32_param_1];
-; SM30-NEXT: ld.param.b32 %r2, [release_sys_i32_param_2];
-; SM30-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;
-; SM30-NEXT: st.param.b32 [func_retval0], %r2;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: release_sys_i32(
+; SM60: {
+; SM60-NEXT: .reg .b32 %r<4>;
+; SM60-NEXT: .reg .b64 %rd<2>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [release_sys_i32_param_0];
+; SM60-NEXT: ld.param.b32 %r1, [release_sys_i32_param_1];
+; SM60-NEXT: ld.param.b32 %r2, [release_sys_i32_param_2];
+; SM60-NEXT: atom.sys.cas.b32 %r3, [%rd1], %r1, %r2;
+; SM60-NEXT: st.param.b32 [func_retval0], %r2;
+; SM60-NEXT: ret;
;
; SM70-LABEL: release_sys_i32(
; SM70: {
@@ -1452,19 +1465,20 @@ define i32 @release_sys_i32(ptr %addr, i32 %cmp, i32 %new) {
}
define i32 @seq_cst_sys_i32(ptr %addr, i32 %cmp, i32 %new) {
-; SM30-LABEL: seq_cst_sys_i32(
-; SM30: {
-; SM30-NEXT: .reg .b32 %r<4>;
-; SM30-NEXT: .reg .b64 %rd<2>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b64 %rd1, [seq_cst_sys_i32_param_0];
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: ld.param.b32 %r1, [seq_cst_sys_i32_param_1];
-; SM30-NEXT: ld.param.b32 %r2, [seq_cst_sys_i32_param_2];
-; SM30-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;
-; SM30-NEXT: st.param.b32 [func_retval0], %r2;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: seq_cst_sys_i32(
+; SM60: {
+; SM60-NEXT: .reg .b32 %r<4>;
+; SM60-NEXT: .reg .b64 %rd<2>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [seq_cst_sys_i32_param_0];
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: ld.param.b32 %r1, [seq_cst_sys_i32_param_1];
+; SM60-NEXT: ld.param.b32 %r2, [seq_cst_sys_i32_param_2];
+; SM60-NEXT: atom.sys.cas.b32 %r3, [%rd1], %r1, %r2;
+; SM60-NEXT: st.param.b32 [func_retval0], %r2;
+; SM60-NEXT: ret;
;
; SM70-LABEL: seq_cst_sys_i32(
; SM70: {
@@ -1499,17 +1513,18 @@ define i32 @seq_cst_sys_i32(ptr %addr, i32 %cmp, i32 %new) {
; CHECK-LABEL: relaxed_sys_i64
define i64 @relaxed_sys_i64(ptr %addr, i64 %cmp, i64 %new) {
-; SM30-LABEL: relaxed_sys_i64(
-; SM30: {
-; SM30-NEXT: .reg .b64 %rd<5>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b64 %rd1, [relaxed_sys_i64_param_0];
-; SM30-NEXT: ld.param.b64 %rd2, [relaxed_sys_i64_param_1];
-; SM30-NEXT: ld.param.b64 %rd3, [relaxed_sys_i64_param_2];
-; SM30-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;
-; SM30-NEXT: st.param.b64 [func_retval0], %rd3;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: relaxed_sys_i64(
+; SM60: {
+; SM60-NEXT: .reg .b64 %rd<5>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [relaxed_sys_i64_param_0];
+; SM60-NEXT: ld.param.b64 %rd2, [relaxed_sys_i64_param_1];
+; SM60-NEXT: ld.param.b64 %rd3, [relaxed_sys_i64_param_2];
+; SM60-NEXT: atom.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3;
+; SM60-NEXT: st.param.b64 [func_retval0], %rd3;
+; SM60-NEXT: ret;
;
; SM70-LABEL: relaxed_sys_i64(
; SM70: {
@@ -1538,17 +1553,18 @@ define i64 @relaxed_sys_i64(ptr %addr, i64 %cmp, i64 %new) {
}
define i64 @acquire_sys_i64(ptr %addr, i64 %cmp, i64 %new) {
-; SM30-LABEL: acquire_sys_i64(
-; SM30: {
-; SM30-NEXT: .reg .b64 %rd<5>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b64 %rd1, [acquire_sys_i64_param_0];
-; SM30-NEXT: ld.param.b64 %rd2, [acquire_sys_i64_param_1];
-; SM30-NEXT: ld.param.b64 %rd3, [acquire_sys_i64_param_2];
-; SM30-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;
-; SM30-NEXT: st.param.b64 [func_retval0], %rd3;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: acquire_sys_i64(
+; SM60: {
+; SM60-NEXT: .reg .b64 %rd<5>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [acquire_sys_i64_param_0];
+; SM60-NEXT: ld.param.b64 %rd2, [acquire_sys_i64_param_1];
+; SM60-NEXT: ld.param.b64 %rd3, [acquire_sys_i64_param_2];
+; SM60-NEXT: atom.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3;
+; SM60-NEXT: st.param.b64 [func_retval0], %rd3;
+; SM60-NEXT: ret;
;
; SM70-LABEL: acquire_sys_i64(
; SM70: {
@@ -1577,17 +1593,18 @@ define i64 @acquire_sys_i64(ptr %addr, i64 %cmp, i64 %new) {
}
define i64 @acq_rel_sys_i64(ptr %addr, i64 %cmp, i64 %new) {
-; SM30-LABEL: acq_rel_sys_i64(
-; SM30: {
-; SM30-NEXT: .reg .b64 %rd<5>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i64_param_0];
-; SM30-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i64_param_1];
-; SM30-NEXT: ld.param.b64 %rd3, [acq_rel_sys_i64_param_2];
-; SM30-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;
-; SM30-NEXT: st.param.b64 [func_retval0], %rd3;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: acq_rel_sys_i64(
+; SM60: {
+; SM60-NEXT: .reg .b64 %rd<5>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i64_param_0];
+; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i64_param_1];
+; SM60-NEXT: ld.param.b64 %rd3, [acq_rel_sys_i64_param_2];
+; SM60-NEXT: atom.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3;
+; SM60-NEXT: st.param.b64 [func_retval0], %rd3;
+; SM60-NEXT: ret;
;
; SM70-LABEL: acq_rel_sys_i64(
; SM70: {
@@ -1616,17 +1633,18 @@ define i64 @acq_rel_sys_i64(ptr %addr, i64 %cmp, i64 %new) {
}
define i64 @release_sys_i64(ptr %addr, i64 %cmp, i64 %new) {
-; SM30-LABEL: release_sys_i64(
-; SM30: {
-; SM30-NEXT: .reg .b64 %rd<5>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b64 %rd1, [release_sys_i64_param_0];
-; SM30-NEXT: ld.param.b64 %rd2, [release_sys_i64_param_1];
-; SM30-NEXT: ld.param.b64 %rd3, [release_sys_i64_param_2];
-; SM30-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;
-; SM30-NEXT: st.param.b64 [func_retval0], %rd3;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: release_sys_i64(
+; SM60: {
+; SM60-NEXT: .reg .b64 %rd<5>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [release_sys_i64_param_0];
+; SM60-NEXT: ld.param.b64 %rd2, [release_sys_i64_param_1];
+; SM60-NEXT: ld.param.b64 %rd3, [release_sys_i64_param_2];
+; SM60-NEXT: atom.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3;
+; SM60-NEXT: st.param.b64 [func_retval0], %rd3;
+; SM60-NEXT: ret;
;
; SM70-LABEL: release_sys_i64(
; SM70: {
@@ -1655,18 +1673,19 @@ define i64 @release_sys_i64(ptr %addr, i64 %cmp, i64 %new) {
}
define i64 @seq_cst_sys_i64(ptr %addr, i64 %cmp, i64 %new) {
-; SM30-LABEL: seq_cst_sys_i64(
-; SM30: {
-; SM30-NEXT: .reg .b64 %rd<5>;
-; SM30-EMPTY:
-; SM30-NEXT: // %bb.0:
-; SM30-NEXT: ld.param.b64 %rd1, [seq_cst_sys_i64_param_0];
-; SM30-NEXT: membar.sys;
-; SM30-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i64_param_1];
-; SM30-NEXT: ld.param.b64 %rd3, [seq_cst_sys_i64_param_2];
-; SM30-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;
-; SM30-NEXT: st.param.b64 [func_retval0], %rd3;
-; SM30-NEXT: ret;
+;
+; SM60-LABEL: seq_cst_sys_i64(
+; SM60: {
+; SM60-NEXT: .reg .b64 %rd<5>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [seq_cst_sys_i64_param_0];
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i64_param_1];
+; SM60-NEXT: ld.param.b64 %rd3, [seq_cst_sys_i64_param_2];
+; SM60-NEXT: atom.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3;
+; SM60-NEXT: st.param.b64 [func_retval0], %rd3;
+; SM60-NEXT: ret;
;
; SM70-LABEL: seq_cst_sys_i64(
; SM70: {
>From 22dcab3c3133ee084d5c66f7ce0c7066d3f728e0 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Fri, 11 Sep 2026 03:04:53 +0000
Subject: [PATCH 2/2] use DiagnosticInfoUnsupported
---
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 6 ++++--
.../NVPTX/atomics-system-scope-sm50.err.ll | 21 +++++++++++++++----
2 files changed, 21 insertions(+), 6 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 68855f9e92819..a9e93e817cb88 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -679,8 +679,10 @@ NVPTX::Scope NVPTXDAGToDAGISel::getAtomicScope(const MemSDNode *N) const {
NVPTX::Scope Scope = resolveScope(Scopes[N->getSyncScopeID()], Subtarget);
if (!Subtarget->hasAtomScope()) {
if (Scope == NVPTX::Scope::System)
- reportFatalUsageError(
- "NVPTX system scope atomics require sm_60 or later");
+ CurDAG->getContext()->diagnose(DiagnosticInfoUnsupported(
+ CurDAG->getMachineFunction().getFunction(),
+ "NVPTX system scope atomics require sm_60 or later",
+ N->getDebugLoc()));
return NVPTX::Scope::DefaultDevice;
}
return Scope;
diff --git a/llvm/test/CodeGen/NVPTX/atomics-system-scope-sm50.err.ll b/llvm/test/CodeGen/NVPTX/atomics-system-scope-sm50.err.ll
index dad55c89f6119..1b15f8f46fd08 100644
--- a/llvm/test/CodeGen/NVPTX/atomics-system-scope-sm50.err.ll
+++ b/llvm/test/CodeGen/NVPTX/atomics-system-scope-sm50.err.ll
@@ -1,19 +1,32 @@
; RUN: split-file %s %t
-; RUN: not llc -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx40 -filetype=null %t/cmpxchg.ll 2>&1 | FileCheck %s
+; RUN: not llc -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx40 -filetype=null %t/cmpxchg.ll 2>&1 | FileCheck %s --check-prefix=LOC
; RUN: not llc -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx40 -filetype=null %t/cmpxchg-partword.ll 2>&1 | FileCheck %s
; RUN: not llc -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx40 -filetype=null %t/atomicrmw.ll 2>&1 | FileCheck %s
; RUN: not llc -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx40 -filetype=null %t/atomicrmw-expand.ll 2>&1 | FileCheck %s
; RUN: llc -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx40 %t/device.ll -o - | FileCheck %s --check-prefix=DEVICE
-; CHECK: LLVM ERROR: NVPTX system scope atomics require sm_60 or later
+; CHECK: error: <unknown>:0:0: in function {{.*}}: NVPTX system scope atomics require sm_60 or later
+; LOC: error: system-atomic.c:7:9: in function cmpxchg_system i32 (ptr, i32, i32): NVPTX system scope atomics require sm_60 or later
;--- cmpxchg.ll
-define i32 @cmpxchg_system(ptr %addr, i32 %cmp, i32 %new) {
- %result = cmpxchg ptr %addr, i32 %cmp, i32 %new monotonic monotonic
+define i32 @cmpxchg_system(ptr %addr, i32 %cmp, i32 %new) !dbg !4 {
+ %result = cmpxchg ptr %addr, i32 %cmp, i32 %new monotonic monotonic, !dbg !7
%value = extractvalue { i32, i1 } %result, 0
ret i32 %value
}
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!2, !3}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C11, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug)
+!1 = !DIFile(filename: "system-atomic.c", directory: "/tmp")
+!2 = !{i32 7, !"Dwarf Version", i32 5}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = distinct !DISubprogram(name: "cmpxchg_system", scope: !1, file: !1, line: 1, type: !5, scopeLine: 1, spFlags: DISPFlagDefinition, unit: !0)
+!5 = !DISubroutineType(types: !6)
+!6 = !{}
+!7 = !DILocation(line: 7, column: 9, scope: !4)
+
;--- cmpxchg-partword.ll
define i8 @cmpxchg_system_i8(ptr %addr, i8 %cmp, i8 %new) {
%result = cmpxchg ptr %addr, i8 %cmp, i8 %new monotonic monotonic
More information about the llvm-commits
mailing list