Skip to content

Commit 6a08c58

Browse files
AlexMacleanAlexMaclean
committed
address comments -- auto-gen tests
1 parent c3ff962 commit 6a08c58

File tree

1 file changed

+202
-41
lines changed

1 file changed

+202
-41
lines changed

‎llvm/test/CodeGen/NVPTX/ldu-ldg.ll‎

Lines changed: 202 additions & 41 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: --version 5
12
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_32 | FileCheck %s
23
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_32 | %ptxas-verify %}
34

@@ -22,145 +23,305 @@ declare double @llvm.nvvm.ldg.global.f.f64.p1(ptr addrspace(1) %ptr, i32 %align)
2223
declare half @llvm.nvvm.ldg.global.f.f16.p1(ptr addrspace(1) %ptr, i32 %align)
2324
declare <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 %align)
2425

25-
; CHECK-LABEL: test_ldu_i8
2626
define i8 @test_ldu_i8(ptr addrspace(1) %ptr) {
27-
; CHECK: ldu.global.u8
27+
; CHECK-LABEL: test_ldu_i8(
28+
; CHECK: {
29+
; CHECK-NEXT: .reg .b16 %rs<2>;
30+
; CHECK-NEXT: .reg .b32 %r<3>;
31+
; CHECK-NEXT: .reg .b64 %rd<2>;
32+
; CHECK-EMPTY:
33+
; CHECK-NEXT: // %bb.0:
34+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_i8_param_0];
35+
; CHECK-NEXT: ldu.global.u8 %rs1, [%rd1];
36+
; CHECK-NEXT: cvt.u32.u16 %r1, %rs1;
37+
; CHECK-NEXT: and.b32 %r2, %r1, 255;
38+
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
39+
; CHECK-NEXT: ret;
2840
%val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)
2941
ret i8 %val
3042
}
3143

32-
; CHECK-LABEL: test_ldu_i16
3344
define i16 @test_ldu_i16(ptr addrspace(1) %ptr) {
34-
; CHECK: ldu.global.u16
45+
; CHECK-LABEL: test_ldu_i16(
46+
; CHECK: {
47+
; CHECK-NEXT: .reg .b16 %rs<2>;
48+
; CHECK-NEXT: .reg .b32 %r<2>;
49+
; CHECK-NEXT: .reg .b64 %rd<2>;
50+
; CHECK-EMPTY:
51+
; CHECK-NEXT: // %bb.0:
52+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_i16_param_0];
53+
; CHECK-NEXT: ldu.global.u16 %rs1, [%rd1];
54+
; CHECK-NEXT: cvt.u32.u16 %r1, %rs1;
55+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
56+
; CHECK-NEXT: ret;
3557
%val = tail call i16 @llvm.nvvm.ldu.global.i.i16.p1(ptr addrspace(1) %ptr, i32 2)
3658
ret i16 %val
3759
}
3860

39-
; CHECK-LABEL: test_ldu_i32
4061
define i32 @test_ldu_i32(ptr addrspace(1) %ptr) {
41-
; CHECK: ldu.global.u32
62+
; CHECK-LABEL: test_ldu_i32(
63+
; CHECK: {
64+
; CHECK-NEXT: .reg .b32 %r<2>;
65+
; CHECK-NEXT: .reg .b64 %rd<2>;
66+
; CHECK-EMPTY:
67+
; CHECK-NEXT: // %bb.0:
68+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_i32_param_0];
69+
; CHECK-NEXT: ldu.global.u32 %r1, [%rd1];
70+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
71+
; CHECK-NEXT: ret;
4272
%val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) %ptr, i32 4)
4373
ret i32 %val
4474
}
4575

