xref: /llvm-project/llvm/test/CodeGen/NVPTX/vector-returns.ll (revision 310e79875752886a7713911e2a1ec14bc75bd4b3) 
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
3
4target triple = "nvptx-nvidia-cuda"
5
6define <3 x i64> @long3() {
7; CHECK-LABEL: long3(
8; CHECK:       {
9; CHECK-NEXT:    .reg .b64 %rd<2>;
10; CHECK-EMPTY:
11; CHECK-NEXT:  // %bb.0:
12; CHECK-NEXT:    mov.b64 %rd1, 0;
13; CHECK-NEXT:    st.param.v2.b64 [func_retval0], {%rd1, %rd1};
14; CHECK-NEXT:    st.param.b64 [func_retval0+16], %rd1;
15; CHECK-NEXT:    ret;
16  ret <3 x i64> zeroinitializer
17}
18
19define <2 x i64> @long2() {
20; CHECK-LABEL: long2(
21; CHECK:       {
22; CHECK-NEXT:    .reg .b64 %rd<2>;
23; CHECK-EMPTY:
24; CHECK-NEXT:  // %bb.0:
25; CHECK-NEXT:    mov.b64 %rd1, 0;
26; CHECK-NEXT:    st.param.v2.b64 [func_retval0], {%rd1, %rd1};
27; CHECK-NEXT:    ret;
28  ret <2 x i64> zeroinitializer
29}
30
31define <1 x i64> @long1() {
32; CHECK-LABEL: long1(
33; CHECK:       {
34; CHECK-NEXT:    .reg .b64 %rd<2>;
35; CHECK-EMPTY:
36; CHECK-NEXT:  // %bb.0:
37; CHECK-NEXT:    mov.b64 %rd1, 0;
38; CHECK-NEXT:    st.param.b64 [func_retval0], %rd1;
39; CHECK-NEXT:    ret;
40  ret <1 x i64> zeroinitializer
41}
42
43define <5 x i32> @int5() {
44; CHECK-LABEL: int5(
45; CHECK:       {
46; CHECK-NEXT:    .reg .b32 %r<2>;
47; CHECK-EMPTY:
48; CHECK-NEXT:  // %bb.0:
49; CHECK-NEXT:    mov.b32 %r1, 0;
50; CHECK-NEXT:    st.param.v4.b32 [func_retval0], {%r1, %r1, %r1, %r1};
51; CHECK-NEXT:    st.param.b32 [func_retval0+16], %r1;
52; CHECK-NEXT:    ret;
53  ret <5 x i32> zeroinitializer
54}
55
56define <4 x i32> @int4() {
57; CHECK-LABEL: int4(
58; CHECK:       {
59; CHECK-NEXT:    .reg .b32 %r<2>;
60; CHECK-EMPTY:
61; CHECK-NEXT:  // %bb.0:
62; CHECK-NEXT:    mov.b32 %r1, 0;
63; CHECK-NEXT:    st.param.v4.b32 [func_retval0], {%r1, %r1, %r1, %r1};
64; CHECK-NEXT:    ret;
65  ret <4 x i32> zeroinitializer
66}
67
68define <3 x i32> @int3() {
69; CHECK-LABEL: int3(
70; CHECK:       {
71; CHECK-NEXT:    .reg .b32 %r<2>;
72; CHECK-EMPTY:
73; CHECK-NEXT:  // %bb.0:
74; CHECK-NEXT:    mov.b32 %r1, 0;
75; CHECK-NEXT:    st.param.v2.b32 [func_retval0], {%r1, %r1};
76; CHECK-NEXT:    st.param.b32 [func_retval0+8], %r1;
77; CHECK-NEXT:    ret;
78  ret <3 x i32> zeroinitializer
79}
80
81define <2 x i32> @int2() {
82; CHECK-LABEL: int2(
83; CHECK:       {
84; CHECK-NEXT:    .reg .b32 %r<2>;
85; CHECK-EMPTY:
86; CHECK-NEXT:  // %bb.0:
87; CHECK-NEXT:    mov.b32 %r1, 0;
88; CHECK-NEXT:    st.param.v2.b32 [func_retval0], {%r1, %r1};
89; CHECK-NEXT:    ret;
90  ret <2 x i32> zeroinitializer
91}
92
93define <1 x i32> @int1() {
94; CHECK-LABEL: int1(
95; CHECK:       {
96; CHECK-NEXT:    .reg .b32 %r<2>;
97; CHECK-EMPTY:
98; CHECK-NEXT:  // %bb.0:
99; CHECK-NEXT:    mov.b32 %r1, 0;
100; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
101; CHECK-NEXT:    ret;
102  ret <1 x i32> zeroinitializer
103}
104
105define <9 x i16> @short9() {
106; CHECK-LABEL: short9(
107; CHECK:       {
108; CHECK-NEXT:    .reg .b16 %rs<2>;
109; CHECK-EMPTY:
110; CHECK-NEXT:  // %bb.0:
111; CHECK-NEXT:    mov.b16 %rs1, 0;
112; CHECK-NEXT:    st.param.v4.b16 [func_retval0], {%rs1, %rs1, %rs1, %rs1};
113; CHECK-NEXT:    st.param.v4.b16 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1};
114; CHECK-NEXT:    st.param.b16 [func_retval0+16], %rs1;
115; CHECK-NEXT:    ret;
116  ret <9 x i16> zeroinitializer
117}
118
119define <8 x i16> @short8() {
120; CHECK-LABEL: short8(
121; CHECK:       {
122; CHECK-NEXT:    .reg .b32 %r<2>;
123; CHECK-EMPTY:
124; CHECK-NEXT:  // %bb.0:
125; CHECK-NEXT:    mov.b32 %r1, 0;
126; CHECK-NEXT:    st.param.v4.b32 [func_retval0], {%r1, %r1, %r1, %r1};
127; CHECK-NEXT:    ret;
128  ret <8 x i16> zeroinitializer
129}
130
131define <7 x i16> @short7() {
132; CHECK-LABEL: short7(
133; CHECK:       {
134; CHECK-NEXT:    .reg .b16 %rs<2>;
135; CHECK-EMPTY:
136; CHECK-NEXT:  // %bb.0:
137; CHECK-NEXT:    mov.b16 %rs1, 0;
138; CHECK-NEXT:    st.param.v4.b16 [func_retval0], {%rs1, %rs1, %rs1, %rs1};
139; CHECK-NEXT:    st.param.v2.b16 [func_retval0+8], {%rs1, %rs1};
140; CHECK-NEXT:    st.param.b16 [func_retval0+12], %rs1;
141; CHECK-NEXT:    ret;
142  ret <7 x i16> zeroinitializer
143}
144
145define <5 x i16> @short5() {
146; CHECK-LABEL: short5(
147; CHECK:       {
148; CHECK-NEXT:    .reg .b16 %rs<2>;
149; CHECK-EMPTY:
150; CHECK-NEXT:  // %bb.0:
151; CHECK-NEXT:    mov.b16 %rs1, 0;
152; CHECK-NEXT:    st.param.v4.b16 [func_retval0], {%rs1, %rs1, %rs1, %rs1};
153; CHECK-NEXT:    st.param.b16 [func_retval0+8], %rs1;
154; CHECK-NEXT:    ret;
155  ret <5 x i16> zeroinitializer
156}
157
158define <4 x i16> @short4() {
159; CHECK-LABEL: short4(
160; CHECK:       {
161; CHECK-NEXT:    .reg .b32 %r<2>;
162; CHECK-EMPTY:
163; CHECK-NEXT:  // %bb.0:
164; CHECK-NEXT:    mov.b32 %r1, 0;
165; CHECK-NEXT:    st.param.v2.b32 [func_retval0], {%r1, %r1};
166; CHECK-NEXT:    ret;
167  ret <4 x i16> zeroinitializer
168}
169
170define <3 x i16> @short3() {
171; CHECK-LABEL: short3(
172; CHECK:       {
173; CHECK-NEXT:    .reg .b16 %rs<2>;
174; CHECK-EMPTY:
175; CHECK-NEXT:  // %bb.0:
176; CHECK-NEXT:    mov.b16 %rs1, 0;
177; CHECK-NEXT:    st.param.v2.b16 [func_retval0], {%rs1, %rs1};
178; CHECK-NEXT:    st.param.b16 [func_retval0+4], %rs1;
179; CHECK-NEXT:    ret;
180  ret <3 x i16> zeroinitializer
181}
182
183define <2 x i16> @short2() {
184; CHECK-LABEL: short2(
185; CHECK:       {
186; CHECK-NEXT:    .reg .b32 %r<2>;
187; CHECK-EMPTY:
188; CHECK-NEXT:  // %bb.0:
189; CHECK-NEXT:    mov.b32 %r1, 0;
190; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
191; CHECK-NEXT:    ret;
192  ret <2 x i16> zeroinitializer
193}
194
195define <1 x i16> @short1() {
196; CHECK-LABEL: short1(
197; CHECK:       {
198; CHECK-NEXT:    .reg .b16 %rs<2>;
199; CHECK-EMPTY:
200; CHECK-NEXT:  // %bb.0:
201; CHECK-NEXT:    mov.b16 %rs1, 0;
202; CHECK-NEXT:    st.param.b16 [func_retval0], %rs1;
203; CHECK-NEXT:    ret;
204  ret <1 x i16> zeroinitializer
205}
206
207define <17 x i8> @byte17() {
208; CHECK-LABEL: byte17(
209; CHECK:       {
210; CHECK-NEXT:    .reg .b16 %rs<2>;
211; CHECK-EMPTY:
212; CHECK-NEXT:  // %bb.0:
213; CHECK-NEXT:    mov.b16 %rs1, 0;
214; CHECK-NEXT:    st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1};
215; CHECK-NEXT:    st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1};
216; CHECK-NEXT:    st.param.v4.b8 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1};
217; CHECK-NEXT:    st.param.v4.b8 [func_retval0+12], {%rs1, %rs1, %rs1, %rs1};
218; CHECK-NEXT:    st.param.b8 [func_retval0+16], %rs1;
219; CHECK-NEXT:    ret;
220  ret <17 x i8> zeroinitializer
221}
222
223define <16 x i8> @byte16() {
224; CHECK-LABEL: byte16(
225; CHECK:       {
226; CHECK-NEXT:    .reg .b32 %r<2>;
227; CHECK-EMPTY:
228; CHECK-NEXT:  // %bb.0:
229; CHECK-NEXT:    mov.b32 %r1, 0;
230; CHECK-NEXT:    st.param.v4.b32 [func_retval0], {%r1, %r1, %r1, %r1};
231; CHECK-NEXT:    ret;
232  ret <16 x i8> zeroinitializer
233}
234
235define <15 x i8> @byte15() {
236; CHECK-LABEL: byte15(
237; CHECK:       {
238; CHECK-NEXT:    .reg .b16 %rs<2>;
239; CHECK-EMPTY:
240; CHECK-NEXT:  // %bb.0:
241; CHECK-NEXT:    mov.b16 %rs1, 0;
242; CHECK-NEXT:    st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1};
243; CHECK-NEXT:    st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1};
244; CHECK-NEXT:    st.param.v4.b8 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1};
245; CHECK-NEXT:    st.param.v2.b8 [func_retval0+12], {%rs1, %rs1};
246; CHECK-NEXT:    st.param.b8 [func_retval0+14], %rs1;
247; CHECK-NEXT:    ret;
248  ret <15 x i8> zeroinitializer
249}
250
251define <9 x i8> @byte9() {
252; CHECK-LABEL: byte9(
253; CHECK:       {
254; CHECK-NEXT:    .reg .b16 %rs<2>;
255; CHECK-EMPTY:
256; CHECK-NEXT:  // %bb.0:
257; CHECK-NEXT:    mov.b16 %rs1, 0;
258; CHECK-NEXT:    st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1};
259; CHECK-NEXT:    st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1};
260; CHECK-NEXT:    st.param.b8 [func_retval0+8], %rs1;
261; CHECK-NEXT:    ret;
262  ret <9 x i8> zeroinitializer
263}
264
265define <8 x i8> @byte8() {
266; CHECK-LABEL: byte8(
267; CHECK:       {
268; CHECK-NEXT:    .reg .b32 %r<2>;
269; CHECK-EMPTY:
270; CHECK-NEXT:  // %bb.0:
271; CHECK-NEXT:    mov.b32 %r1, 0;
272; CHECK-NEXT:    st.param.v2.b32 [func_retval0], {%r1, %r1};
273; CHECK-NEXT:    ret;
274  ret <8 x i8> zeroinitializer
275}
276
277define <7 x i8> @byte7() {
278; CHECK-LABEL: byte7(
279; CHECK:       {
280; CHECK-NEXT:    .reg .b16 %rs<2>;
281; CHECK-EMPTY:
282; CHECK-NEXT:  // %bb.0:
283; CHECK-NEXT:    mov.b16 %rs1, 0;
284; CHECK-NEXT:    st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1};
285; CHECK-NEXT:    st.param.v2.b8 [func_retval0+4], {%rs1, %rs1};
286; CHECK-NEXT:    st.param.b8 [func_retval0+6], %rs1;
287; CHECK-NEXT:    ret;
288  ret <7 x i8> zeroinitializer
289}
290
291define <5 x i8> @byte5() {
292; CHECK-LABEL: byte5(
293; CHECK:       {
294; CHECK-NEXT:    .reg .b16 %rs<2>;
295; CHECK-EMPTY:
296; CHECK-NEXT:  // %bb.0:
297; CHECK-NEXT:    mov.b16 %rs1, 0;
298; CHECK-NEXT:    st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1};
299; CHECK-NEXT:    st.param.b8 [func_retval0+4], %rs1;
300; CHECK-NEXT:    ret;
301  ret <5 x i8> zeroinitializer
302}
303
304define <4 x i8> @byte4() {
305; CHECK-LABEL: byte4(
306; CHECK:       {
307; CHECK-NEXT:    .reg .b32 %r<2>;
308; CHECK-EMPTY:
309; CHECK-NEXT:  // %bb.0:
310; CHECK-NEXT:    mov.b32 %r1, 0;
311; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
312; CHECK-NEXT:    ret;
313  ret <4 x i8> zeroinitializer
314}
315
316define <3 x i8> @byte3() {
317; CHECK-LABEL: byte3(
318; CHECK:       {
319; CHECK-NEXT:    .reg .b32 %r<2>;
320; CHECK-EMPTY:
321; CHECK-NEXT:  // %bb.0:
322; CHECK-NEXT:    mov.b32 %r1, 0;
323; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
324; CHECK-NEXT:    ret;
325  ret <3 x i8> zeroinitializer
326}
327
328define <2 x i8> @byte2() {
329; CHECK-LABEL: byte2(
330; CHECK:       {
331; CHECK-NEXT:    .reg .b32 %r<2>;
332; CHECK-EMPTY:
333; CHECK-NEXT:  // %bb.0:
334; CHECK-NEXT:    mov.b32 %r1, 0;
335; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
336; CHECK-NEXT:    ret;
337  ret <2 x i8> zeroinitializer
338}
339
340define <1 x i8> @byte1() {
341; CHECK-LABEL: byte1(
342; CHECK:       {
343; CHECK-NEXT:    .reg .b16 %rs<2>;
344; CHECK-EMPTY:
345; CHECK-NEXT:  // %bb.0:
346; CHECK-NEXT:    mov.b16 %rs1, 0;
347; CHECK-NEXT:    st.param.b8 [func_retval0], %rs1;
348; CHECK-NEXT:    ret;
349  ret <1 x i8> zeroinitializer
350}
351
352define <17 x i1> @bit17() {
353; CHECK-LABEL: bit17(
354; CHECK:       {
355; CHECK-NEXT:    .reg .b16 %rs<2>;
356; CHECK-EMPTY:
357; CHECK-NEXT:  // %bb.0:
358; CHECK-NEXT:    mov.b16 %rs1, 0;
359; CHECK-NEXT:    st.param.v4.b8 [func_retval0], {%rs1, %rs1, %rs1, %rs1};
360; CHECK-NEXT:    st.param.v4.b8 [func_retval0+4], {%rs1, %rs1, %rs1, %rs1};
361; CHECK-NEXT:    st.param.v4.b8 [func_retval0+8], {%rs1, %rs1, %rs1, %rs1};
362; CHECK-NEXT:    st.param.v4.b8 [func_retval0+12], {%rs1, %rs1, %rs1, %rs1};
363; CHECK-NEXT:    st.param.b8 [func_retval0+16], %rs1;
364; CHECK-NEXT:    ret;
365  ret <17 x i1> zeroinitializer
366}
367
368define <16 x i1> @bit16() {
369; CHECK-LABEL: bit16(
370; CHECK:       {
371; CHECK-NEXT:    .reg .b16 %rs<2>;
372; CHECK-EMPTY:
373; CHECK-NEXT:  // %bb.0:
374; CHECK-NEXT:    mov.b16 %rs1, 0;
375; CHECK-NEXT:    st.param.v2.b8 [func_retval0], {%rs1, %rs1};
376; CHECK-NEXT:    st.param.v2.b8 [func_retval0+2], {%rs1, %rs1};
377; CHECK-NEXT:    st.param.v2.b8 [func_retval0+4], {%rs1, %rs1};
378; CHECK-NEXT:    st.param.v2.b8 [func_retval0+6], {%rs1, %rs1};
379; CHECK-NEXT:    st.param.v2.b8 [func_retval0+8], {%rs1, %rs1};
380; CHECK-NEXT:    st.param.v2.b8 [func_retval0+10], {%rs1, %rs1};
381; CHECK-NEXT:    st.param.v2.b8 [func_retval0+12], {%rs1, %rs1};
382; CHECK-NEXT:    st.param.v2.b8 [func_retval0+14], {%rs1, %rs1};
383; CHECK-NEXT:    ret;
384  ret <16 x i1> zeroinitializer
385}
386
387define <15 x i1> @bit15() {
388; CHECK-LABEL: bit15(
389; CHECK:       {
390; CHECK-NEXT:    .reg .b16 %rs<2>;
391; CHECK-EMPTY:
392; CHECK-NEXT:  // %bb.0:
393; CHECK-NEXT:    mov.b16 %rs1, 0;
394; CHECK-NEXT:    st.param.v2.b8 [func_retval0], {%rs1, %rs1};
395; CHECK-NEXT:    st.param.v2.b8 [func_retval0+2], {%rs1, %rs1};
396; CHECK-NEXT:    st.param.v2.b8 [func_retval0+4], {%rs1, %rs1};
397; CHECK-NEXT:    st.param.v2.b8 [func_retval0+6], {%rs1, %rs1};
398; CHECK-NEXT:    st.param.v2.b8 [func_retval0+8], {%rs1, %rs1};
399; CHECK-NEXT:    st.param.v2.b8 [func_retval0+10], {%rs1, %rs1};
400; CHECK-NEXT:    st.param.v2.b8 [func_retval0+12], {%rs1, %rs1};
401; CHECK-NEXT:    st.param.b8 [func_retval0+14], %rs1;
402; CHECK-NEXT:    ret;
403  ret <15 x i1> zeroinitializer
404}
405
406define <9 x i1> @bit9() {
407; CHECK-LABEL: bit9(
408; CHECK:       {
409; CHECK-NEXT:    .reg .b16 %rs<2>;
410; CHECK-EMPTY:
411; CHECK-NEXT:  // %bb.0:
412; CHECK-NEXT:    mov.b16 %rs1, 0;
413; CHECK-NEXT:    st.param.v2.b8 [func_retval0], {%rs1, %rs1};
414; CHECK-NEXT:    st.param.v2.b8 [func_retval0+2], {%rs1, %rs1};
415; CHECK-NEXT:    st.param.v2.b8 [func_retval0+4], {%rs1, %rs1};
416; CHECK-NEXT:    st.param.v2.b8 [func_retval0+6], {%rs1, %rs1};
417; CHECK-NEXT:    st.param.b8 [func_retval0+8], %rs1;
418; CHECK-NEXT:    ret;
419  ret <9 x i1> zeroinitializer
420}
421
422define <8 x i1> @bit8() {
423; CHECK-LABEL: bit8(
424; CHECK:       {
425; CHECK-NEXT:    .reg .b16 %rs<2>;
426; CHECK-EMPTY:
427; CHECK-NEXT:  // %bb.0:
428; CHECK-NEXT:    mov.b16 %rs1, 0;
429; CHECK-NEXT:    st.param.b8 [func_retval0], %rs1;
430; CHECK-NEXT:    st.param.b8 [func_retval0+1], %rs1;
431; CHECK-NEXT:    st.param.b8 [func_retval0+2], %rs1;
432; CHECK-NEXT:    st.param.b8 [func_retval0+3], %rs1;
433; CHECK-NEXT:    st.param.b8 [func_retval0+4], %rs1;
434; CHECK-NEXT:    st.param.b8 [func_retval0+5], %rs1;
435; CHECK-NEXT:    st.param.b8 [func_retval0+6], %rs1;
436; CHECK-NEXT:    st.param.b8 [func_retval0+7], %rs1;
437; CHECK-NEXT:    ret;
438  ret <8 x i1> zeroinitializer
439}
440
441define <7 x i1> @bit7() {
442; CHECK-LABEL: bit7(
443; CHECK:       {
444; CHECK-NEXT:    .reg .b16 %rs<2>;
445; CHECK-EMPTY:
446; CHECK-NEXT:  // %bb.0:
447; CHECK-NEXT:    mov.b16 %rs1, 0;
448; CHECK-NEXT:    st.param.b8 [func_retval0], %rs1;
449; CHECK-NEXT:    st.param.b8 [func_retval0+1], %rs1;
450; CHECK-NEXT:    st.param.b8 [func_retval0+2], %rs1;
451; CHECK-NEXT:    st.param.b8 [func_retval0+3], %rs1;
452; CHECK-NEXT:    st.param.b8 [func_retval0+4], %rs1;
453; CHECK-NEXT:    st.param.b8 [func_retval0+5], %rs1;
454; CHECK-NEXT:    st.param.b8 [func_retval0+6], %rs1;
455; CHECK-NEXT:    ret;
456  ret <7 x i1> zeroinitializer
457}
458
459define <5 x i1> @bit5() {
460; CHECK-LABEL: bit5(
461; CHECK:       {
462; CHECK-NEXT:    .reg .b16 %rs<2>;
463; CHECK-EMPTY:
464; CHECK-NEXT:  // %bb.0:
465; CHECK-NEXT:    mov.b16 %rs1, 0;
466; CHECK-NEXT:    st.param.b8 [func_retval0], %rs1;
467; CHECK-NEXT:    st.param.b8 [func_retval0+1], %rs1;
468; CHECK-NEXT:    st.param.b8 [func_retval0+2], %rs1;
469; CHECK-NEXT:    st.param.b8 [func_retval0+3], %rs1;
470; CHECK-NEXT:    st.param.b8 [func_retval0+4], %rs1;
471; CHECK-NEXT:    ret;
472  ret <5 x i1> zeroinitializer
473}
474
475define <4 x i1> @bit4() {
476; CHECK-LABEL: bit4(
477; CHECK:       {
478; CHECK-NEXT:    .reg .b16 %rs<2>;
479; CHECK-EMPTY:
480; CHECK-NEXT:  // %bb.0:
481; CHECK-NEXT:    mov.b16 %rs1, 0;
482; CHECK-NEXT:    st.param.b8 [func_retval0], %rs1;
483; CHECK-NEXT:    st.param.b8 [func_retval0+1], %rs1;
484; CHECK-NEXT:    st.param.b8 [func_retval0+2], %rs1;
485; CHECK-NEXT:    st.param.b8 [func_retval0+3], %rs1;
486; CHECK-NEXT:    ret;
487  ret <4 x i1> zeroinitializer
488}
489
490define <3 x i1> @bit3() {
491; CHECK-LABEL: bit3(
492; CHECK:       {
493; CHECK-NEXT:    .reg .b16 %rs<2>;
494; CHECK-EMPTY:
495; CHECK-NEXT:  // %bb.0:
496; CHECK-NEXT:    mov.b16 %rs1, 0;
497; CHECK-NEXT:    st.param.b8 [func_retval0], %rs1;
498; CHECK-NEXT:    st.param.b8 [func_retval0+1], %rs1;
499; CHECK-NEXT:    st.param.b8 [func_retval0+2], %rs1;
500; CHECK-NEXT:    ret;
501  ret <3 x i1> zeroinitializer
502}
503
504define <2 x i1> @bit2() {
505; CHECK-LABEL: bit2(
506; CHECK:       {
507; CHECK-NEXT:    .reg .b16 %rs<2>;
508; CHECK-EMPTY:
509; CHECK-NEXT:  // %bb.0:
510; CHECK-NEXT:    mov.b16 %rs1, 0;
511; CHECK-NEXT:    st.param.b8 [func_retval0], %rs1;
512; CHECK-NEXT:    st.param.b8 [func_retval0+1], %rs1;
513; CHECK-NEXT:    ret;
514  ret <2 x i1> zeroinitializer
515}
516
517define <1 x i1> @bit1() {
518; CHECK-LABEL: bit1(
519; CHECK:       {
520; CHECK-NEXT:    .reg .b16 %rs<2>;
521; CHECK-EMPTY:
522; CHECK-NEXT:  // %bb.0:
523; CHECK-NEXT:    mov.b16 %rs1, 0;
524; CHECK-NEXT:    st.param.b8 [func_retval0], %rs1;
525; CHECK-NEXT:    ret;
526  ret <1 x i1> zeroinitializer
527}
528