Skip to content

Commit 51ecddf

Browse files
DenisGZMDenisGZM
committed
Autogenerated test for cmpxchg
1 parent a33afbb commit 51ecddf

File tree

2 files changed

+186
-6
lines changed

2 files changed

+186
-6
lines changed

‎llvm/test/CodeGen/NVPTX/cmpxchg.ll

Lines changed: 185 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,4 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --default-march nvptx64 --version 5
12
; RUN: llc < %s -march=nvptx64 -mcpu=sm_32 | FileCheck %s --check-prefixes=SM30,CHECK
23
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_32 | %ptxas-verify %}
34
; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | FileCheck %s --check-prefixes=SM70,CHECK
@@ -8,30 +9,208 @@
89

910
; CHECK-LABEL: relaxed_sys_i8
1011
define i8 @relaxed_sys_i8(ptr %addr, i8 %cmp, i8 %new) {
11-
; SM30: atom.cas.b32
12-
; SM70: atom.cas.b16
12+
; SM30-LABEL: relaxed_sys_i8(
13+
; SM30: {
14+
; SM30-NEXT: .reg .pred %p<3>;
15+
; SM30-NEXT: .reg .b16 %rs<2>;
16+
; SM30-NEXT: .reg .b32 %r<21>;
17+
; SM30-NEXT: .reg .b64 %rd<3>;
18+
; SM30-EMPTY:
19+
; SM30-NEXT: // %bb.0:
20+
; SM30-NEXT: ld.param.u8 %rs1, [relaxed_sys_i8_param_2];
21+
; SM30-NEXT: ld.param.u64 %rd2, [relaxed_sys_i8_param_0];
22+
; SM30-NEXT: and.b64 %rd1, %rd2, -4;
23+
; SM30-NEXT: cvt.u32.u64 %r9, %rd2;
24+
; SM30-NEXT: and.b32 %r10, %r9, 3;
25+
; SM30-NEXT: shl.b32 %r1, %r10, 3;
26+
; SM30-NEXT: mov.b32 %r11, 255;
27+
; SM30-NEXT: shl.b32 %r12, %r11, %r1;
28+
; SM30-NEXT: not.b32 %r2, %r12;
29+
; SM30-NEXT: cvt.u32.u16 %r13, %rs1;
30+
; SM30-NEXT: and.b32 %r14, %r13, 255;
31+
; SM30-NEXT: shl.b32 %r3, %r14, %r1;
32+
; SM30-NEXT: ld.param.u8 %r15, [relaxed_sys_i8_param_1];
33+
; SM30-NEXT: shl.b32 %r4, %r15, %r1;
34+
; SM30-NEXT: ld.u32 %r16, [%rd1];
35+
; SM30-NEXT: and.b32 %r20, %r16, %r2;
36+
; SM30-NEXT: $L__BB0_1: // %partword.cmpxchg.loop
37+
; SM30-NEXT: // =>This Inner Loop Header: Depth=1
38+
; SM30-NEXT: or.b32 %r17, %r20, %r3;
39+
; SM30-NEXT: or.b32 %r18, %r20, %r4;
40+
; SM30-NEXT: atom.cas.b32 %r7, [%rd1], %r18, %r17;
41+
; SM30-NEXT: setp.eq.s32 %p1, %r7, %r18;
42+
; SM30-NEXT: @%p1 bra $L__BB0_3;
43+
; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
44+
; SM30-NEXT: // in Loop: Header=BB0_1 Depth=1
45+
; SM30-NEXT: and.b32 %r8, %r7, %r2;
46+
; SM30-NEXT: setp.ne.s32 %p2, %r20, %r8;
47+
; SM30-NEXT: mov.u32 %r20, %r8;
48+
; SM30-NEXT: @%p2 bra $L__BB0_1;
49+
; SM30-NEXT: $L__BB0_3: // %partword.cmpxchg.end
50+
; SM30-NEXT: st.param.b32 [func_retval0+0], %r13;
51+
; SM30-NEXT: ret;
52+
;
53+
; SM70-LABEL: relaxed_sys_i8(
54+
; SM70: {
55+
; SM70-NEXT: .reg .pred %p<3>;
56+
; SM70-NEXT: .reg .b16 %rs<17>;
57+
; SM70-NEXT: .reg .b32 %r<3>;
58+
; SM70-NEXT: .reg .b64 %rd<5>;
59+
; SM70-EMPTY:
60+
; SM70-NEXT: // %bb.0:
61+
; SM70-NEXT: ld.param.u8 %rs9, [relaxed_sys_i8_param_2];
62+
; SM70-NEXT: ld.param.u64 %rd2, [relaxed_sys_i8_param_0];
63+
; SM70-NEXT: and.b64 %rd1, %rd2, -2;
64+
; SM70-NEXT: ld.param.u8 %rs10, [relaxed_sys_i8_param_1];
65+
; SM70-NEXT: and.b64 %rd3, %rd2, 1;
66+
; SM70-NEXT: shl.b64 %rd4, %rd3, 3;
67+
; SM70-NEXT: cvt.u32.u64 %r1, %rd4;
68+
; SM70-NEXT: mov.u16 %rs11, 255;
69+
; SM70-NEXT: shl.b16 %rs12, %rs11, %r1;
70+
; SM70-NEXT: not.b16 %rs2, %rs12;
71+
; SM70-NEXT: shl.b16 %rs3, %rs9, %r1;
72+
; SM70-NEXT: shl.b16 %rs4, %rs10, %r1;
73+
; SM70-NEXT: ld.u16 %rs13, [%rd1];
74+
; SM70-NEXT: and.b16 %rs16, %rs13, %rs2;
75+
; SM70-NEXT: $L__BB0_1: // %partword.cmpxchg.loop
76+
; SM70-NEXT: // =>This Inner Loop Header: Depth=1
77+
; SM70-NEXT: or.b16 %rs14, %rs16, %rs3;
78+
; SM70-NEXT: or.b16 %rs15, %rs16, %rs4;
79+
; SM70-NEXT: atom.cas.b16 %rs7, [%rd1], %rs15, %rs14;
80+
; SM70-NEXT: setp.eq.s16 %p1, %rs7, %rs15;
81+
; SM70-NEXT: @%p1 bra $L__BB0_3;
82+
; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure
83+
; SM70-NEXT: // in Loop: Header=BB0_1 Depth=1
84+
; SM70-NEXT: and.b16 %rs8, %rs7, %rs2;
85+
; SM70-NEXT: setp.ne.s16 %p2, %rs16, %rs8;
86+
; SM70-NEXT: mov.u16 %rs16, %rs8;
87+
; SM70-NEXT: @%p2 bra $L__BB0_1;
88+
; SM70-NEXT: $L__BB0_3: // %partword.cmpxchg.end
89+
; SM70-NEXT: cvt.u32.u16 %r2, %rs9;
90+
; SM70-NEXT: st.param.b32 [func_retval0+0], %r2;
91+
; SM70-NEXT: ret;
1392
%pairold = cmpxchg ptr %addr, i8 %cmp, i8 %new seq_cst seq_cst
1493
ret i8 %new
1594
}
1695

1796
; CHECK-LABEL: relaxed_sys_i16
1897
define i16 @relaxed_sys_i16(ptr %addr, i16 %cmp, i16 %new) {
19-
; SM30: atom.cas.b32
20-
; SM70: atom.cas.b16
98+
; SM30-LABEL: relaxed_sys_i16(
99+
; SM30: {
100+
; SM30-NEXT: .reg .pred %p<3>;
101+
; SM30-NEXT: .reg .b16 %rs<2>;
102+
; SM30-NEXT: .reg .b32 %r<20>;
103+
; SM30-NEXT: .reg .b64 %rd<3>;
104+
; SM30-EMPTY:
105+
; SM30-NEXT: // %bb.0:
106+
; SM30-NEXT: ld.param.u16 %rs1, [relaxed_sys_i16_param_2];
107+
; SM30-NEXT: ld.param.u64 %rd2, [relaxed_sys_i16_param_0];
108+
; SM30-NEXT: and.b64 %rd1, %rd2, -4;
109+
; SM30-NEXT: ld.param.u16 %r9, [relaxed_sys_i16_param_1];
110+
; SM30-NEXT: cvt.u32.u64 %r10, %rd2;
111+
; SM30-NEXT: and.b32 %r11, %r10, 3;
112+
; SM30-NEXT: shl.b32 %r1, %r11, 3;
113+
; SM30-NEXT: mov.b32 %r12, 65535;
114+
; SM30-NEXT: shl.b32 %r13, %r12, %r1;
115+
; SM30-NEXT: not.b32 %r2, %r13;
116+
; SM30-NEXT: cvt.u32.u16 %r14, %rs1;
117+
; SM30-NEXT: shl.b32 %r3, %r14, %r1;
118+
; SM30-NEXT: shl.b32 %r4, %r9, %r1;
119+
; SM30-NEXT: ld.u32 %r15, [%rd1];
120+
; SM30-NEXT: and.b32 %r19, %r15, %r2;
121+
; SM30-NEXT: $L__BB1_1: // %partword.cmpxchg.loop
122+
; SM30-NEXT: // =>This Inner Loop Header: Depth=1
123+
; SM30-NEXT: or.b32 %r16, %r19, %r3;
124+
; SM30-NEXT: or.b32 %r17, %r19, %r4;
125+
; SM30-NEXT: atom.cas.b32 %r7, [%rd1], %r17, %r16;
126+
; SM30-NEXT: setp.eq.s32 %p1, %r7, %r17;
127+
; SM30-NEXT: @%p1 bra $L__BB1_3;
128+
; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure
129+
; SM30-NEXT: // in Loop: Header=BB1_1 Depth=1
130+
; SM30-NEXT: and.b32 %r8, %r7, %r2;
131+
; SM30-NEXT: setp.ne.s32 %p2, %r19, %r8;
132+
; SM30-NEXT: mov.u32 %r19, %r8;
133+
; SM30-NEXT: @%p2 bra $L__BB1_1;
134+
; SM30-NEXT: $L__BB1_3: // %partword.cmpxchg.end
135+
; SM30-NEXT: st.param.b32 [func_retval0+0], %r14;
136+
; SM30-NEXT: ret;
137+
;
138+
; SM70-LABEL: relaxed_sys_i16(
139+
; SM70: {
140+
; SM70-NEXT: .reg .b16 %rs<4>;
141+
; SM70-NEXT: .reg .b32 %r<2>;
142+
; SM70-NEXT: .reg .b64 %rd<2>;
143+
; SM70-EMPTY:
144+
; SM70-NEXT: // %bb.0:
145+
; SM70-NEXT: ld.param.u64 %rd1, [relaxed_sys_i16_param_0];
146+
; SM70-NEXT: ld.param.u16 %rs1, [relaxed_sys_i16_param_1];
147+
; SM70-NEXT: ld.param.u16 %rs2, [relaxed_sys_i16_param_2];
148+
; SM70-NEXT: atom.cas.b16 %rs3, [%rd1], %rs1, %rs2;
149+
; SM70-NEXT: cvt.u32.u16 %r1, %rs2;
150+
; SM70-NEXT: st.param.b32 [func_retval0+0], %r1;
151+
; SM70-NEXT: ret;
21152
%pairold = cmpxchg ptr %addr, i16 %cmp, i16 %new seq_cst seq_cst
22153
ret i16 %new
23154
}
24155

25156
; CHECK-LABEL: relaxed_sys_i32
26157
define i32 @relaxed_sys_i32(ptr %addr, i32 %cmp, i32 %new) {
27-
; CHECK: atom.cas.b32
158+
; SM30-LABEL: relaxed_sys_i32(
159+
; SM30: {
160+
; SM30-NEXT: .reg .b32 %r<4>;
161+
; SM30-NEXT: .reg .b64 %rd<2>;
162+
; SM30-EMPTY:
163+
; SM30-NEXT: // %bb.0:
164+
; SM30-NEXT: ld.param.u64 %rd1, [relaxed_sys_i32_param_0];
165+
; SM30-NEXT: ld.param.u32 %r1, [relaxed_sys_i32_param_1];
166+
; SM30-NEXT: ld.param.u32 %r2, [relaxed_sys_i32_param_2];
167+
; SM30-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;
168+
; SM30-NEXT: st.param.b32 [func_retval0+0], %r2;
169+
; SM30-NEXT: ret;
170+
;
171+
; SM70-LABEL: relaxed_sys_i32(
172+
; SM70: {
173+
; SM70-NEXT: .reg .b32 %r<4>;
174+
; SM70-NEXT: .reg .b64 %rd<2>;
175+
; SM70-EMPTY:
176+
; SM70-NEXT: // %bb.0:
177+
; SM70-NEXT: ld.param.u64 %rd1, [relaxed_sys_i32_param_0];
178+
; SM70-NEXT: ld.param.u32 %r1, [relaxed_sys_i32_param_1];
179+
; SM70-NEXT: ld.param.u32 %r2, [relaxed_sys_i32_param_2];
180+
; SM70-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;
181+
; SM70-NEXT: st.param.b32 [func_retval0+0], %r2;
182+
; SM70-NEXT: ret;
28183
%pairold = cmpxchg ptr %addr, i32 %cmp, i32 %new seq_cst seq_cst
29184
ret i32 %new
30185
}
31186

32187
; CHECK-LABEL: relaxed_sys_i64
33188
define i64 @relaxed_sys_i64(ptr %addr, i64 %cmp, i64 %new) {
34-
; CHECK: atom.cas.b64
189+
; SM30-LABEL: relaxed_sys_i64(
190+
; SM30: {
191+
; SM30-NEXT: .reg .b64 %rd<5>;
192+
; SM30-EMPTY:
193+
; SM30-NEXT: // %bb.0:
194+
; SM30-NEXT: ld.param.u64 %rd1, [relaxed_sys_i64_param_0];
195+
; SM30-NEXT: ld.param.u64 %rd2, [relaxed_sys_i64_param_1];
196+
; SM30-NEXT: ld.param.u64 %rd3, [relaxed_sys_i64_param_2];
197+
; SM30-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;
198+
; SM30-NEXT: st.param.b64 [func_retval0+0], %rd3;
199+
; SM30-NEXT: ret;
200+
;
201+
; SM70-LABEL: relaxed_sys_i64(
202+
; SM70: {
203+
; SM70-NEXT: .reg .b64 %rd<5>;
204+
; SM70-EMPTY:
205+
; SM70-NEXT: // %bb.0:
206+
; SM70-NEXT: ld.param.u64 %rd1, [relaxed_sys_i64_param_0];
207+
; SM70-NEXT: ld.param.u64 %rd2, [relaxed_sys_i64_param_1];
208+
; SM70-NEXT: ld.param.u64 %rd3, [relaxed_sys_i64_param_2];
209+
; SM70-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;
210+
; SM70-NEXT: st.param.b64 [func_retval0+0], %rd3;
211+
; SM70-NEXT: ret;
35212
%pairold = cmpxchg ptr %addr, i64 %cmp, i64 %new seq_cst seq_cst
36213
ret i64 %new
37214
}
215+
;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
216+
; CHECK: {{.*}}

‎llvm/utils/UpdateTestChecks/common.py

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -636,6 +636,7 @@ def get_triple_from_march(march):
636636
"amdgcn": "amdgcn",
637637
"r600": "r600",
638638
"mips": "mips",
639+
"nvptx64": "nvptx64",
639640
"sparc": "sparc",
640641
"hexagon": "hexagon",
641642
"ve": "ve",

0 commit comments

Comments
 (0)