46-
; CHECK-LABEL: test_ldu_i64
4776
define i64 @test_ldu_i64(ptr addrspace(1) %ptr) {
48-
; CHECK: ldu.global.u64
77+
; CHECK-LABEL: test_ldu_i64(
78+
; CHECK: {
79+
; CHECK-NEXT: .reg .b64 %rd<3>;
80+
; CHECK-EMPTY:
81+
; CHECK-NEXT: // %bb.0:
82+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_i64_param_0];
83+
; CHECK-NEXT: ldu.global.u64 %rd2, [%rd1];
84+
; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
85+
; CHECK-NEXT: ret;
4986
%val = tail call i64 @llvm.nvvm.ldu.global.i.i64.p1(ptr addrspace(1) %ptr, i32 8)
5087
ret i64 %val
5188
}
5289

53-
; CHECK-LABEL: test_ldu_p
5490
define ptr @test_ldu_p(ptr addrspace(1) %ptr) {
55-
; CHECK: ldu.global.u64
91+
; CHECK-LABEL: test_ldu_p(
92+
; CHECK: {
93+
; CHECK-NEXT: .reg .b64 %rd<3>;
94+
; CHECK-EMPTY:
95+
; CHECK-NEXT: // %bb.0:
96+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_p_param_0];
97+
; CHECK-NEXT: ldu.global.u64 %rd2, [%rd1];
98+
; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
99+
; CHECK-NEXT: ret;
56100
%val = tail call ptr @llvm.nvvm.ldu.global.p.p1(ptr addrspace(1) %ptr, i32 8)
57101
ret ptr %val
58102
}
59103

60-
61-
; CHECK-LABEL: test_ldu_f32
62104
define float @test_ldu_f32(ptr addrspace(1) %ptr) {
63-
; CHECK: ldu.global.f32
105+
; CHECK-LABEL: test_ldu_f32(
106+
; CHECK: {
107+
; CHECK-NEXT: .reg .f32 %f<2>;
108+
; CHECK-NEXT: .reg .b64 %rd<2>;
109+
; CHECK-EMPTY:
110+
; CHECK-NEXT: // %bb.0:
111+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_f32_param_0];
112+
; CHECK-NEXT: ldu.global.f32 %f1, [%rd1];
113+
; CHECK-NEXT: st.param.f32 [func_retval0], %f1;
114+
; CHECK-NEXT: ret;
64115
%val = tail call float @llvm.nvvm.ldu.global.f.f32.p1(ptr addrspace(1) %ptr, i32 4)
65116
ret float %val
66117
}
67118

68-
; CHECK-LABEL: test_ldu_f64
69119
define double @test_ldu_f64(ptr addrspace(1) %ptr) {
70-
; CHECK: ldu.global.f64
120+
; CHECK-LABEL: test_ldu_f64(
121+
; CHECK: {
122+
; CHECK-NEXT: .reg .b64 %rd<2>;
123+
; CHECK-NEXT: .reg .f64 %fd<2>;
124+
; CHECK-EMPTY:
125+
; CHECK-NEXT: // %bb.0:
126+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_f64_param_0];
127+
; CHECK-NEXT: ldu.global.f64 %fd1, [%rd1];
128+
; CHECK-NEXT: st.param.f64 [func_retval0], %fd1;
129+
; CHECK-NEXT: ret;
71130
%val = tail call double @llvm.nvvm.ldu.global.f.f64.p1(ptr addrspace(1) %ptr, i32 8)
72131
ret double %val
73132
}
74133

75-
; CHECK-LABEL: test_ldu_f16
76134
define half @test_ldu_f16(ptr addrspace(1) %ptr) {
77-
; CHECK: ldu.global.u16
135+
; CHECK-LABEL: test_ldu_f16(
136+
; CHECK: {
137+
; CHECK-NEXT: .reg .b16 %rs<2>;
138+
; CHECK-NEXT: .reg .b64 %rd<2>;
139+
; CHECK-EMPTY:
140+
; CHECK-NEXT: // %bb.0:
141+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_f16_param_0];
142+
; CHECK-NEXT: ldu.global.u16 %rs1, [%rd1];
143+
; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;
144+
; CHECK-NEXT: ret;
78145
%val = tail call half @llvm.nvvm.ldu.global.f.f16.p1(ptr addrspace(1) %ptr, i32 2)
79146
ret half %val
80147
}
81148

82-
; CHECK-LABEL: test_ldu_v2f16
83149
define <2 x half> @test_ldu_v2f16(ptr addrspace(1) %ptr) {
84-
; CHECK: ldu.global.u32
150+
; CHECK-LABEL: test_ldu_v2f16(
151+
; CHECK: {
152+
; CHECK-NEXT: .reg .b32 %r<2>;
153+
; CHECK-NEXT: .reg .b64 %rd<2>;
154+
; CHECK-EMPTY:
155+
; CHECK-NEXT: // %bb.0:
156+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_v2f16_param_0];
157+
; CHECK-NEXT: ldu.global.u32 %r1, [%rd1];
158+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
159+
; CHECK-NEXT: ret;
85160
%val = tail call <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)
86161
ret <2 x half> %val
87162
}
88163

89-
; CHECK-LABEL: test_ldg_i8
90164
define i8 @test_ldg_i8(ptr addrspace(1) %ptr) {
91-
; CHECK: ld.global.nc.u8
165+
; CHECK-LABEL: test_ldg_i8(
166+
; CHECK: {
167+
; CHECK-NEXT: .reg .b16 %rs<2>;
168+
; CHECK-NEXT: .reg .b32 %r<2>;
169+
; CHECK-NEXT: .reg .b64 %rd<2>;
170+
; CHECK-EMPTY:
171+
; CHECK-NEXT: // %bb.0:
172+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_i8_param_0];
173+
; CHECK-NEXT: ld.global.nc.u8 %rs1, [%rd1];
174+
; CHECK-NEXT: cvt.u32.u8 %r1, %rs1;
175+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
176+
; CHECK-NEXT: ret;
92177
%val = tail call i8 @llvm.nvvm.ldg.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)
93178
ret i8 %val
94179
}
95180

96-
; CHECK-LABEL: test_ldg_i16
97181
define i16 @test_ldg_i16(ptr addrspace(1) %ptr) {
98-
; CHECK: ld.global.nc.u16
182+
; CHECK-LABEL: test_ldg_i16(
183+
; CHECK: {
184+
; CHECK-NEXT: .reg .b16 %rs<2>;
185+
; CHECK-NEXT: .reg .b32 %r<2>;
186+
; CHECK-NEXT: .reg .b64 %rd<2>;
187+
; CHECK-EMPTY:
188+
; CHECK-NEXT: // %bb.0:
189+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_i16_param_0];
190+
; CHECK-NEXT: ld.global.nc.u16 %rs1, [%rd1];
191+
; CHECK-NEXT: cvt.u32.u16 %r1, %rs1;
192+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
193+
; CHECK-NEXT: ret;
99194
%val = tail call i16 @llvm.nvvm.ldg.global.i.i16.p1(ptr addrspace(1) %ptr, i32 2)
100195
ret i16 %val
101196
}
102197

103-
; CHECK-LABEL: test_ldg_i32
104198
define i32 @test_ldg_i32(ptr addrspace(1) %ptr) {
105-
; CHECK: ld.global.nc.u32
199+
; CHECK-LABEL: test_ldg_i32(
200+
; CHECK: {
201+
; CHECK-NEXT: .reg .b32 %r<2>;
202+
; CHECK-NEXT: .reg .b64 %rd<2>;
203+
; CHECK-EMPTY:
204+
; CHECK-NEXT: // %bb.0:
205+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_i32_param_0];
206+
; CHECK-NEXT: ld.global.nc.u32 %r1, [%rd1];
207+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
208+
; CHECK-NEXT: ret;
106209
%val = tail call i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) %ptr, i32 4)
107210
ret i32 %val
108211
}
109212

110-
; CHECK-LABEL: test_ldg_i64
111213
define i64 @test_ldg_i64(ptr addrspace(1) %ptr) {
112-
; CHECK: ld.global.nc.u64
214+
; CHECK-LABEL: test_ldg_i64(
215+
; CHECK: {
216+
; CHECK-NEXT: .reg .b64 %rd<3>;
217+
; CHECK-EMPTY:
218+
; CHECK-NEXT: // %bb.0:
219+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_i64_param_0];
220+
; CHECK-NEXT: ld.global.nc.u64 %rd2, [%rd1];
221+
; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
222+
; CHECK-NEXT: ret;
113223
%val = tail call i64 @llvm.nvvm.ldg.global.i.i64.p1(ptr addrspace(1) %ptr, i32 8)
114224
ret i64 %val
115225
}
116226

117-
; CHECK-LABEL: test_ldg_p
118227
define ptr @test_ldg_p(ptr addrspace(1) %ptr) {
119-
; CHECK: ld.global.nc.u64
228+
; CHECK-LABEL: test_ldg_p(
229+
; CHECK: {
230+
; CHECK-NEXT: .reg .b64 %rd<3>;
231+
; CHECK-EMPTY:
232+
; CHECK-NEXT: // %bb.0:
233+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_p_param_0];
234+
; CHECK-NEXT: ld.global.nc.u64 %rd2, [%rd1];
235+
; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
236+
; CHECK-NEXT: ret;
120237
%val = tail call ptr @llvm.nvvm.ldg.global.p.p1(ptr addrspace(1) %ptr, i32 8)
121238
ret ptr %val
122239
}
123240

124-
; CHECK-LABEL: test_ldg_f32
125241
define float @test_ldg_f32(ptr addrspace(1) %ptr) {
126-
; CHECK: ld.global.nc.f32
242+
; CHECK-LABEL: test_ldg_f32(
243+
; CHECK: {
244+
; CHECK-NEXT: .reg .f32 %f<2>;
245+
; CHECK-NEXT: .reg .b64 %rd<2>;
246+
; CHECK-EMPTY:
247+
; CHECK-NEXT: // %bb.0:
248+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_f32_param_0];
249+
; CHECK-NEXT: ld.global.nc.f32 %f1, [%rd1];
250+
; CHECK-NEXT: st.param.f32 [func_retval0], %f1;
251+
; CHECK-NEXT: ret;
127252
%val = tail call float @llvm.nvvm.ldg.global.f.f32.p1(ptr addrspace(1) %ptr, i32 4)
128253
ret float %val
129254
}
130255

131-
; CHECK-LABEL: test_ldg_f64
132256
define double @test_ldg_f64(ptr addrspace(1) %ptr) {
133-
; CHECK: ld.global.nc.f64
257+
; CHECK-LABEL: test_ldg_f64(
258+
; CHECK: {
259+
; CHECK-NEXT: .reg .b64 %rd<2>;
260+
; CHECK-NEXT: .reg .f64 %fd<2>;
261+
; CHECK-EMPTY:
262+
; CHECK-NEXT: // %bb.0:
263+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_f64_param_0];
264+
; CHECK-NEXT: ld.global.nc.f64 %fd1, [%rd1];
265+
; CHECK-NEXT: st.param.f64 [func_retval0], %fd1;
266+
; CHECK-NEXT: ret;
134267
%val = tail call double @llvm.nvvm.ldg.global.f.f64.p1(ptr addrspace(1) %ptr, i32 8)
135268
ret double %val
136269
}
137270

138-
; CHECK-LABEL: test_ldg_f16
139271
define half @test_ldg_f16(ptr addrspace(1) %ptr) {
140-
; CHECK: ld.global.nc.u16
272+
; CHECK-LABEL: test_ldg_f16(
273+
; CHECK: {
274+
; CHECK-NEXT: .reg .b16 %rs<2>;
275+
; CHECK-NEXT: .reg .b64 %rd<2>;
276+
; CHECK-EMPTY:
277+
; CHECK-NEXT: // %bb.0:
278+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_f16_param_0];
279+
; CHECK-NEXT: ld.global.nc.u16 %rs1, [%rd1];
280+
; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;
281+
; CHECK-NEXT: ret;
141282
%val = tail call half @llvm.nvvm.ldg.global.f.f16.p1(ptr addrspace(1) %ptr, i32 2)
142283
ret half %val
143284
}
144285

145-
; CHECK-LABEL: test_ldg_v2f16
146286
define <2 x half> @test_ldg_v2f16(ptr addrspace(1) %ptr) {
147-
; CHECK: ld.global.nc.u32
287+
; CHECK-LABEL: test_ldg_v2f16(
288+
; CHECK: {
289+
; CHECK-NEXT: .reg .b32 %r<2>;
290+
; CHECK-NEXT: .reg .b64 %rd<2>;
291+
; CHECK-EMPTY:
292+
; CHECK-NEXT: // %bb.0:
293+
; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_v2f16_param_0];
294+
; CHECK-NEXT: ld.global.nc.u32 %r1, [%rd1];
295+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
296+
; CHECK-NEXT: ret;
148297
%val = tail call <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)
149298
ret <2 x half> %val
150299
}
151300

152301
@g = addrspace(1) global i32 0
153302

154-
; CHECK-LABEL: test_ldg_asi
155303
define i32 @test_ldg_asi() {
156-
; CHECK: ld.global.nc.u32 %r{{.*}}, [g+4]
304+
; CHECK-LABEL: test_ldg_asi(
305+
; CHECK: {
306+
; CHECK-NEXT: .reg .b32 %r<2>;
307+
; CHECK-EMPTY:
308+
; CHECK-NEXT: // %bb.0:
309+
; CHECK-NEXT: ld.global.nc.u32 %r1, [g+4];
310+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
311+
; CHECK-NEXT: ret;
157312
%val = tail call i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) getelementptr (i8, ptr addrspace(1) @g, i32 4), i32 4)
158313
ret i32 %val
159314
}
160315

161-
; CHECK-LABEL: test_lug_asi
162316
define i32 @test_lug_asi() {
163-
; CHECK: ldu.global.u32 %r{{.*}}, [g+4]
317+
; CHECK-LABEL: test_lug_asi(
318+
; CHECK: {
319+
; CHECK-NEXT: .reg .b32 %r<2>;
320+
; CHECK-EMPTY:
321+
; CHECK-NEXT: // %bb.0:
322+
; CHECK-NEXT: ldu.global.u32 %r1, [g+4];
323+
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
324+
; CHECK-NEXT: ret;
164325
%val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) getelementptr (i8, ptr addrspace(1) @g, i32 4), i32 4)
165326
ret i32 %val
166327
}

0 commit comments

Comments
 (0)