Skip to content

Commit 369dbeb

Browse files
AlexMacleanAlexMaclean
committed
address comments
1 parent 62671f5 commit 369dbeb

File tree

1 file changed

+185
-43
lines changed

1 file changed

+185
-43
lines changed

‎llvm/test/CodeGen/NVPTX/local-stack-frame.ll‎

Lines changed: 185 additions & 43 deletions
Original file line numberDiff line numberDiff line change
@@ -1,35 +1,96 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
12
; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=PTX32
23
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=PTX64
34
; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -mtriple=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
45
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
56

67
; Ensure we access the local stack properly
78

8-
; PTX32: mov.u32 %SPL, __local_depot{{[0-9]+}};
9-
; PTX32: ld.param.u32 %r{{[0-9]+}}, [foo_param_0];
10-
; PTX32: add.u32 %r[[SP_REG:[0-9]+]], %SPL, 0;
11-
; PTX32: st.local.u32 [%r[[SP_REG]]], %r{{[0-9]+}};
12-
; PTX64: mov.u64 %SPL, __local_depot{{[0-9]+}};
13-
; PTX64: ld.param.u32 %r{{[0-9]+}}, [foo_param_0];
14-
; PTX64: add.u64 %rd[[SP_REG:[0-9]+]], %SPL, 0;
15-
; PTX64: st.local.u32 [%rd[[SP_REG]]], %r{{[0-9]+}};
169
define void @foo(i32 %a) {
10+
; PTX32-LABEL: foo(
11+
; PTX32: {
12+
; PTX32-NEXT: .local .align 4 .b8 __local_depot0[4];
13+
; PTX32-NEXT: .reg .b32 %SP;
14+
; PTX32-NEXT: .reg .b32 %SPL;
15+
; PTX32-NEXT: .reg .b32 %r<4>;
16+
; PTX32-EMPTY:
17+
; PTX32-NEXT: // %bb.0:
18+
; PTX32-NEXT: mov.u32 %SPL, __local_depot0;
19+
; PTX32-NEXT: ld.param.u32 %r1, [foo_param_0];
20+
; PTX32-NEXT: add.u32 %r3, %SPL, 0;
21+
; PTX32-NEXT: st.local.u32 [%r3], %r1;
22+
; PTX32-NEXT: ret;
23+
;
24+
; PTX64-LABEL: foo(
25+
; PTX64: {
26+
; PTX64-NEXT: .local .align 4 .b8 __local_depot0[4];
27+
; PTX64-NEXT: .reg .b64 %SP;
28+
; PTX64-NEXT: .reg .b64 %SPL;
29+
; PTX64-NEXT: .reg .b32 %r<2>;
30+
; PTX64-NEXT: .reg .b64 %rd<3>;
31+
; PTX64-EMPTY:
32+
; PTX64-NEXT: // %bb.0:
33+
; PTX64-NEXT: mov.u64 %SPL, __local_depot0;
34+
; PTX64-NEXT: ld.param.u32 %r1, [foo_param_0];
35+
; PTX64-NEXT: add.u64 %rd2, %SPL, 0;
36+
; PTX64-NEXT: st.local.u32 [%rd2], %r1;
37+
; PTX64-NEXT: ret;
1738
%local = alloca i32, align 4
1839
store volatile i32 %a, ptr %local
1940
ret void
2041
}
2142

22-
; PTX32: mov.u32 %SPL, __local_depot{{[0-9]+}};
23-
; PTX32: cvta.local.u32 %SP, %SPL;
24-
; PTX32: ld.param.u32 %r{{[0-9]+}}, [foo2_param_0];
25-
; PTX32: add.u32 %r[[SP_REG:[0-9]+]], %SPL, 0;
26-
; PTX32: st.local.u32 [%r[[SP_REG]]], %r{{[0-9]+}};
27-
; PTX64: mov.u64 %SPL, __local_depot{{[0-9]+}};
28-
; PTX64: cvta.local.u64 %SP, %SPL;
29-
; PTX64: ld.param.u32 %r{{[0-9]+}}, [foo2_param_0];
30-
; PTX64: add.u64 %rd[[SP_REG:[0-9]+]], %SPL, 0;
31-
; PTX64: st.local.u32 [%rd[[SP_REG]]], %r{{[0-9]+}};
3243
define ptx_kernel void @foo2(i32 %a) {
44+
; PTX32-LABEL: foo2(
45+
; PTX32: {
46+
; PTX32-NEXT: .local .align 4 .b8 __local_depot1[4];
47+
; PTX32-NEXT: .reg .b32 %SP;
48+
; PTX32-NEXT: .reg .b32 %SPL;
49+
; PTX32-NEXT: .reg .b32 %r<4>;
50+
; PTX32-EMPTY:
51+
; PTX32-NEXT: // %bb.0:
52+
; PTX32-NEXT: mov.u32 %SPL, __local_depot1;
53+
; PTX32-NEXT: cvta.local.u32 %SP, %SPL;
54+
; PTX32-NEXT: ld.param.u32 %r1, [foo2_param_0];
55+
; PTX32-NEXT: add.u32 %r2, %SP, 0;
56+
; PTX32-NEXT: add.u32 %r3, %SPL, 0;
57+
; PTX32-NEXT: st.local.u32 [%r3], %r1;
58+
; PTX32-NEXT: { // callseq 0, 0
59+
; PTX32-NEXT: .param .b32 param0;
60+
; PTX32-NEXT: st.param.b32 [param0], %r2;
61+
; PTX32-NEXT: call.uni
62+
; PTX32-NEXT: bar,
63+
; PTX32-NEXT: (
64+
; PTX32-NEXT: param0
65+
; PTX32-NEXT: );
66+
; PTX32-NEXT: } // callseq 0
67+
; PTX32-NEXT: ret;
68+
;
69+
; PTX64-LABEL: foo2(
70+
; PTX64: {
71+
; PTX64-NEXT: .local .align 4 .b8 __local_depot1[4];
72+
; PTX64-NEXT: .reg .b64 %SP;
73+
; PTX64-NEXT: .reg .b64 %SPL;
74+
; PTX64-NEXT: .reg .b32 %r<2>;
75+
; PTX64-NEXT: .reg .b64 %rd<3>;
76+
; PTX64-EMPTY:
77+
; PTX64-NEXT: // %bb.0:
78+
; PTX64-NEXT: mov.u64 %SPL, __local_depot1;
79+
; PTX64-NEXT: cvta.local.u64 %SP, %SPL;
80+
; PTX64-NEXT: ld.param.u32 %r1, [foo2_param_0];
81+
; PTX64-NEXT: add.u64 %rd1, %SP, 0;
82+
; PTX64-NEXT: add.u64 %rd2, %SPL, 0;
83+
; PTX64-NEXT: st.local.u32 [%rd2], %r1;
84+
; PTX64-NEXT: { // callseq 0, 0
85+
; PTX64-NEXT: .param .b64 param0;
86+
; PTX64-NEXT: st.param.b64 [param0], %rd1;
87+
; PTX64-NEXT: call.uni
88+
; PTX64-NEXT: bar,
89+
; PTX64-NEXT: (
90+
; PTX64-NEXT: param0
91+
; PTX64-NEXT: );
92+
; PTX64-NEXT: } // callseq 0
93+
; PTX64-NEXT: ret;
3394
%local = alloca i32, align 4
3495
store i32 %a, ptr %local
3596
call void @bar(ptr %local)
@@ -38,39 +99,120 @@ define ptx_kernel void @foo2(i32 %a) {
3899

39100
declare void @bar(ptr %a)
40101

41-
42-
; PTX32: mov.u32 %SPL, __local_depot{{[0-9]+}};
43-
; PTX32-NOT: cvta.local.u32 %SP, %SPL;
44-
; PTX32: ld.param.u32 %r{{[0-9]+}}, [foo3_param_0];
45-
; PTX32: add.u32 %r{{[0-9]+}}, %SPL, 0;
46-
; PTX32: st.local.u32 [%r{{[0-9]+}}], %r{{[0-9]+}};
47-
; PTX64: mov.u64 %SPL, __local_depot{{[0-9]+}};
48-
; PTX64-NOT: cvta.local.u64 %SP, %SPL;
49-
; PTX64: ld.param.u32 %r{{[0-9]+}}, [foo3_param_0];
50-
; PTX64: add.u64 %rd{{[0-9]+}}, %SPL, 0;
51-
; PTX64: st.local.u32 [%rd{{[0-9]+}}], %r{{[0-9]+}};
52102
define void @foo3(i32 %a) {
103+
; PTX32-LABEL: foo3(
104+
; PTX32: {
105+
; PTX32-NEXT: .local .align 4 .b8 __local_depot2[12];
106+
; PTX32-NEXT: .reg .b32 %SP;
107+
; PTX32-NEXT: .reg .b32 %SPL;
108+
; PTX32-NEXT: .reg .b32 %r<6>;
109+
; PTX32-EMPTY:
110+
; PTX32-NEXT: // %bb.0:
111+
; PTX32-NEXT: mov.u32 %SPL, __local_depot2;
112+
; PTX32-NEXT: ld.param.u32 %r1, [foo3_param_0];
113+
; PTX32-NEXT: add.u32 %r3, %SPL, 0;
114+
; PTX32-NEXT: shl.b32 %r4, %r1, 2;
115+
; PTX32-NEXT: add.s32 %r5, %r3, %r4;
116+
; PTX32-NEXT: st.local.u32 [%r5], %r1;
117+
; PTX32-NEXT: ret;
118+
;
119+
; PTX64-LABEL: foo3(
120+
; PTX64: {
121+
; PTX64-NEXT: .local .align 4 .b8 __local_depot2[12];
122+
; PTX64-NEXT: .reg .b64 %SP;
123+
; PTX64-NEXT: .reg .b64 %SPL;
124+
; PTX64-NEXT: .reg .b32 %r<2>;
125+
; PTX64-NEXT: .reg .b64 %rd<5>;
126+
; PTX64-EMPTY:
127+
; PTX64-NEXT: // %bb.0:
128+
; PTX64-NEXT: mov.u64 %SPL, __local_depot2;
129+
; PTX64-NEXT: ld.param.u32 %r1, [foo3_param_0];
130+
; PTX64-NEXT: add.u64 %rd2, %SPL, 0;
131+
; PTX64-NEXT: mul.wide.s32 %rd3, %r1, 4;
132+
; PTX64-NEXT: add.s64 %rd4, %rd2, %rd3;
133+
; PTX64-NEXT: st.local.u32 [%rd4], %r1;
134+
; PTX64-NEXT: ret;
53135
%local = alloca [3 x i32], align 4
54136
%1 = getelementptr inbounds i32, ptr %local, i32 %a
55137
store i32 %a, ptr %1
56138
ret void
57139
}
58140

59-
; PTX32: cvta.local.u32 %SP, %SPL;
60-
; PTX32: add.u32 {{%r[0-9]+}}, %SP, 0;
61-
; PTX32: add.u32 {{%r[0-9]+}}, %SPL, 0;
62-
; PTX32: add.u32 {{%r[0-9]+}}, %SP, 4;
63-
; PTX32: add.u32 {{%r[0-9]+}}, %SPL, 4;
64-
; PTX32: st.local.u32 [{{%r[0-9]+}}], {{%r[0-9]+}}
65-
; PTX32: st.local.u32 [{{%r[0-9]+}}], {{%r[0-9]+}}
66-
; PTX64: cvta.local.u64 %SP, %SPL;
67-
; PTX64: add.u64 {{%rd[0-9]+}}, %SP, 0;
68-
; PTX64: add.u64 {{%rd[0-9]+}}, %SPL, 0;
69-
; PTX64: add.u64 {{%rd[0-9]+}}, %SP, 4;
70-
; PTX64: add.u64 {{%rd[0-9]+}}, %SPL, 4;
71-
; PTX64: st.local.u32 [{{%rd[0-9]+}}], {{%r[0-9]+}}
72-
; PTX64: st.local.u32 [{{%rd[0-9]+}}], {{%r[0-9]+}}
73141
define void @foo4() {
142+
; PTX32-LABEL: foo4(
143+
; PTX32: {
144+
; PTX32-NEXT: .local .align 4 .b8 __local_depot3[8];
145+
; PTX32-NEXT: .reg .b32 %SP;
146+
; PTX32-NEXT: .reg .b32 %SPL;
147+
; PTX32-NEXT: .reg .b32 %r<6>;
148+
; PTX32-EMPTY:
149+
; PTX32-NEXT: // %bb.0:
150+
; PTX32-NEXT: mov.u32 %SPL, __local_depot3;
151+
; PTX32-NEXT: cvta.local.u32 %SP, %SPL;
152+
; PTX32-NEXT: add.u32 %r1, %SP, 0;
153+
; PTX32-NEXT: add.u32 %r2, %SPL, 0;
154+
; PTX32-NEXT: add.u32 %r3, %SP, 4;
155+
; PTX32-NEXT: add.u32 %r4, %SPL, 4;
156+
; PTX32-NEXT: mov.b32 %r5, 0;
157+
; PTX32-NEXT: st.local.u32 [%r2], %r5;
158+
; PTX32-NEXT: st.local.u32 [%r4], %r5;
159+
; PTX32-NEXT: { // callseq 1, 0
160+
; PTX32-NEXT: .param .b32 param0;
161+
; PTX32-NEXT: st.param.b32 [param0], %r1;
162+
; PTX32-NEXT: call.uni
163+
; PTX32-NEXT: bar,
164+
; PTX32-NEXT: (
165+
; PTX32-NEXT: param0
166+
; PTX32-NEXT: );
167+
; PTX32-NEXT: } // callseq 1
168+
; PTX32-NEXT: { // callseq 2, 0
169+
; PTX32-NEXT: .param .b32 param0;
170+
; PTX32-NEXT: st.param.b32 [param0], %r3;
171+
; PTX32-NEXT: call.uni
172+
; PTX32-NEXT: bar,
173+
; PTX32-NEXT: (
174+
; PTX32-NEXT: param0
175+
; PTX32-NEXT: );
176+
; PTX32-NEXT: } // callseq 2
177+
; PTX32-NEXT: ret;
178+
;
179+
; PTX64-LABEL: foo4(
180+
; PTX64: {
181+
; PTX64-NEXT: .local .align 4 .b8 __local_depot3[8];
182+
; PTX64-NEXT: .reg .b64 %SP;
183+
; PTX64-NEXT: .reg .b64 %SPL;
184+
; PTX64-NEXT: .reg .b32 %r<2>;
185+
; PTX64-NEXT: .reg .b64 %rd<5>;
186+
; PTX64-EMPTY:
187+
; PTX64-NEXT: // %bb.0:
188+
; PTX64-NEXT: mov.u64 %SPL, __local_depot3;
189+
; PTX64-NEXT: cvta.local.u64 %SP, %SPL;
190+
; PTX64-NEXT: add.u64 %rd1, %SP, 0;
191+
; PTX64-NEXT: add.u64 %rd2, %SPL, 0;
192+
; PTX64-NEXT: add.u64 %rd3, %SP, 4;
193+
; PTX64-NEXT: add.u64 %rd4, %SPL, 4;
194+
; PTX64-NEXT: mov.b32 %r1, 0;
195+
; PTX64-NEXT: st.local.u32 [%rd2], %r1;
196+
; PTX64-NEXT: st.local.u32 [%rd4], %r1;
197+
; PTX64-NEXT: { // callseq 1, 0
198+
; PTX64-NEXT: .param .b64 param0;
199+
; PTX64-NEXT: st.param.b64 [param0], %rd1;
200+
; PTX64-NEXT: call.uni
201+
; PTX64-NEXT: bar,
202+
; PTX64-NEXT: (
203+
; PTX64-NEXT: param0
204+
; PTX64-NEXT: );
205+
; PTX64-NEXT: } // callseq 1
206+
; PTX64-NEXT: { // callseq 2, 0
207+
; PTX64-NEXT: .param .b64 param0;
208+
; PTX64-NEXT: st.param.b64 [param0], %rd3;
209+
; PTX64-NEXT: call.uni
210+
; PTX64-NEXT: bar,
211+
; PTX64-NEXT: (
212+
; PTX64-NEXT: param0
213+
; PTX64-NEXT: );
214+
; PTX64-NEXT: } // callseq 2
215+
; PTX64-NEXT: ret;
74216
%A = alloca i32
75217
%B = alloca i32
76218
store i32 0, ptr %A

0 commit comments

Comments
 (0)