xref: /llvm-project/llvm/test/CodeGen/NVPTX/st-param-imm.ll (revision b279f6b098d3849f7f1c1f539b108307d5f8ae2d)
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
2; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s
3; RUN: llc < %s -mtriple=nvptx | FileCheck %s
4; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -mtriple=nvptx -verify-machineinstrs | %ptxas-verify %}
5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -verify-machineinstrs | %ptxas-verify %}
6
7target triple = "nvptx64-nvidia-cuda"
8
9%struct.A = type { i8, i16 }
10%struct.char2 = type { i8, i8 }
11%struct.char4 = type { i8, i8, i8, i8 }
12%struct.short2 = type { i16, i16 }
13%struct.short4 = type { i16, i16, i16, i16 }
14%struct.int2 = type { i32, i32 }
15%struct.int4 = type { i32, i32, i32, i32 }
16%struct.longlong2 = type { i64, i64 }
17%struct.float2 = type { float, float }
18%struct.float4 = type { float, float, float, float }
19%struct.double2 = type { double, double }
20
21define void @st_param_i8_i16() {
22; CHECK-LABEL: st_param_i8_i16(
23; CHECK:       {
24; CHECK-EMPTY:
25; CHECK-EMPTY:
26; CHECK-NEXT:  // %bb.0:
27; CHECK-NEXT:    { // callseq 0, 0
28; CHECK-NEXT:    .param .align 2 .b8 param0[4];
29; CHECK-NEXT:    st.param.b8 [param0], 1;
30; CHECK-NEXT:    st.param.b16 [param0+2], 2;
31; CHECK-NEXT:    call.uni
32; CHECK-NEXT:    call_i8_i16,
33; CHECK-NEXT:    (
34; CHECK-NEXT:    param0
35; CHECK-NEXT:    );
36; CHECK-NEXT:    } // callseq 0
37; CHECK-NEXT:    ret;
38  call void @call_i8_i16(%struct.A { i8 1, i16 2 })
39  ret void
40}
41
42define void @st_param_i32() {
43; CHECK-LABEL: st_param_i32(
44; CHECK:       {
45; CHECK-EMPTY:
46; CHECK-EMPTY:
47; CHECK-NEXT:  // %bb.0:
48; CHECK-NEXT:    { // callseq 1, 0
49; CHECK-NEXT:    .param .b32 param0;
50; CHECK-NEXT:    st.param.b32 [param0], 3;
51; CHECK-NEXT:    call.uni
52; CHECK-NEXT:    call_i32,
53; CHECK-NEXT:    (
54; CHECK-NEXT:    param0
55; CHECK-NEXT:    );
56; CHECK-NEXT:    } // callseq 1
57; CHECK-NEXT:    ret;
58  call void @call_i32(i32 3)
59  ret void
60}
61
62define void @st_param_i64() {
63; CHECK-LABEL: st_param_i64(
64; CHECK:       {
65; CHECK-EMPTY:
66; CHECK-EMPTY:
67; CHECK-NEXT:  // %bb.0:
68; CHECK-NEXT:    { // callseq 2, 0
69; CHECK-NEXT:    .param .b64 param0;
70; CHECK-NEXT:    st.param.b64 [param0], 4;
71; CHECK-NEXT:    call.uni
72; CHECK-NEXT:    call_i64,
73; CHECK-NEXT:    (
74; CHECK-NEXT:    param0
75; CHECK-NEXT:    );
76; CHECK-NEXT:    } // callseq 2
77; CHECK-NEXT:    ret;
78  call void @call_i64(i64 4)
79  ret void
80}
81
82define void @st_param_f32() {
83; CHECK-LABEL: st_param_f32(
84; CHECK:       {
85; CHECK-EMPTY:
86; CHECK-EMPTY:
87; CHECK-NEXT:  // %bb.0:
88; CHECK-NEXT:    { // callseq 3, 0
89; CHECK-NEXT:    .param .b32 param0;
90; CHECK-NEXT:    st.param.f32 [param0], 0f40A00000;
91; CHECK-NEXT:    call.uni
92; CHECK-NEXT:    call_f32,
93; CHECK-NEXT:    (
94; CHECK-NEXT:    param0
95; CHECK-NEXT:    );
96; CHECK-NEXT:    } // callseq 3
97; CHECK-NEXT:    ret;
98  call void @call_f32(float 5.0)
99  ret void
100}
101
102define void @st_param_f64() {
103; CHECK-LABEL: st_param_f64(
104; CHECK:       {
105; CHECK-EMPTY:
106; CHECK-EMPTY:
107; CHECK-NEXT:  // %bb.0:
108; CHECK-NEXT:    { // callseq 4, 0
109; CHECK-NEXT:    .param .b64 param0;
110; CHECK-NEXT:    st.param.f64 [param0], 0d4018000000000000;
111; CHECK-NEXT:    call.uni
112; CHECK-NEXT:    call_f64,
113; CHECK-NEXT:    (
114; CHECK-NEXT:    param0
115; CHECK-NEXT:    );
116; CHECK-NEXT:    } // callseq 4
117; CHECK-NEXT:    ret;
118  call void @call_f64(double 6.0)
119  ret void
120}
121
122declare void @call_i8_i16(%struct.A)
123declare void @call_i32(i32)
124declare void @call_i64(i64)
125declare void @call_f32(float)
126declare void @call_f64(double)
127
128define void @st_param_v2_i8_ii() {
129; CHECK-LABEL: st_param_v2_i8_ii(
130; CHECK:       {
131; CHECK-EMPTY:
132; CHECK-EMPTY:
133; CHECK-NEXT:  // %bb.0:
134; CHECK-NEXT:    { // callseq 5, 0
135; CHECK-NEXT:    .param .align 2 .b8 param0[2];
136; CHECK-NEXT:    st.param.v2.b8 [param0], {1, 2};
137; CHECK-NEXT:    call.uni
138; CHECK-NEXT:    call_v2_i8,
139; CHECK-NEXT:    (
140; CHECK-NEXT:    param0
141; CHECK-NEXT:    );
142; CHECK-NEXT:    } // callseq 5
143; CHECK-NEXT:    ret;
144  call void @call_v2_i8(%struct.char2 { i8 1, i8 2 })
145  ret void
146}
147define void @st_param_v2_i8_ir(i8 %val) {
148; CHECK-LABEL: st_param_v2_i8_ir(
149; CHECK:       {
150; CHECK-NEXT:    .reg .b16 %rs<2>;
151; CHECK-EMPTY:
152; CHECK-NEXT:  // %bb.0:
153; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v2_i8_ir_param_0];
154; CHECK-NEXT:    { // callseq 6, 0
155; CHECK-NEXT:    .param .align 2 .b8 param0[2];
156; CHECK-NEXT:    st.param.v2.b8 [param0], {1, %rs1};
157; CHECK-NEXT:    call.uni
158; CHECK-NEXT:    call_v2_i8,
159; CHECK-NEXT:    (
160; CHECK-NEXT:    param0
161; CHECK-NEXT:    );
162; CHECK-NEXT:    } // callseq 6
163; CHECK-NEXT:    ret;
164  %struct.ir0 = insertvalue %struct.char2 poison, i8 1, 0
165  %struct.ir1 = insertvalue %struct.char2 %struct.ir0, i8 %val, 1
166  call void @call_v2_i8(%struct.char2 %struct.ir1)
167  ret void
168}
169define void @st_param_v2_i8_ri(i8 %val) {
170; CHECK-LABEL: st_param_v2_i8_ri(
171; CHECK:       {
172; CHECK-NEXT:    .reg .b16 %rs<2>;
173; CHECK-EMPTY:
174; CHECK-NEXT:  // %bb.0:
175; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v2_i8_ri_param_0];
176; CHECK-NEXT:    { // callseq 7, 0
177; CHECK-NEXT:    .param .align 2 .b8 param0[2];
178; CHECK-NEXT:    st.param.v2.b8 [param0], {%rs1, 2};
179; CHECK-NEXT:    call.uni
180; CHECK-NEXT:    call_v2_i8,
181; CHECK-NEXT:    (
182; CHECK-NEXT:    param0
183; CHECK-NEXT:    );
184; CHECK-NEXT:    } // callseq 7
185; CHECK-NEXT:    ret;
186  %struct.ri0 = insertvalue %struct.char2 poison, i8 %val, 0
187  %struct.ri1 = insertvalue %struct.char2 %struct.ri0, i8 2, 1
188  call void @call_v2_i8(%struct.char2 %struct.ri1)
189  ret void
190}
191
192define void @st_param_v2_i16_ii() {
193; CHECK-LABEL: st_param_v2_i16_ii(
194; CHECK:       {
195; CHECK-EMPTY:
196; CHECK-EMPTY:
197; CHECK-NEXT:  // %bb.0:
198; CHECK-NEXT:    { // callseq 8, 0
199; CHECK-NEXT:    .param .align 4 .b8 param0[4];
200; CHECK-NEXT:    st.param.v2.b16 [param0], {1, 2};
201; CHECK-NEXT:    call.uni
202; CHECK-NEXT:    call_v2_i16,
203; CHECK-NEXT:    (
204; CHECK-NEXT:    param0
205; CHECK-NEXT:    );
206; CHECK-NEXT:    } // callseq 8
207; CHECK-NEXT:    ret;
208  call void @call_v2_i16(%struct.short2 { i16 1, i16 2 })
209  ret void
210}
211define void @st_param_v2_i16_ir(i16 %val) {
212; CHECK-LABEL: st_param_v2_i16_ir(
213; CHECK:       {
214; CHECK-NEXT:    .reg .b16 %rs<2>;
215; CHECK-EMPTY:
216; CHECK-NEXT:  // %bb.0:
217; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v2_i16_ir_param_0];
218; CHECK-NEXT:    { // callseq 9, 0
219; CHECK-NEXT:    .param .align 4 .b8 param0[4];
220; CHECK-NEXT:    st.param.v2.b16 [param0], {1, %rs1};
221; CHECK-NEXT:    call.uni
222; CHECK-NEXT:    call_v2_i16,
223; CHECK-NEXT:    (
224; CHECK-NEXT:    param0
225; CHECK-NEXT:    );
226; CHECK-NEXT:    } // callseq 9
227; CHECK-NEXT:    ret;
228  %struct.ir0 = insertvalue %struct.short2 poison, i16 1, 0
229  %struct.ir1 = insertvalue %struct.short2 %struct.ir0, i16 %val, 1
230  call void @call_v2_i16(%struct.short2 %struct.ir1)
231  ret void
232}
233define void @st_param_v2_i16_ri(i16 %val) {
234; CHECK-LABEL: st_param_v2_i16_ri(
235; CHECK:       {
236; CHECK-NEXT:    .reg .b16 %rs<2>;
237; CHECK-EMPTY:
238; CHECK-NEXT:  // %bb.0:
239; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v2_i16_ri_param_0];
240; CHECK-NEXT:    { // callseq 10, 0
241; CHECK-NEXT:    .param .align 4 .b8 param0[4];
242; CHECK-NEXT:    st.param.v2.b16 [param0], {%rs1, 2};
243; CHECK-NEXT:    call.uni
244; CHECK-NEXT:    call_v2_i16,
245; CHECK-NEXT:    (
246; CHECK-NEXT:    param0
247; CHECK-NEXT:    );
248; CHECK-NEXT:    } // callseq 10
249; CHECK-NEXT:    ret;
250  %struct.ri0 = insertvalue %struct.short2 poison, i16 %val, 0
251  %struct.ri1 = insertvalue %struct.short2 %struct.ri0, i16 2, 1
252  call void @call_v2_i16(%struct.short2 %struct.ri1)
253  ret void
254}
255
256define void @st_param_v2_i32_ii() {
257; CHECK-LABEL: st_param_v2_i32_ii(
258; CHECK:       {
259; CHECK-EMPTY:
260; CHECK-EMPTY:
261; CHECK-NEXT:  // %bb.0:
262; CHECK-NEXT:    { // callseq 11, 0
263; CHECK-NEXT:    .param .align 8 .b8 param0[8];
264; CHECK-NEXT:    st.param.v2.b32 [param0], {1, 2};
265; CHECK-NEXT:    call.uni
266; CHECK-NEXT:    call_v2_i32,
267; CHECK-NEXT:    (
268; CHECK-NEXT:    param0
269; CHECK-NEXT:    );
270; CHECK-NEXT:    } // callseq 11
271; CHECK-NEXT:    ret;
272  call void @call_v2_i32(%struct.int2 { i32 1, i32 2 })
273  ret void
274}
275define void @st_param_v2_i32_ir(i32 %val) {
276; CHECK-LABEL: st_param_v2_i32_ir(
277; CHECK:       {
278; CHECK-NEXT:    .reg .b32 %r<2>;
279; CHECK-EMPTY:
280; CHECK-NEXT:  // %bb.0:
281; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v2_i32_ir_param_0];
282; CHECK-NEXT:    { // callseq 12, 0
283; CHECK-NEXT:    .param .align 8 .b8 param0[8];
284; CHECK-NEXT:    st.param.v2.b32 [param0], {1, %r1};
285; CHECK-NEXT:    call.uni
286; CHECK-NEXT:    call_v2_i32,
287; CHECK-NEXT:    (
288; CHECK-NEXT:    param0
289; CHECK-NEXT:    );
290; CHECK-NEXT:    } // callseq 12
291; CHECK-NEXT:    ret;
292  %struct.ir0 = insertvalue %struct.int2 poison, i32 1, 0
293  %struct.ir1 = insertvalue %struct.int2 %struct.ir0, i32 %val, 1
294  call void @call_v2_i32(%struct.int2 %struct.ir1)
295  ret void
296}
297define void @st_param_v2_i32_ri(i32 %val) {
298; CHECK-LABEL: st_param_v2_i32_ri(
299; CHECK:       {
300; CHECK-NEXT:    .reg .b32 %r<2>;
301; CHECK-EMPTY:
302; CHECK-NEXT:  // %bb.0:
303; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v2_i32_ri_param_0];
304; CHECK-NEXT:    { // callseq 13, 0
305; CHECK-NEXT:    .param .align 8 .b8 param0[8];
306; CHECK-NEXT:    st.param.v2.b32 [param0], {%r1, 2};
307; CHECK-NEXT:    call.uni
308; CHECK-NEXT:    call_v2_i32,
309; CHECK-NEXT:    (
310; CHECK-NEXT:    param0
311; CHECK-NEXT:    );
312; CHECK-NEXT:    } // callseq 13
313; CHECK-NEXT:    ret;
314  %struct.ri0 = insertvalue %struct.int2 poison, i32 %val, 0
315  %struct.ri1 = insertvalue %struct.int2 %struct.ri0, i32 2, 1
316  call void @call_v2_i32(%struct.int2 %struct.ri1)
317  ret void
318}
319
320define void @st_param_v2_i64_ii() {
321; CHECK-LABEL: st_param_v2_i64_ii(
322; CHECK:       {
323; CHECK-EMPTY:
324; CHECK-EMPTY:
325; CHECK-NEXT:  // %bb.0:
326; CHECK-NEXT:    { // callseq 14, 0
327; CHECK-NEXT:    .param .align 16 .b8 param0[16];
328; CHECK-NEXT:    st.param.v2.b64 [param0], {1, 2};
329; CHECK-NEXT:    call.uni
330; CHECK-NEXT:    call_v2_i64,
331; CHECK-NEXT:    (
332; CHECK-NEXT:    param0
333; CHECK-NEXT:    );
334; CHECK-NEXT:    } // callseq 14
335; CHECK-NEXT:    ret;
336  call void @call_v2_i64(%struct.longlong2 { i64 1, i64 2 })
337  ret void
338}
339define void @st_param_v2_i64_ir(i64 %val) {
340; CHECK-LABEL: st_param_v2_i64_ir(
341; CHECK:       {
342; CHECK-NEXT:    .reg .b64 %rd<2>;
343; CHECK-EMPTY:
344; CHECK-NEXT:  // %bb.0:
345; CHECK-NEXT:    ld.param.u64 %rd1, [st_param_v2_i64_ir_param_0];
346; CHECK-NEXT:    { // callseq 15, 0
347; CHECK-NEXT:    .param .align 16 .b8 param0[16];
348; CHECK-NEXT:    st.param.v2.b64 [param0], {1, %rd1};
349; CHECK-NEXT:    call.uni
350; CHECK-NEXT:    call_v2_i64,
351; CHECK-NEXT:    (
352; CHECK-NEXT:    param0
353; CHECK-NEXT:    );
354; CHECK-NEXT:    } // callseq 15
355; CHECK-NEXT:    ret;
356  %struct.ir0 = insertvalue %struct.longlong2 poison, i64 1, 0
357  %struct.ir1 = insertvalue %struct.longlong2 %struct.ir0, i64 %val, 1
358  call void @call_v2_i64(%struct.longlong2 %struct.ir1)
359  ret void
360}
361define void @st_param_v2_i64_ri(i64 %val) {
362; CHECK-LABEL: st_param_v2_i64_ri(
363; CHECK:       {
364; CHECK-NEXT:    .reg .b64 %rd<2>;
365; CHECK-EMPTY:
366; CHECK-NEXT:  // %bb.0:
367; CHECK-NEXT:    ld.param.u64 %rd1, [st_param_v2_i64_ri_param_0];
368; CHECK-NEXT:    { // callseq 16, 0
369; CHECK-NEXT:    .param .align 16 .b8 param0[16];
370; CHECK-NEXT:    st.param.v2.b64 [param0], {%rd1, 2};
371; CHECK-NEXT:    call.uni
372; CHECK-NEXT:    call_v2_i64,
373; CHECK-NEXT:    (
374; CHECK-NEXT:    param0
375; CHECK-NEXT:    );
376; CHECK-NEXT:    } // callseq 16
377; CHECK-NEXT:    ret;
378  %struct.ri0 = insertvalue %struct.longlong2 poison, i64 %val, 0
379  %struct.ri1 = insertvalue %struct.longlong2 %struct.ri0, i64 2, 1
380  call void @call_v2_i64(%struct.longlong2 %struct.ri1)
381  ret void
382}
383
384define void @st_param_v2_f32_ii(float %val) {
385; CHECK-LABEL: st_param_v2_f32_ii(
386; CHECK:       {
387; CHECK-EMPTY:
388; CHECK-EMPTY:
389; CHECK-NEXT:  // %bb.0:
390; CHECK-NEXT:    { // callseq 17, 0
391; CHECK-NEXT:    .param .align 8 .b8 param0[8];
392; CHECK-NEXT:    st.param.v2.f32 [param0], {0f3F800000, 0f40000000};
393; CHECK-NEXT:    call.uni
394; CHECK-NEXT:    call_v2_f32,
395; CHECK-NEXT:    (
396; CHECK-NEXT:    param0
397; CHECK-NEXT:    );
398; CHECK-NEXT:    } // callseq 17
399; CHECK-NEXT:    ret;
400  call void @call_v2_f32(%struct.float2 { float 1.0, float 2.0 })
401  ret void
402}
403define void @st_param_v2_f32_ir(float %val) {
404; CHECK-LABEL: st_param_v2_f32_ir(
405; CHECK:       {
406; CHECK-NEXT:    .reg .f32 %f<2>;
407; CHECK-EMPTY:
408; CHECK-NEXT:  // %bb.0:
409; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v2_f32_ir_param_0];
410; CHECK-NEXT:    { // callseq 18, 0
411; CHECK-NEXT:    .param .align 8 .b8 param0[8];
412; CHECK-NEXT:    st.param.v2.f32 [param0], {0f3F800000, %f1};
413; CHECK-NEXT:    call.uni
414; CHECK-NEXT:    call_v2_f32,
415; CHECK-NEXT:    (
416; CHECK-NEXT:    param0
417; CHECK-NEXT:    );
418; CHECK-NEXT:    } // callseq 18
419; CHECK-NEXT:    ret;
420  %struct.ir0 = insertvalue %struct.float2 poison, float 1.0, 0
421  %struct.ir1 = insertvalue %struct.float2 %struct.ir0, float %val, 1
422  call void @call_v2_f32(%struct.float2 %struct.ir1)
423  ret void
424}
425define void @st_param_v2_f32_ri(float %val) {
426; CHECK-LABEL: st_param_v2_f32_ri(
427; CHECK:       {
428; CHECK-NEXT:    .reg .f32 %f<2>;
429; CHECK-EMPTY:
430; CHECK-NEXT:  // %bb.0:
431; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v2_f32_ri_param_0];
432; CHECK-NEXT:    { // callseq 19, 0
433; CHECK-NEXT:    .param .align 8 .b8 param0[8];
434; CHECK-NEXT:    st.param.v2.f32 [param0], {%f1, 0f40000000};
435; CHECK-NEXT:    call.uni
436; CHECK-NEXT:    call_v2_f32,
437; CHECK-NEXT:    (
438; CHECK-NEXT:    param0
439; CHECK-NEXT:    );
440; CHECK-NEXT:    } // callseq 19
441; CHECK-NEXT:    ret;
442  %struct.ri0 = insertvalue %struct.float2 poison, float %val, 0
443  %struct.ri1 = insertvalue %struct.float2 %struct.ri0, float 2.0, 1
444  call void @call_v2_f32(%struct.float2 %struct.ri1)
445  ret void
446}
447
448define void @st_param_v2_f64_ii(double %val) {
449; CHECK-LABEL: st_param_v2_f64_ii(
450; CHECK:       {
451; CHECK-EMPTY:
452; CHECK-EMPTY:
453; CHECK-NEXT:  // %bb.0:
454; CHECK-NEXT:    { // callseq 20, 0
455; CHECK-NEXT:    .param .align 16 .b8 param0[16];
456; CHECK-NEXT:    st.param.v2.f64 [param0], {0d3FF0000000000000, 0d4000000000000000};
457; CHECK-NEXT:    call.uni
458; CHECK-NEXT:    call_v2_f64,
459; CHECK-NEXT:    (
460; CHECK-NEXT:    param0
461; CHECK-NEXT:    );
462; CHECK-NEXT:    } // callseq 20
463; CHECK-NEXT:    ret;
464  call void @call_v2_f64(%struct.double2 { double 1.0, double 2.0 })
465  ret void
466}
467define void @st_param_v2_f64_ir(double %val) {
468; CHECK-LABEL: st_param_v2_f64_ir(
469; CHECK:       {
470; CHECK-NEXT:    .reg .f64 %fd<2>;
471; CHECK-EMPTY:
472; CHECK-NEXT:  // %bb.0:
473; CHECK-NEXT:    ld.param.f64 %fd1, [st_param_v2_f64_ir_param_0];
474; CHECK-NEXT:    { // callseq 21, 0
475; CHECK-NEXT:    .param .align 16 .b8 param0[16];
476; CHECK-NEXT:    st.param.v2.f64 [param0], {0d3FF0000000000000, %fd1};
477; CHECK-NEXT:    call.uni
478; CHECK-NEXT:    call_v2_f64,
479; CHECK-NEXT:    (
480; CHECK-NEXT:    param0
481; CHECK-NEXT:    );
482; CHECK-NEXT:    } // callseq 21
483; CHECK-NEXT:    ret;
484  %struct.ir0 = insertvalue %struct.double2 poison, double 1.0, 0
485  %struct.ir1 = insertvalue %struct.double2 %struct.ir0, double %val, 1
486  call void @call_v2_f64(%struct.double2 %struct.ir1)
487  ret void
488}
489define void @st_param_v2_f64_ri(double %val) {
490; CHECK-LABEL: st_param_v2_f64_ri(
491; CHECK:       {
492; CHECK-NEXT:    .reg .f64 %fd<2>;
493; CHECK-EMPTY:
494; CHECK-NEXT:  // %bb.0:
495; CHECK-NEXT:    ld.param.f64 %fd1, [st_param_v2_f64_ri_param_0];
496; CHECK-NEXT:    { // callseq 22, 0
497; CHECK-NEXT:    .param .align 16 .b8 param0[16];
498; CHECK-NEXT:    st.param.v2.f64 [param0], {%fd1, 0d4000000000000000};
499; CHECK-NEXT:    call.uni
500; CHECK-NEXT:    call_v2_f64,
501; CHECK-NEXT:    (
502; CHECK-NEXT:    param0
503; CHECK-NEXT:    );
504; CHECK-NEXT:    } // callseq 22
505; CHECK-NEXT:    ret;
506  %struct.ri0 = insertvalue %struct.double2 poison, double %val, 0
507  %struct.ri1 = insertvalue %struct.double2 %struct.ri0, double 2.0, 1
508  call void @call_v2_f64(%struct.double2 %struct.ri1)
509  ret void
510}
511
512declare void @call_v2_i8(%struct.char2 alignstack(2))
513declare void @call_v2_i16(%struct.short2 alignstack(4))
514declare void @call_v2_i32(%struct.int2 alignstack(8))
515declare void @call_v2_i64(%struct.longlong2 alignstack(16))
516declare void @call_v2_f32(%struct.float2 alignstack(8))
517declare void @call_v2_f64(%struct.double2 alignstack(16))
518
519define void @st_param_v4_i8_iiii() {
520; CHECK-LABEL: st_param_v4_i8_iiii(
521; CHECK:       {
522; CHECK-EMPTY:
523; CHECK-EMPTY:
524; CHECK-NEXT:  // %bb.0:
525; CHECK-NEXT:    { // callseq 23, 0
526; CHECK-NEXT:    .param .align 4 .b8 param0[4];
527; CHECK-NEXT:    st.param.v4.b8 [param0], {1, 2, 3, 4};
528; CHECK-NEXT:    call.uni
529; CHECK-NEXT:    call_v4_i8,
530; CHECK-NEXT:    (
531; CHECK-NEXT:    param0
532; CHECK-NEXT:    );
533; CHECK-NEXT:    } // callseq 23
534; CHECK-NEXT:    ret;
535  call void @call_v4_i8(%struct.char4 { i8 1, i8 2, i8 3, i8 4 })
536  ret void
537}
538define void @st_param_v4_i8_irrr(i8 %b, i8 %c, i8 %d) {
539; CHECK-LABEL: st_param_v4_i8_irrr(
540; CHECK:       {
541; CHECK-NEXT:    .reg .b16 %rs<4>;
542; CHECK-EMPTY:
543; CHECK-NEXT:  // %bb.0:
544; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_irrr_param_0];
545; CHECK-NEXT:    ld.param.u8 %rs2, [st_param_v4_i8_irrr_param_1];
546; CHECK-NEXT:    ld.param.u8 %rs3, [st_param_v4_i8_irrr_param_2];
547; CHECK-NEXT:    { // callseq 24, 0
548; CHECK-NEXT:    .param .align 4 .b8 param0[4];
549; CHECK-NEXT:    st.param.v4.b8 [param0], {1, %rs1, %rs2, %rs3};
550; CHECK-NEXT:    call.uni
551; CHECK-NEXT:    call_v4_i8,
552; CHECK-NEXT:    (
553; CHECK-NEXT:    param0
554; CHECK-NEXT:    );
555; CHECK-NEXT:    } // callseq 24
556; CHECK-NEXT:    ret;
557  %struct.irrr0 = insertvalue %struct.char4 poison, i8 1, 0
558  %struct.irrr1 = insertvalue %struct.char4 %struct.irrr0, i8 %b, 1
559  %struct.irrr2 = insertvalue %struct.char4 %struct.irrr1, i8 %c, 2
560  %struct.irrr3 = insertvalue %struct.char4 %struct.irrr2, i8 %d, 3
561  call void @call_v4_i8(%struct.char4 %struct.irrr3)
562  ret void
563}
564define void @st_param_v4_i8_rirr(i8 %a, i8 %c, i8 %d) {
565; CHECK-LABEL: st_param_v4_i8_rirr(
566; CHECK:       {
567; CHECK-NEXT:    .reg .b16 %rs<4>;
568; CHECK-EMPTY:
569; CHECK-NEXT:  // %bb.0:
570; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_rirr_param_0];
571; CHECK-NEXT:    ld.param.u8 %rs2, [st_param_v4_i8_rirr_param_1];
572; CHECK-NEXT:    ld.param.u8 %rs3, [st_param_v4_i8_rirr_param_2];
573; CHECK-NEXT:    { // callseq 25, 0
574; CHECK-NEXT:    .param .align 4 .b8 param0[4];
575; CHECK-NEXT:    st.param.v4.b8 [param0], {%rs1, 2, %rs2, %rs3};
576; CHECK-NEXT:    call.uni
577; CHECK-NEXT:    call_v4_i8,
578; CHECK-NEXT:    (
579; CHECK-NEXT:    param0
580; CHECK-NEXT:    );
581; CHECK-NEXT:    } // callseq 25
582; CHECK-NEXT:    ret;
583  %struct.rirr0 = insertvalue %struct.char4 poison, i8 %a, 0
584  %struct.rirr1 = insertvalue %struct.char4 %struct.rirr0, i8 2, 1
585  %struct.rirr2 = insertvalue %struct.char4 %struct.rirr1, i8 %c, 2
586  %struct.rirr3 = insertvalue %struct.char4 %struct.rirr2, i8 %d, 3
587  call void @call_v4_i8(%struct.char4 %struct.rirr3)
588  ret void
589}
590define void @st_param_v4_i8_rrir(i8 %a, i8 %b, i8 %d) {
591; CHECK-LABEL: st_param_v4_i8_rrir(
592; CHECK:       {
593; CHECK-NEXT:    .reg .b16 %rs<4>;
594; CHECK-EMPTY:
595; CHECK-NEXT:  // %bb.0:
596; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_rrir_param_0];
597; CHECK-NEXT:    ld.param.u8 %rs2, [st_param_v4_i8_rrir_param_1];
598; CHECK-NEXT:    ld.param.u8 %rs3, [st_param_v4_i8_rrir_param_2];
599; CHECK-NEXT:    { // callseq 26, 0
600; CHECK-NEXT:    .param .align 4 .b8 param0[4];
601; CHECK-NEXT:    st.param.v4.b8 [param0], {%rs1, %rs2, 3, %rs3};
602; CHECK-NEXT:    call.uni
603; CHECK-NEXT:    call_v4_i8,
604; CHECK-NEXT:    (
605; CHECK-NEXT:    param0
606; CHECK-NEXT:    );
607; CHECK-NEXT:    } // callseq 26
608; CHECK-NEXT:    ret;
609  %struct.rrir0 = insertvalue %struct.char4 poison, i8 %a, 0
610  %struct.rrir1 = insertvalue %struct.char4 %struct.rrir0, i8 %b, 1
611  %struct.rrir2 = insertvalue %struct.char4 %struct.rrir1, i8 3, 2
612  %struct.rrir3 = insertvalue %struct.char4 %struct.rrir2, i8 %d, 3
613  call void @call_v4_i8(%struct.char4 %struct.rrir3)
614  ret void
615}
616define void @st_param_v4_i8_rrri(i8 %a, i8 %b, i8 %c) {
617; CHECK-LABEL: st_param_v4_i8_rrri(
618; CHECK:       {
619; CHECK-NEXT:    .reg .b16 %rs<4>;
620; CHECK-EMPTY:
621; CHECK-NEXT:  // %bb.0:
622; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_rrri_param_0];
623; CHECK-NEXT:    ld.param.u8 %rs2, [st_param_v4_i8_rrri_param_1];
624; CHECK-NEXT:    ld.param.u8 %rs3, [st_param_v4_i8_rrri_param_2];
625; CHECK-NEXT:    { // callseq 27, 0
626; CHECK-NEXT:    .param .align 4 .b8 param0[4];
627; CHECK-NEXT:    st.param.v4.b8 [param0], {%rs1, %rs2, %rs3, 4};
628; CHECK-NEXT:    call.uni
629; CHECK-NEXT:    call_v4_i8,
630; CHECK-NEXT:    (
631; CHECK-NEXT:    param0
632; CHECK-NEXT:    );
633; CHECK-NEXT:    } // callseq 27
634; CHECK-NEXT:    ret;
635  %struct.rrri0 = insertvalue %struct.char4 poison, i8 %a, 0
636  %struct.rrri1 = insertvalue %struct.char4 %struct.rrri0, i8 %b, 1
637  %struct.rrri2 = insertvalue %struct.char4 %struct.rrri1, i8 %c, 2
638  %struct.rrri3 = insertvalue %struct.char4 %struct.rrri2, i8 4, 3
639  call void @call_v4_i8(%struct.char4 %struct.rrri3)
640  ret void
641}
642define void @st_param_v4_i8_iirr(i8 %c, i8 %d) {
643; CHECK-LABEL: st_param_v4_i8_iirr(
644; CHECK:       {
645; CHECK-NEXT:    .reg .b16 %rs<3>;
646; CHECK-EMPTY:
647; CHECK-NEXT:  // %bb.0:
648; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_iirr_param_0];
649; CHECK-NEXT:    ld.param.u8 %rs2, [st_param_v4_i8_iirr_param_1];
650; CHECK-NEXT:    { // callseq 28, 0
651; CHECK-NEXT:    .param .align 4 .b8 param0[4];
652; CHECK-NEXT:    st.param.v4.b8 [param0], {1, 2, %rs1, %rs2};
653; CHECK-NEXT:    call.uni
654; CHECK-NEXT:    call_v4_i8,
655; CHECK-NEXT:    (
656; CHECK-NEXT:    param0
657; CHECK-NEXT:    );
658; CHECK-NEXT:    } // callseq 28
659; CHECK-NEXT:    ret;
660  %struct.iirr0 = insertvalue %struct.char4 poison, i8 1, 0
661  %struct.iirr1 = insertvalue %struct.char4 %struct.iirr0, i8 2, 1
662  %struct.iirr2 = insertvalue %struct.char4 %struct.iirr1, i8 %c, 2
663  %struct.iirr3 = insertvalue %struct.char4 %struct.iirr2, i8 %d, 3
664  call void @call_v4_i8(%struct.char4 %struct.iirr3)
665  ret void
666}
667define void @st_param_v4_i8_irir(i8 %b, i8 %d) {
668; CHECK-LABEL: st_param_v4_i8_irir(
669; CHECK:       {
670; CHECK-NEXT:    .reg .b16 %rs<3>;
671; CHECK-EMPTY:
672; CHECK-NEXT:  // %bb.0:
673; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_irir_param_0];
674; CHECK-NEXT:    ld.param.u8 %rs2, [st_param_v4_i8_irir_param_1];
675; CHECK-NEXT:    { // callseq 29, 0
676; CHECK-NEXT:    .param .align 4 .b8 param0[4];
677; CHECK-NEXT:    st.param.v4.b8 [param0], {1, %rs1, 3, %rs2};
678; CHECK-NEXT:    call.uni
679; CHECK-NEXT:    call_v4_i8,
680; CHECK-NEXT:    (
681; CHECK-NEXT:    param0
682; CHECK-NEXT:    );
683; CHECK-NEXT:    } // callseq 29
684; CHECK-NEXT:    ret;
685  %struct.irir0 = insertvalue %struct.char4 poison, i8 1, 0
686  %struct.irir1 = insertvalue %struct.char4 %struct.irir0, i8 %b, 1
687  %struct.irir2 = insertvalue %struct.char4 %struct.irir1, i8 3, 2
688  %struct.irir3 = insertvalue %struct.char4 %struct.irir2, i8 %d, 3
689  call void @call_v4_i8(%struct.char4 %struct.irir3)
690  ret void
691}
692define void @st_param_v4_i8_irri(i8 %b, i8 %c) {
693; CHECK-LABEL: st_param_v4_i8_irri(
694; CHECK:       {
695; CHECK-NEXT:    .reg .b16 %rs<3>;
696; CHECK-EMPTY:
697; CHECK-NEXT:  // %bb.0:
698; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_irri_param_0];
699; CHECK-NEXT:    ld.param.u8 %rs2, [st_param_v4_i8_irri_param_1];
700; CHECK-NEXT:    { // callseq 30, 0
701; CHECK-NEXT:    .param .align 4 .b8 param0[4];
702; CHECK-NEXT:    st.param.v4.b8 [param0], {1, %rs1, %rs2, 4};
703; CHECK-NEXT:    call.uni
704; CHECK-NEXT:    call_v4_i8,
705; CHECK-NEXT:    (
706; CHECK-NEXT:    param0
707; CHECK-NEXT:    );
708; CHECK-NEXT:    } // callseq 30
709; CHECK-NEXT:    ret;
710  %struct.irri0 = insertvalue %struct.char4 poison, i8 1, 0
711  %struct.irri1 = insertvalue %struct.char4 %struct.irri0, i8 %b, 1
712  %struct.irri2 = insertvalue %struct.char4 %struct.irri1, i8 %c, 2
713  %struct.irri3 = insertvalue %struct.char4 %struct.irri2, i8 4, 3
714  call void @call_v4_i8(%struct.char4 %struct.irri3)
715  ret void
716}
717define void @st_param_v4_i8_riir(i8 %a, i8 %d) {
718; CHECK-LABEL: st_param_v4_i8_riir(
719; CHECK:       {
720; CHECK-NEXT:    .reg .b16 %rs<3>;
721; CHECK-EMPTY:
722; CHECK-NEXT:  // %bb.0:
723; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_riir_param_0];
724; CHECK-NEXT:    ld.param.u8 %rs2, [st_param_v4_i8_riir_param_1];
725; CHECK-NEXT:    { // callseq 31, 0
726; CHECK-NEXT:    .param .align 4 .b8 param0[4];
727; CHECK-NEXT:    st.param.v4.b8 [param0], {%rs1, 2, 3, %rs2};
728; CHECK-NEXT:    call.uni
729; CHECK-NEXT:    call_v4_i8,
730; CHECK-NEXT:    (
731; CHECK-NEXT:    param0
732; CHECK-NEXT:    );
733; CHECK-NEXT:    } // callseq 31
734; CHECK-NEXT:    ret;
735  %struct.riir0 = insertvalue %struct.char4 poison, i8 %a, 0
736  %struct.riir1 = insertvalue %struct.char4 %struct.riir0, i8 2, 1
737  %struct.riir2 = insertvalue %struct.char4 %struct.riir1, i8 3, 2
738  %struct.riir3 = insertvalue %struct.char4 %struct.riir2, i8 %d, 3
739  call void @call_v4_i8(%struct.char4 %struct.riir3)
740  ret void
741}
742define void @st_param_v4_i8_riri(i8 %a, i8 %c) {
743; CHECK-LABEL: st_param_v4_i8_riri(
744; CHECK:       {
745; CHECK-NEXT:    .reg .b16 %rs<3>;
746; CHECK-EMPTY:
747; CHECK-NEXT:  // %bb.0:
748; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_riri_param_0];
749; CHECK-NEXT:    ld.param.u8 %rs2, [st_param_v4_i8_riri_param_1];
750; CHECK-NEXT:    { // callseq 32, 0
751; CHECK-NEXT:    .param .align 4 .b8 param0[4];
752; CHECK-NEXT:    st.param.v4.b8 [param0], {%rs1, 2, %rs2, 4};
753; CHECK-NEXT:    call.uni
754; CHECK-NEXT:    call_v4_i8,
755; CHECK-NEXT:    (
756; CHECK-NEXT:    param0
757; CHECK-NEXT:    );
758; CHECK-NEXT:    } // callseq 32
759; CHECK-NEXT:    ret;
760  %struct.riri0 = insertvalue %struct.char4 poison, i8 %a, 0
761  %struct.riri1 = insertvalue %struct.char4 %struct.riri0, i8 2, 1
762  %struct.riri2 = insertvalue %struct.char4 %struct.riri1, i8 %c, 2
763  %struct.riri3 = insertvalue %struct.char4 %struct.riri2, i8 4, 3
764  call void @call_v4_i8(%struct.char4 %struct.riri3)
765  ret void
766}
767define void @st_param_v4_i8_rrii(i8 %a, i8 %b) {
768; CHECK-LABEL: st_param_v4_i8_rrii(
769; CHECK:       {
770; CHECK-NEXT:    .reg .b16 %rs<3>;
771; CHECK-EMPTY:
772; CHECK-NEXT:  // %bb.0:
773; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_rrii_param_0];
774; CHECK-NEXT:    ld.param.u8 %rs2, [st_param_v4_i8_rrii_param_1];
775; CHECK-NEXT:    { // callseq 33, 0
776; CHECK-NEXT:    .param .align 4 .b8 param0[4];
777; CHECK-NEXT:    st.param.v4.b8 [param0], {%rs1, %rs2, 3, 4};
778; CHECK-NEXT:    call.uni
779; CHECK-NEXT:    call_v4_i8,
780; CHECK-NEXT:    (
781; CHECK-NEXT:    param0
782; CHECK-NEXT:    );
783; CHECK-NEXT:    } // callseq 33
784; CHECK-NEXT:    ret;
785  %struct.rrii0 = insertvalue %struct.char4 poison, i8 %a, 0
786  %struct.rrii1 = insertvalue %struct.char4 %struct.rrii0, i8 %b, 1
787  %struct.rrii2 = insertvalue %struct.char4 %struct.rrii1, i8 3, 2
788  %struct.rrii3 = insertvalue %struct.char4 %struct.rrii2, i8 4, 3
789  call void @call_v4_i8(%struct.char4 %struct.rrii3)
790  ret void
791}
792define void @st_param_v4_i8_iiir(i8 %d) {
793; CHECK-LABEL: st_param_v4_i8_iiir(
794; CHECK:       {
795; CHECK-NEXT:    .reg .b16 %rs<2>;
796; CHECK-EMPTY:
797; CHECK-NEXT:  // %bb.0:
798; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_iiir_param_0];
799; CHECK-NEXT:    { // callseq 34, 0
800; CHECK-NEXT:    .param .align 4 .b8 param0[4];
801; CHECK-NEXT:    st.param.v4.b8 [param0], {1, 2, 3, %rs1};
802; CHECK-NEXT:    call.uni
803; CHECK-NEXT:    call_v4_i8,
804; CHECK-NEXT:    (
805; CHECK-NEXT:    param0
806; CHECK-NEXT:    );
807; CHECK-NEXT:    } // callseq 34
808; CHECK-NEXT:    ret;
809  %struct.iiir0 = insertvalue %struct.char4 poison, i8 1, 0
810  %struct.iiir1 = insertvalue %struct.char4 %struct.iiir0, i8 2, 1
811  %struct.iiir2 = insertvalue %struct.char4 %struct.iiir1, i8 3, 2
812  %struct.iiir3 = insertvalue %struct.char4 %struct.iiir2, i8 %d, 3
813  call void @call_v4_i8(%struct.char4 %struct.iiir3)
814  ret void
815}
816define void @st_param_v4_i8_iiri(i8 %c) {
817; CHECK-LABEL: st_param_v4_i8_iiri(
818; CHECK:       {
819; CHECK-NEXT:    .reg .b16 %rs<2>;
820; CHECK-EMPTY:
821; CHECK-NEXT:  // %bb.0:
822; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_iiri_param_0];
823; CHECK-NEXT:    { // callseq 35, 0
824; CHECK-NEXT:    .param .align 4 .b8 param0[4];
825; CHECK-NEXT:    st.param.v4.b8 [param0], {1, 2, %rs1, 4};
826; CHECK-NEXT:    call.uni
827; CHECK-NEXT:    call_v4_i8,
828; CHECK-NEXT:    (
829; CHECK-NEXT:    param0
830; CHECK-NEXT:    );
831; CHECK-NEXT:    } // callseq 35
832; CHECK-NEXT:    ret;
833  %struct.iiri0 = insertvalue %struct.char4 poison, i8 1, 0
834  %struct.iiri1 = insertvalue %struct.char4 %struct.iiri0, i8 2, 1
835  %struct.iiri2 = insertvalue %struct.char4 %struct.iiri1, i8 %c, 2
836  %struct.iiri3 = insertvalue %struct.char4 %struct.iiri2, i8 4, 3
837  call void @call_v4_i8(%struct.char4 %struct.iiri3)
838  ret void
839}
840define void @st_param_v4_i8_irii(i8 %b) {
841; CHECK-LABEL: st_param_v4_i8_irii(
842; CHECK:       {
843; CHECK-NEXT:    .reg .b16 %rs<2>;
844; CHECK-EMPTY:
845; CHECK-NEXT:  // %bb.0:
846; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_irii_param_0];
847; CHECK-NEXT:    { // callseq 36, 0
848; CHECK-NEXT:    .param .align 4 .b8 param0[4];
849; CHECK-NEXT:    st.param.v4.b8 [param0], {1, %rs1, 3, 4};
850; CHECK-NEXT:    call.uni
851; CHECK-NEXT:    call_v4_i8,
852; CHECK-NEXT:    (
853; CHECK-NEXT:    param0
854; CHECK-NEXT:    );
855; CHECK-NEXT:    } // callseq 36
856; CHECK-NEXT:    ret;
857  %struct.irii0 = insertvalue %struct.char4 poison, i8 1, 0
858  %struct.irii1 = insertvalue %struct.char4 %struct.irii0, i8 %b, 1
859  %struct.irii2 = insertvalue %struct.char4 %struct.irii1, i8 3, 2
860  %struct.irii3 = insertvalue %struct.char4 %struct.irii2, i8 4, 3
861  call void @call_v4_i8(%struct.char4 %struct.irii3)
862  ret void
863}
864define void @st_param_v4_i8_riii(i8 %a) {
865; CHECK-LABEL: st_param_v4_i8_riii(
866; CHECK:       {
867; CHECK-NEXT:    .reg .b16 %rs<2>;
868; CHECK-EMPTY:
869; CHECK-NEXT:  // %bb.0:
870; CHECK-NEXT:    ld.param.u8 %rs1, [st_param_v4_i8_riii_param_0];
871; CHECK-NEXT:    { // callseq 37, 0
872; CHECK-NEXT:    .param .align 4 .b8 param0[4];
873; CHECK-NEXT:    st.param.v4.b8 [param0], {%rs1, 2, 3, 4};
874; CHECK-NEXT:    call.uni
875; CHECK-NEXT:    call_v4_i8,
876; CHECK-NEXT:    (
877; CHECK-NEXT:    param0
878; CHECK-NEXT:    );
879; CHECK-NEXT:    } // callseq 37
880; CHECK-NEXT:    ret;
881  %struct.riii0 = insertvalue %struct.char4 poison, i8 %a, 0
882  %struct.riii1 = insertvalue %struct.char4 %struct.riii0, i8 2, 1
883  %struct.riii2 = insertvalue %struct.char4 %struct.riii1, i8 3, 2
884  %struct.riii3 = insertvalue %struct.char4 %struct.riii2, i8 4, 3
885  call void @call_v4_i8(%struct.char4 %struct.riii3)
886  ret void
887}
888
889define void @st_param_v4_i16_iiii() {
890; CHECK-LABEL: st_param_v4_i16_iiii(
891; CHECK:       {
892; CHECK-EMPTY:
893; CHECK-EMPTY:
894; CHECK-NEXT:  // %bb.0:
895; CHECK-NEXT:    { // callseq 38, 0
896; CHECK-NEXT:    .param .align 8 .b8 param0[8];
897; CHECK-NEXT:    st.param.v4.b16 [param0], {1, 2, 3, 4};
898; CHECK-NEXT:    call.uni
899; CHECK-NEXT:    call_v4_i16,
900; CHECK-NEXT:    (
901; CHECK-NEXT:    param0
902; CHECK-NEXT:    );
903; CHECK-NEXT:    } // callseq 38
904; CHECK-NEXT:    ret;
905  call void @call_v4_i16(%struct.short4 { i16 1, i16 2, i16 3, i16 4 })
906  ret void
907}
908define void @st_param_v4_i16_irrr(i16 %b, i16 %c, i16 %d) {
909; CHECK-LABEL: st_param_v4_i16_irrr(
910; CHECK:       {
911; CHECK-NEXT:    .reg .b16 %rs<4>;
912; CHECK-EMPTY:
913; CHECK-NEXT:  // %bb.0:
914; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_irrr_param_0];
915; CHECK-NEXT:    ld.param.u16 %rs2, [st_param_v4_i16_irrr_param_1];
916; CHECK-NEXT:    ld.param.u16 %rs3, [st_param_v4_i16_irrr_param_2];
917; CHECK-NEXT:    { // callseq 39, 0
918; CHECK-NEXT:    .param .align 8 .b8 param0[8];
919; CHECK-NEXT:    st.param.v4.b16 [param0], {1, %rs1, %rs2, %rs3};
920; CHECK-NEXT:    call.uni
921; CHECK-NEXT:    call_v4_i16,
922; CHECK-NEXT:    (
923; CHECK-NEXT:    param0
924; CHECK-NEXT:    );
925; CHECK-NEXT:    } // callseq 39
926; CHECK-NEXT:    ret;
927  %struct.irrr0 = insertvalue %struct.short4 poison, i16 1, 0
928  %struct.irrr1 = insertvalue %struct.short4 %struct.irrr0, i16 %b, 1
929  %struct.irrr2 = insertvalue %struct.short4 %struct.irrr1, i16 %c, 2
930  %struct.irrr3 = insertvalue %struct.short4 %struct.irrr2, i16 %d, 3
931  call void @call_v4_i16(%struct.short4 %struct.irrr3)
932  ret void
933}
934define void @st_param_v4_i16_rirr(i16 %a, i16 %c, i16 %d) {
935; CHECK-LABEL: st_param_v4_i16_rirr(
936; CHECK:       {
937; CHECK-NEXT:    .reg .b16 %rs<4>;
938; CHECK-EMPTY:
939; CHECK-NEXT:  // %bb.0:
940; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_rirr_param_0];
941; CHECK-NEXT:    ld.param.u16 %rs2, [st_param_v4_i16_rirr_param_1];
942; CHECK-NEXT:    ld.param.u16 %rs3, [st_param_v4_i16_rirr_param_2];
943; CHECK-NEXT:    { // callseq 40, 0
944; CHECK-NEXT:    .param .align 8 .b8 param0[8];
945; CHECK-NEXT:    st.param.v4.b16 [param0], {%rs1, 2, %rs2, %rs3};
946; CHECK-NEXT:    call.uni
947; CHECK-NEXT:    call_v4_i16,
948; CHECK-NEXT:    (
949; CHECK-NEXT:    param0
950; CHECK-NEXT:    );
951; CHECK-NEXT:    } // callseq 40
952; CHECK-NEXT:    ret;
953  %struct.rirr0 = insertvalue %struct.short4 poison, i16 %a, 0
954  %struct.rirr1 = insertvalue %struct.short4 %struct.rirr0, i16 2, 1
955  %struct.rirr2 = insertvalue %struct.short4 %struct.rirr1, i16 %c, 2
956  %struct.rirr3 = insertvalue %struct.short4 %struct.rirr2, i16 %d, 3
957  call void @call_v4_i16(%struct.short4 %struct.rirr3)
958  ret void
959}
960define void @st_param_v4_i16_rrir(i16 %a, i16 %b, i16 %d) {
961; CHECK-LABEL: st_param_v4_i16_rrir(
962; CHECK:       {
963; CHECK-NEXT:    .reg .b16 %rs<4>;
964; CHECK-EMPTY:
965; CHECK-NEXT:  // %bb.0:
966; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_rrir_param_0];
967; CHECK-NEXT:    ld.param.u16 %rs2, [st_param_v4_i16_rrir_param_1];
968; CHECK-NEXT:    ld.param.u16 %rs3, [st_param_v4_i16_rrir_param_2];
969; CHECK-NEXT:    { // callseq 41, 0
970; CHECK-NEXT:    .param .align 8 .b8 param0[8];
971; CHECK-NEXT:    st.param.v4.b16 [param0], {%rs1, %rs2, 3, %rs3};
972; CHECK-NEXT:    call.uni
973; CHECK-NEXT:    call_v4_i16,
974; CHECK-NEXT:    (
975; CHECK-NEXT:    param0
976; CHECK-NEXT:    );
977; CHECK-NEXT:    } // callseq 41
978; CHECK-NEXT:    ret;
979  %struct.rrir0 = insertvalue %struct.short4 poison, i16 %a, 0
980  %struct.rrir1 = insertvalue %struct.short4 %struct.rrir0, i16 %b, 1
981  %struct.rrir2 = insertvalue %struct.short4 %struct.rrir1, i16 3, 2
982  %struct.rrir3 = insertvalue %struct.short4 %struct.rrir2, i16 %d, 3
983  call void @call_v4_i16(%struct.short4 %struct.rrir3)
984  ret void
985}
986define void @st_param_v4_i16_rrri(i16 %a, i16 %b, i16 %c) {
987; CHECK-LABEL: st_param_v4_i16_rrri(
988; CHECK:       {
989; CHECK-NEXT:    .reg .b16 %rs<4>;
990; CHECK-EMPTY:
991; CHECK-NEXT:  // %bb.0:
992; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_rrri_param_0];
993; CHECK-NEXT:    ld.param.u16 %rs2, [st_param_v4_i16_rrri_param_1];
994; CHECK-NEXT:    ld.param.u16 %rs3, [st_param_v4_i16_rrri_param_2];
995; CHECK-NEXT:    { // callseq 42, 0
996; CHECK-NEXT:    .param .align 8 .b8 param0[8];
997; CHECK-NEXT:    st.param.v4.b16 [param0], {%rs1, %rs2, %rs3, 4};
998; CHECK-NEXT:    call.uni
999; CHECK-NEXT:    call_v4_i16,
1000; CHECK-NEXT:    (
1001; CHECK-NEXT:    param0
1002; CHECK-NEXT:    );
1003; CHECK-NEXT:    } // callseq 42
1004; CHECK-NEXT:    ret;
1005  %struct.rrri0 = insertvalue %struct.short4 poison, i16 %a, 0
1006  %struct.rrri1 = insertvalue %struct.short4 %struct.rrri0, i16 %b, 1
1007  %struct.rrri2 = insertvalue %struct.short4 %struct.rrri1, i16 %c, 2
1008  %struct.rrri3 = insertvalue %struct.short4 %struct.rrri2, i16 4, 3
1009  call void @call_v4_i16(%struct.short4 %struct.rrri3)
1010  ret void
1011}
1012define void @st_param_v4_i16_iirr(i16 %c, i16 %d) {
1013; CHECK-LABEL: st_param_v4_i16_iirr(
1014; CHECK:       {
1015; CHECK-NEXT:    .reg .b16 %rs<3>;
1016; CHECK-EMPTY:
1017; CHECK-NEXT:  // %bb.0:
1018; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_iirr_param_0];
1019; CHECK-NEXT:    ld.param.u16 %rs2, [st_param_v4_i16_iirr_param_1];
1020; CHECK-NEXT:    { // callseq 43, 0
1021; CHECK-NEXT:    .param .align 8 .b8 param0[8];
1022; CHECK-NEXT:    st.param.v4.b16 [param0], {1, 2, %rs1, %rs2};
1023; CHECK-NEXT:    call.uni
1024; CHECK-NEXT:    call_v4_i16,
1025; CHECK-NEXT:    (
1026; CHECK-NEXT:    param0
1027; CHECK-NEXT:    );
1028; CHECK-NEXT:    } // callseq 43
1029; CHECK-NEXT:    ret;
1030  %struct.iirr0 = insertvalue %struct.short4 poison, i16 1, 0
1031  %struct.iirr1 = insertvalue %struct.short4 %struct.iirr0, i16 2, 1
1032  %struct.iirr2 = insertvalue %struct.short4 %struct.iirr1, i16 %c, 2
1033  %struct.iirr3 = insertvalue %struct.short4 %struct.iirr2, i16 %d, 3
1034  call void @call_v4_i16(%struct.short4 %struct.iirr3)
1035  ret void
1036}
1037define void @st_param_v4_i16_irir(i16 %b, i16 %d) {
1038; CHECK-LABEL: st_param_v4_i16_irir(
1039; CHECK:       {
1040; CHECK-NEXT:    .reg .b16 %rs<3>;
1041; CHECK-EMPTY:
1042; CHECK-NEXT:  // %bb.0:
1043; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_irir_param_0];
1044; CHECK-NEXT:    ld.param.u16 %rs2, [st_param_v4_i16_irir_param_1];
1045; CHECK-NEXT:    { // callseq 44, 0
1046; CHECK-NEXT:    .param .align 8 .b8 param0[8];
1047; CHECK-NEXT:    st.param.v4.b16 [param0], {1, %rs1, 3, %rs2};
1048; CHECK-NEXT:    call.uni
1049; CHECK-NEXT:    call_v4_i16,
1050; CHECK-NEXT:    (
1051; CHECK-NEXT:    param0
1052; CHECK-NEXT:    );
1053; CHECK-NEXT:    } // callseq 44
1054; CHECK-NEXT:    ret;
1055  %struct.irir0 = insertvalue %struct.short4 poison, i16 1, 0
1056  %struct.irir1 = insertvalue %struct.short4 %struct.irir0, i16 %b, 1
1057  %struct.irir2 = insertvalue %struct.short4 %struct.irir1, i16 3, 2
1058  %struct.irir3 = insertvalue %struct.short4 %struct.irir2, i16 %d, 3
1059  call void @call_v4_i16(%struct.short4 %struct.irir3)
1060  ret void
1061}
1062define void @st_param_v4_i16_irri(i16 %b, i16 %c) {
1063; CHECK-LABEL: st_param_v4_i16_irri(
1064; CHECK:       {
1065; CHECK-NEXT:    .reg .b16 %rs<3>;
1066; CHECK-EMPTY:
1067; CHECK-NEXT:  // %bb.0:
1068; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_irri_param_0];
1069; CHECK-NEXT:    ld.param.u16 %rs2, [st_param_v4_i16_irri_param_1];
1070; CHECK-NEXT:    { // callseq 45, 0
1071; CHECK-NEXT:    .param .align 8 .b8 param0[8];
1072; CHECK-NEXT:    st.param.v4.b16 [param0], {1, %rs1, %rs2, 4};
1073; CHECK-NEXT:    call.uni
1074; CHECK-NEXT:    call_v4_i16,
1075; CHECK-NEXT:    (
1076; CHECK-NEXT:    param0
1077; CHECK-NEXT:    );
1078; CHECK-NEXT:    } // callseq 45
1079; CHECK-NEXT:    ret;
1080  %struct.irri0 = insertvalue %struct.short4 poison, i16 1, 0
1081  %struct.irri1 = insertvalue %struct.short4 %struct.irri0, i16 %b, 1
1082  %struct.irri2 = insertvalue %struct.short4 %struct.irri1, i16 %c, 2
1083  %struct.irri3 = insertvalue %struct.short4 %struct.irri2, i16 4, 3
1084  call void @call_v4_i16(%struct.short4 %struct.irri3)
1085  ret void
1086}
1087define void @st_param_v4_i16_riir(i16 %a, i16 %d) {
1088; CHECK-LABEL: st_param_v4_i16_riir(
1089; CHECK:       {
1090; CHECK-NEXT:    .reg .b16 %rs<3>;
1091; CHECK-EMPTY:
1092; CHECK-NEXT:  // %bb.0:
1093; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_riir_param_0];
1094; CHECK-NEXT:    ld.param.u16 %rs2, [st_param_v4_i16_riir_param_1];
1095; CHECK-NEXT:    { // callseq 46, 0
1096; CHECK-NEXT:    .param .align 8 .b8 param0[8];
1097; CHECK-NEXT:    st.param.v4.b16 [param0], {%rs1, 2, 3, %rs2};
1098; CHECK-NEXT:    call.uni
1099; CHECK-NEXT:    call_v4_i16,
1100; CHECK-NEXT:    (
1101; CHECK-NEXT:    param0
1102; CHECK-NEXT:    );
1103; CHECK-NEXT:    } // callseq 46
1104; CHECK-NEXT:    ret;
1105  %struct.riir0 = insertvalue %struct.short4 poison, i16 %a, 0
1106  %struct.riir1 = insertvalue %struct.short4 %struct.riir0, i16 2, 1
1107  %struct.riir2 = insertvalue %struct.short4 %struct.riir1, i16 3, 2
1108  %struct.riir3 = insertvalue %struct.short4 %struct.riir2, i16 %d, 3
1109  call void @call_v4_i16(%struct.short4 %struct.riir3)
1110  ret void
1111}
1112define void @st_param_v4_i16_riri(i16 %a, i16 %c) {
1113; CHECK-LABEL: st_param_v4_i16_riri(
1114; CHECK:       {
1115; CHECK-NEXT:    .reg .b16 %rs<3>;
1116; CHECK-EMPTY:
1117; CHECK-NEXT:  // %bb.0:
1118; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_riri_param_0];
1119; CHECK-NEXT:    ld.param.u16 %rs2, [st_param_v4_i16_riri_param_1];
1120; CHECK-NEXT:    { // callseq 47, 0
1121; CHECK-NEXT:    .param .align 8 .b8 param0[8];
1122; CHECK-NEXT:    st.param.v4.b16 [param0], {%rs1, 2, %rs2, 4};
1123; CHECK-NEXT:    call.uni
1124; CHECK-NEXT:    call_v4_i16,
1125; CHECK-NEXT:    (
1126; CHECK-NEXT:    param0
1127; CHECK-NEXT:    );
1128; CHECK-NEXT:    } // callseq 47
1129; CHECK-NEXT:    ret;
1130  %struct.riri0 = insertvalue %struct.short4 poison, i16 %a, 0
1131  %struct.riri1 = insertvalue %struct.short4 %struct.riri0, i16 2, 1
1132  %struct.riri2 = insertvalue %struct.short4 %struct.riri1, i16 %c, 2
1133  %struct.riri3 = insertvalue %struct.short4 %struct.riri2, i16 4, 3
1134  call void @call_v4_i16(%struct.short4 %struct.riri3)
1135  ret void
1136}
1137define void @st_param_v4_i16_rrii(i16 %a, i16 %b) {
1138; CHECK-LABEL: st_param_v4_i16_rrii(
1139; CHECK:       {
1140; CHECK-NEXT:    .reg .b16 %rs<3>;
1141; CHECK-EMPTY:
1142; CHECK-NEXT:  // %bb.0:
1143; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_rrii_param_0];
1144; CHECK-NEXT:    ld.param.u16 %rs2, [st_param_v4_i16_rrii_param_1];
1145; CHECK-NEXT:    { // callseq 48, 0
1146; CHECK-NEXT:    .param .align 8 .b8 param0[8];
1147; CHECK-NEXT:    st.param.v4.b16 [param0], {%rs1, %rs2, 3, 4};
1148; CHECK-NEXT:    call.uni
1149; CHECK-NEXT:    call_v4_i16,
1150; CHECK-NEXT:    (
1151; CHECK-NEXT:    param0
1152; CHECK-NEXT:    );
1153; CHECK-NEXT:    } // callseq 48
1154; CHECK-NEXT:    ret;
1155  %struct.rrii0 = insertvalue %struct.short4 poison, i16 %a, 0
1156  %struct.rrii1 = insertvalue %struct.short4 %struct.rrii0, i16 %b, 1
1157  %struct.rrii2 = insertvalue %struct.short4 %struct.rrii1, i16 3, 2
1158  %struct.rrii3 = insertvalue %struct.short4 %struct.rrii2, i16 4, 3
1159  call void @call_v4_i16(%struct.short4 %struct.rrii3)
1160  ret void
1161}
1162define void @st_param_v4_i16_iiir(i16 %d) {
1163; CHECK-LABEL: st_param_v4_i16_iiir(
1164; CHECK:       {
1165; CHECK-NEXT:    .reg .b16 %rs<2>;
1166; CHECK-EMPTY:
1167; CHECK-NEXT:  // %bb.0:
1168; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_iiir_param_0];
1169; CHECK-NEXT:    { // callseq 49, 0
1170; CHECK-NEXT:    .param .align 8 .b8 param0[8];
1171; CHECK-NEXT:    st.param.v4.b16 [param0], {1, 2, 3, %rs1};
1172; CHECK-NEXT:    call.uni
1173; CHECK-NEXT:    call_v4_i16,
1174; CHECK-NEXT:    (
1175; CHECK-NEXT:    param0
1176; CHECK-NEXT:    );
1177; CHECK-NEXT:    } // callseq 49
1178; CHECK-NEXT:    ret;
1179  %struct.iiir0 = insertvalue %struct.short4 poison, i16 1, 0
1180  %struct.iiir1 = insertvalue %struct.short4 %struct.iiir0, i16 2, 1
1181  %struct.iiir2 = insertvalue %struct.short4 %struct.iiir1, i16 3, 2
1182  %struct.iiir3 = insertvalue %struct.short4 %struct.iiir2, i16 %d, 3
1183  call void @call_v4_i16(%struct.short4 %struct.iiir3)
1184  ret void
1185}
1186define void @st_param_v4_i16_iiri(i16 %c) {
1187; CHECK-LABEL: st_param_v4_i16_iiri(
1188; CHECK:       {
1189; CHECK-NEXT:    .reg .b16 %rs<2>;
1190; CHECK-EMPTY:
1191; CHECK-NEXT:  // %bb.0:
1192; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_iiri_param_0];
1193; CHECK-NEXT:    { // callseq 50, 0
1194; CHECK-NEXT:    .param .align 8 .b8 param0[8];
1195; CHECK-NEXT:    st.param.v4.b16 [param0], {1, 2, %rs1, 4};
1196; CHECK-NEXT:    call.uni
1197; CHECK-NEXT:    call_v4_i16,
1198; CHECK-NEXT:    (
1199; CHECK-NEXT:    param0
1200; CHECK-NEXT:    );
1201; CHECK-NEXT:    } // callseq 50
1202; CHECK-NEXT:    ret;
1203  %struct.iiri0 = insertvalue %struct.short4 poison, i16 1, 0
1204  %struct.iiri1 = insertvalue %struct.short4 %struct.iiri0, i16 2, 1
1205  %struct.iiri2 = insertvalue %struct.short4 %struct.iiri1, i16 %c, 2
1206  %struct.iiri3 = insertvalue %struct.short4 %struct.iiri2, i16 4, 3
1207  call void @call_v4_i16(%struct.short4 %struct.iiri3)
1208  ret void
1209}
1210define void @st_param_v4_i16_irii(i16 %b) {
1211; CHECK-LABEL: st_param_v4_i16_irii(
1212; CHECK:       {
1213; CHECK-NEXT:    .reg .b16 %rs<2>;
1214; CHECK-EMPTY:
1215; CHECK-NEXT:  // %bb.0:
1216; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_irii_param_0];
1217; CHECK-NEXT:    { // callseq 51, 0
1218; CHECK-NEXT:    .param .align 8 .b8 param0[8];
1219; CHECK-NEXT:    st.param.v4.b16 [param0], {1, %rs1, 3, 4};
1220; CHECK-NEXT:    call.uni
1221; CHECK-NEXT:    call_v4_i16,
1222; CHECK-NEXT:    (
1223; CHECK-NEXT:    param0
1224; CHECK-NEXT:    );
1225; CHECK-NEXT:    } // callseq 51
1226; CHECK-NEXT:    ret;
1227  %struct.irii0 = insertvalue %struct.short4 poison, i16 1, 0
1228  %struct.irii1 = insertvalue %struct.short4 %struct.irii0, i16 %b, 1
1229  %struct.irii2 = insertvalue %struct.short4 %struct.irii1, i16 3, 2
1230  %struct.irii3 = insertvalue %struct.short4 %struct.irii2, i16 4, 3
1231  call void @call_v4_i16(%struct.short4 %struct.irii3)
1232  ret void
1233}
1234define void @st_param_v4_i16_riii(i16 %a) {
1235; CHECK-LABEL: st_param_v4_i16_riii(
1236; CHECK:       {
1237; CHECK-NEXT:    .reg .b16 %rs<2>;
1238; CHECK-EMPTY:
1239; CHECK-NEXT:  // %bb.0:
1240; CHECK-NEXT:    ld.param.u16 %rs1, [st_param_v4_i16_riii_param_0];
1241; CHECK-NEXT:    { // callseq 52, 0
1242; CHECK-NEXT:    .param .align 8 .b8 param0[8];
1243; CHECK-NEXT:    st.param.v4.b16 [param0], {%rs1, 2, 3, 4};
1244; CHECK-NEXT:    call.uni
1245; CHECK-NEXT:    call_v4_i16,
1246; CHECK-NEXT:    (
1247; CHECK-NEXT:    param0
1248; CHECK-NEXT:    );
1249; CHECK-NEXT:    } // callseq 52
1250; CHECK-NEXT:    ret;
1251  %struct.riii0 = insertvalue %struct.short4 poison, i16 %a, 0
1252  %struct.riii1 = insertvalue %struct.short4 %struct.riii0, i16 2, 1
1253  %struct.riii2 = insertvalue %struct.short4 %struct.riii1, i16 3, 2
1254  %struct.riii3 = insertvalue %struct.short4 %struct.riii2, i16 4, 3
1255  call void @call_v4_i16(%struct.short4 %struct.riii3)
1256  ret void
1257}
1258
1259define void @st_param_v4_i32_iiii() {
1260; CHECK-LABEL: st_param_v4_i32_iiii(
1261; CHECK:       {
1262; CHECK-EMPTY:
1263; CHECK-EMPTY:
1264; CHECK-NEXT:  // %bb.0:
1265; CHECK-NEXT:    { // callseq 53, 0
1266; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1267; CHECK-NEXT:    st.param.v4.b32 [param0], {1, 2, 3, 4};
1268; CHECK-NEXT:    call.uni
1269; CHECK-NEXT:    call_v4_i32,
1270; CHECK-NEXT:    (
1271; CHECK-NEXT:    param0
1272; CHECK-NEXT:    );
1273; CHECK-NEXT:    } // callseq 53
1274; CHECK-NEXT:    ret;
1275  call void @call_v4_i32(%struct.int4 { i32 1, i32 2, i32 3, i32 4 })
1276  ret void
1277}
1278define void @st_param_v4_i32_irrr(i32 %b, i32 %c, i32 %d) {
1279; CHECK-LABEL: st_param_v4_i32_irrr(
1280; CHECK:       {
1281; CHECK-NEXT:    .reg .b32 %r<4>;
1282; CHECK-EMPTY:
1283; CHECK-NEXT:  // %bb.0:
1284; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_irrr_param_0];
1285; CHECK-NEXT:    ld.param.u32 %r2, [st_param_v4_i32_irrr_param_1];
1286; CHECK-NEXT:    ld.param.u32 %r3, [st_param_v4_i32_irrr_param_2];
1287; CHECK-NEXT:    { // callseq 54, 0
1288; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1289; CHECK-NEXT:    st.param.v4.b32 [param0], {1, %r1, %r2, %r3};
1290; CHECK-NEXT:    call.uni
1291; CHECK-NEXT:    call_v4_i32,
1292; CHECK-NEXT:    (
1293; CHECK-NEXT:    param0
1294; CHECK-NEXT:    );
1295; CHECK-NEXT:    } // callseq 54
1296; CHECK-NEXT:    ret;
1297  %struct.irrr0 = insertvalue %struct.int4 poison, i32 1, 0
1298  %struct.irrr1 = insertvalue %struct.int4 %struct.irrr0, i32 %b, 1
1299  %struct.irrr2 = insertvalue %struct.int4 %struct.irrr1, i32 %c, 2
1300  %struct.irrr3 = insertvalue %struct.int4 %struct.irrr2, i32 %d, 3
1301  call void @call_v4_i32(%struct.int4 %struct.irrr3)
1302  ret void
1303}
1304define void @st_param_v4_i32_rirr(i32 %a, i32 %c, i32 %d) {
1305; CHECK-LABEL: st_param_v4_i32_rirr(
1306; CHECK:       {
1307; CHECK-NEXT:    .reg .b32 %r<4>;
1308; CHECK-EMPTY:
1309; CHECK-NEXT:  // %bb.0:
1310; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_rirr_param_0];
1311; CHECK-NEXT:    ld.param.u32 %r2, [st_param_v4_i32_rirr_param_1];
1312; CHECK-NEXT:    ld.param.u32 %r3, [st_param_v4_i32_rirr_param_2];
1313; CHECK-NEXT:    { // callseq 55, 0
1314; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1315; CHECK-NEXT:    st.param.v4.b32 [param0], {%r1, 2, %r2, %r3};
1316; CHECK-NEXT:    call.uni
1317; CHECK-NEXT:    call_v4_i32,
1318; CHECK-NEXT:    (
1319; CHECK-NEXT:    param0
1320; CHECK-NEXT:    );
1321; CHECK-NEXT:    } // callseq 55
1322; CHECK-NEXT:    ret;
1323  %struct.rirr0 = insertvalue %struct.int4 poison, i32 %a, 0
1324  %struct.rirr1 = insertvalue %struct.int4 %struct.rirr0, i32 2, 1
1325  %struct.rirr2 = insertvalue %struct.int4 %struct.rirr1, i32 %c, 2
1326  %struct.rirr3 = insertvalue %struct.int4 %struct.rirr2, i32 %d, 3
1327  call void @call_v4_i32(%struct.int4 %struct.rirr3)
1328  ret void
1329}
1330define void @st_param_v4_i32_rrir(i32 %a, i32 %b, i32 %d) {
1331; CHECK-LABEL: st_param_v4_i32_rrir(
1332; CHECK:       {
1333; CHECK-NEXT:    .reg .b32 %r<4>;
1334; CHECK-EMPTY:
1335; CHECK-NEXT:  // %bb.0:
1336; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_rrir_param_0];
1337; CHECK-NEXT:    ld.param.u32 %r2, [st_param_v4_i32_rrir_param_1];
1338; CHECK-NEXT:    ld.param.u32 %r3, [st_param_v4_i32_rrir_param_2];
1339; CHECK-NEXT:    { // callseq 56, 0
1340; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1341; CHECK-NEXT:    st.param.v4.b32 [param0], {%r1, %r2, 3, %r3};
1342; CHECK-NEXT:    call.uni
1343; CHECK-NEXT:    call_v4_i32,
1344; CHECK-NEXT:    (
1345; CHECK-NEXT:    param0
1346; CHECK-NEXT:    );
1347; CHECK-NEXT:    } // callseq 56
1348; CHECK-NEXT:    ret;
1349  %struct.rrir0 = insertvalue %struct.int4 poison, i32 %a, 0
1350  %struct.rrir1 = insertvalue %struct.int4 %struct.rrir0, i32 %b, 1
1351  %struct.rrir2 = insertvalue %struct.int4 %struct.rrir1, i32 3, 2
1352  %struct.rrir3 = insertvalue %struct.int4 %struct.rrir2, i32 %d, 3
1353  call void @call_v4_i32(%struct.int4 %struct.rrir3)
1354  ret void
1355}
1356define void @st_param_v4_i32_rrri(i32 %a, i32 %b, i32 %c) {
1357; CHECK-LABEL: st_param_v4_i32_rrri(
1358; CHECK:       {
1359; CHECK-NEXT:    .reg .b32 %r<4>;
1360; CHECK-EMPTY:
1361; CHECK-NEXT:  // %bb.0:
1362; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_rrri_param_0];
1363; CHECK-NEXT:    ld.param.u32 %r2, [st_param_v4_i32_rrri_param_1];
1364; CHECK-NEXT:    ld.param.u32 %r3, [st_param_v4_i32_rrri_param_2];
1365; CHECK-NEXT:    { // callseq 57, 0
1366; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1367; CHECK-NEXT:    st.param.v4.b32 [param0], {%r1, %r2, %r3, 4};
1368; CHECK-NEXT:    call.uni
1369; CHECK-NEXT:    call_v4_i32,
1370; CHECK-NEXT:    (
1371; CHECK-NEXT:    param0
1372; CHECK-NEXT:    );
1373; CHECK-NEXT:    } // callseq 57
1374; CHECK-NEXT:    ret;
1375  %struct.rrri0 = insertvalue %struct.int4 poison, i32 %a, 0
1376  %struct.rrri1 = insertvalue %struct.int4 %struct.rrri0, i32 %b, 1
1377  %struct.rrri2 = insertvalue %struct.int4 %struct.rrri1, i32 %c, 2
1378  %struct.rrri3 = insertvalue %struct.int4 %struct.rrri2, i32 4, 3
1379  call void @call_v4_i32(%struct.int4 %struct.rrri3)
1380  ret void
1381}
1382define void @st_param_v4_i32_iirr(i32 %c, i32 %d) {
1383; CHECK-LABEL: st_param_v4_i32_iirr(
1384; CHECK:       {
1385; CHECK-NEXT:    .reg .b32 %r<3>;
1386; CHECK-EMPTY:
1387; CHECK-NEXT:  // %bb.0:
1388; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_iirr_param_0];
1389; CHECK-NEXT:    ld.param.u32 %r2, [st_param_v4_i32_iirr_param_1];
1390; CHECK-NEXT:    { // callseq 58, 0
1391; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1392; CHECK-NEXT:    st.param.v4.b32 [param0], {1, 2, %r1, %r2};
1393; CHECK-NEXT:    call.uni
1394; CHECK-NEXT:    call_v4_i32,
1395; CHECK-NEXT:    (
1396; CHECK-NEXT:    param0
1397; CHECK-NEXT:    );
1398; CHECK-NEXT:    } // callseq 58
1399; CHECK-NEXT:    ret;
1400  %struct.iirr0 = insertvalue %struct.int4 poison, i32 1, 0
1401  %struct.iirr1 = insertvalue %struct.int4 %struct.iirr0, i32 2, 1
1402  %struct.iirr2 = insertvalue %struct.int4 %struct.iirr1, i32 %c, 2
1403  %struct.iirr3 = insertvalue %struct.int4 %struct.iirr2, i32 %d, 3
1404  call void @call_v4_i32(%struct.int4 %struct.iirr3)
1405  ret void
1406}
1407define void @st_param_v4_i32_irir(i32 %b, i32 %d) {
1408; CHECK-LABEL: st_param_v4_i32_irir(
1409; CHECK:       {
1410; CHECK-NEXT:    .reg .b32 %r<3>;
1411; CHECK-EMPTY:
1412; CHECK-NEXT:  // %bb.0:
1413; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_irir_param_0];
1414; CHECK-NEXT:    ld.param.u32 %r2, [st_param_v4_i32_irir_param_1];
1415; CHECK-NEXT:    { // callseq 59, 0
1416; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1417; CHECK-NEXT:    st.param.v4.b32 [param0], {1, %r1, 3, %r2};
1418; CHECK-NEXT:    call.uni
1419; CHECK-NEXT:    call_v4_i32,
1420; CHECK-NEXT:    (
1421; CHECK-NEXT:    param0
1422; CHECK-NEXT:    );
1423; CHECK-NEXT:    } // callseq 59
1424; CHECK-NEXT:    ret;
1425  %struct.irir0 = insertvalue %struct.int4 poison, i32 1, 0
1426  %struct.irir1 = insertvalue %struct.int4 %struct.irir0, i32 %b, 1
1427  %struct.irir2 = insertvalue %struct.int4 %struct.irir1, i32 3, 2
1428  %struct.irir3 = insertvalue %struct.int4 %struct.irir2, i32 %d, 3
1429  call void @call_v4_i32(%struct.int4 %struct.irir3)
1430  ret void
1431}
1432define void @st_param_v4_i32_irri(i32 %b, i32 %c) {
1433; CHECK-LABEL: st_param_v4_i32_irri(
1434; CHECK:       {
1435; CHECK-NEXT:    .reg .b32 %r<3>;
1436; CHECK-EMPTY:
1437; CHECK-NEXT:  // %bb.0:
1438; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_irri_param_0];
1439; CHECK-NEXT:    ld.param.u32 %r2, [st_param_v4_i32_irri_param_1];
1440; CHECK-NEXT:    { // callseq 60, 0
1441; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1442; CHECK-NEXT:    st.param.v4.b32 [param0], {1, %r1, %r2, 4};
1443; CHECK-NEXT:    call.uni
1444; CHECK-NEXT:    call_v4_i32,
1445; CHECK-NEXT:    (
1446; CHECK-NEXT:    param0
1447; CHECK-NEXT:    );
1448; CHECK-NEXT:    } // callseq 60
1449; CHECK-NEXT:    ret;
1450  %struct.irri0 = insertvalue %struct.int4 poison, i32 1, 0
1451  %struct.irri1 = insertvalue %struct.int4 %struct.irri0, i32 %b, 1
1452  %struct.irri2 = insertvalue %struct.int4 %struct.irri1, i32 %c, 2
1453  %struct.irri3 = insertvalue %struct.int4 %struct.irri2, i32 4, 3
1454  call void @call_v4_i32(%struct.int4 %struct.irri3)
1455  ret void
1456}
1457define void @st_param_v4_i32_riir(i32 %a, i32 %d) {
1458; CHECK-LABEL: st_param_v4_i32_riir(
1459; CHECK:       {
1460; CHECK-NEXT:    .reg .b32 %r<3>;
1461; CHECK-EMPTY:
1462; CHECK-NEXT:  // %bb.0:
1463; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_riir_param_0];
1464; CHECK-NEXT:    ld.param.u32 %r2, [st_param_v4_i32_riir_param_1];
1465; CHECK-NEXT:    { // callseq 61, 0
1466; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1467; CHECK-NEXT:    st.param.v4.b32 [param0], {%r1, 2, 3, %r2};
1468; CHECK-NEXT:    call.uni
1469; CHECK-NEXT:    call_v4_i32,
1470; CHECK-NEXT:    (
1471; CHECK-NEXT:    param0
1472; CHECK-NEXT:    );
1473; CHECK-NEXT:    } // callseq 61
1474; CHECK-NEXT:    ret;
1475  %struct.riir0 = insertvalue %struct.int4 poison, i32 %a, 0
1476  %struct.riir1 = insertvalue %struct.int4 %struct.riir0, i32 2, 1
1477  %struct.riir2 = insertvalue %struct.int4 %struct.riir1, i32 3, 2
1478  %struct.riir3 = insertvalue %struct.int4 %struct.riir2, i32 %d, 3
1479  call void @call_v4_i32(%struct.int4 %struct.riir3)
1480  ret void
1481}
1482define void @st_param_v4_i32_riri(i32 %a, i32 %c) {
1483; CHECK-LABEL: st_param_v4_i32_riri(
1484; CHECK:       {
1485; CHECK-NEXT:    .reg .b32 %r<3>;
1486; CHECK-EMPTY:
1487; CHECK-NEXT:  // %bb.0:
1488; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_riri_param_0];
1489; CHECK-NEXT:    ld.param.u32 %r2, [st_param_v4_i32_riri_param_1];
1490; CHECK-NEXT:    { // callseq 62, 0
1491; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1492; CHECK-NEXT:    st.param.v4.b32 [param0], {%r1, 2, %r2, 4};
1493; CHECK-NEXT:    call.uni
1494; CHECK-NEXT:    call_v4_i32,
1495; CHECK-NEXT:    (
1496; CHECK-NEXT:    param0
1497; CHECK-NEXT:    );
1498; CHECK-NEXT:    } // callseq 62
1499; CHECK-NEXT:    ret;
1500  %struct.riri0 = insertvalue %struct.int4 poison, i32 %a, 0
1501  %struct.riri1 = insertvalue %struct.int4 %struct.riri0, i32 2, 1
1502  %struct.riri2 = insertvalue %struct.int4 %struct.riri1, i32 %c, 2
1503  %struct.riri3 = insertvalue %struct.int4 %struct.riri2, i32 4, 3
1504  call void @call_v4_i32(%struct.int4 %struct.riri3)
1505  ret void
1506}
1507define void @st_param_v4_i32_rrii(i32 %a, i32 %b) {
1508; CHECK-LABEL: st_param_v4_i32_rrii(
1509; CHECK:       {
1510; CHECK-NEXT:    .reg .b32 %r<3>;
1511; CHECK-EMPTY:
1512; CHECK-NEXT:  // %bb.0:
1513; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_rrii_param_0];
1514; CHECK-NEXT:    ld.param.u32 %r2, [st_param_v4_i32_rrii_param_1];
1515; CHECK-NEXT:    { // callseq 63, 0
1516; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1517; CHECK-NEXT:    st.param.v4.b32 [param0], {%r1, %r2, 3, 4};
1518; CHECK-NEXT:    call.uni
1519; CHECK-NEXT:    call_v4_i32,
1520; CHECK-NEXT:    (
1521; CHECK-NEXT:    param0
1522; CHECK-NEXT:    );
1523; CHECK-NEXT:    } // callseq 63
1524; CHECK-NEXT:    ret;
1525  %struct.rrii0 = insertvalue %struct.int4 poison, i32 %a, 0
1526  %struct.rrii1 = insertvalue %struct.int4 %struct.rrii0, i32 %b, 1
1527  %struct.rrii2 = insertvalue %struct.int4 %struct.rrii1, i32 3, 2
1528  %struct.rrii3 = insertvalue %struct.int4 %struct.rrii2, i32 4, 3
1529  call void @call_v4_i32(%struct.int4 %struct.rrii3)
1530  ret void
1531}
1532define void @st_param_v4_i32_iiir(i32 %d) {
1533; CHECK-LABEL: st_param_v4_i32_iiir(
1534; CHECK:       {
1535; CHECK-NEXT:    .reg .b32 %r<2>;
1536; CHECK-EMPTY:
1537; CHECK-NEXT:  // %bb.0:
1538; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_iiir_param_0];
1539; CHECK-NEXT:    { // callseq 64, 0
1540; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1541; CHECK-NEXT:    st.param.v4.b32 [param0], {1, 2, 3, %r1};
1542; CHECK-NEXT:    call.uni
1543; CHECK-NEXT:    call_v4_i32,
1544; CHECK-NEXT:    (
1545; CHECK-NEXT:    param0
1546; CHECK-NEXT:    );
1547; CHECK-NEXT:    } // callseq 64
1548; CHECK-NEXT:    ret;
1549  %struct.iiir0 = insertvalue %struct.int4 poison, i32 1, 0
1550  %struct.iiir1 = insertvalue %struct.int4 %struct.iiir0, i32 2, 1
1551  %struct.iiir2 = insertvalue %struct.int4 %struct.iiir1, i32 3, 2
1552  %struct.iiir3 = insertvalue %struct.int4 %struct.iiir2, i32 %d, 3
1553  call void @call_v4_i32(%struct.int4 %struct.iiir3)
1554  ret void
1555}
1556define void @st_param_v4_i32_iiri(i32 %c) {
1557; CHECK-LABEL: st_param_v4_i32_iiri(
1558; CHECK:       {
1559; CHECK-NEXT:    .reg .b32 %r<2>;
1560; CHECK-EMPTY:
1561; CHECK-NEXT:  // %bb.0:
1562; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_iiri_param_0];
1563; CHECK-NEXT:    { // callseq 65, 0
1564; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1565; CHECK-NEXT:    st.param.v4.b32 [param0], {1, 2, %r1, 4};
1566; CHECK-NEXT:    call.uni
1567; CHECK-NEXT:    call_v4_i32,
1568; CHECK-NEXT:    (
1569; CHECK-NEXT:    param0
1570; CHECK-NEXT:    );
1571; CHECK-NEXT:    } // callseq 65
1572; CHECK-NEXT:    ret;
1573  %struct.iiri0 = insertvalue %struct.int4 poison, i32 1, 0
1574  %struct.iiri1 = insertvalue %struct.int4 %struct.iiri0, i32 2, 1
1575  %struct.iiri2 = insertvalue %struct.int4 %struct.iiri1, i32 %c, 2
1576  %struct.iiri3 = insertvalue %struct.int4 %struct.iiri2, i32 4, 3
1577  call void @call_v4_i32(%struct.int4 %struct.iiri3)
1578  ret void
1579}
1580define void @st_param_v4_i32_irii(i32 %b) {
1581; CHECK-LABEL: st_param_v4_i32_irii(
1582; CHECK:       {
1583; CHECK-NEXT:    .reg .b32 %r<2>;
1584; CHECK-EMPTY:
1585; CHECK-NEXT:  // %bb.0:
1586; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_irii_param_0];
1587; CHECK-NEXT:    { // callseq 66, 0
1588; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1589; CHECK-NEXT:    st.param.v4.b32 [param0], {1, %r1, 3, 4};
1590; CHECK-NEXT:    call.uni
1591; CHECK-NEXT:    call_v4_i32,
1592; CHECK-NEXT:    (
1593; CHECK-NEXT:    param0
1594; CHECK-NEXT:    );
1595; CHECK-NEXT:    } // callseq 66
1596; CHECK-NEXT:    ret;
1597  %struct.irii0 = insertvalue %struct.int4 poison, i32 1, 0
1598  %struct.irii1 = insertvalue %struct.int4 %struct.irii0, i32 %b, 1
1599  %struct.irii2 = insertvalue %struct.int4 %struct.irii1, i32 3, 2
1600  %struct.irii3 = insertvalue %struct.int4 %struct.irii2, i32 4, 3
1601  call void @call_v4_i32(%struct.int4 %struct.irii3)
1602  ret void
1603}
1604define void @st_param_v4_i32_riii(i32 %a) {
1605; CHECK-LABEL: st_param_v4_i32_riii(
1606; CHECK:       {
1607; CHECK-NEXT:    .reg .b32 %r<2>;
1608; CHECK-EMPTY:
1609; CHECK-NEXT:  // %bb.0:
1610; CHECK-NEXT:    ld.param.u32 %r1, [st_param_v4_i32_riii_param_0];
1611; CHECK-NEXT:    { // callseq 67, 0
1612; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1613; CHECK-NEXT:    st.param.v4.b32 [param0], {%r1, 2, 3, 4};
1614; CHECK-NEXT:    call.uni
1615; CHECK-NEXT:    call_v4_i32,
1616; CHECK-NEXT:    (
1617; CHECK-NEXT:    param0
1618; CHECK-NEXT:    );
1619; CHECK-NEXT:    } // callseq 67
1620; CHECK-NEXT:    ret;
1621  %struct.riii0 = insertvalue %struct.int4 poison, i32 %a, 0
1622  %struct.riii1 = insertvalue %struct.int4 %struct.riii0, i32 2, 1
1623  %struct.riii2 = insertvalue %struct.int4 %struct.riii1, i32 3, 2
1624  %struct.riii3 = insertvalue %struct.int4 %struct.riii2, i32 4, 3
1625  call void @call_v4_i32(%struct.int4 %struct.riii3)
1626  ret void
1627}
1628
1629define void @st_param_v4_f32_iiii() {
1630; CHECK-LABEL: st_param_v4_f32_iiii(
1631; CHECK:       {
1632; CHECK-EMPTY:
1633; CHECK-EMPTY:
1634; CHECK-NEXT:  // %bb.0:
1635; CHECK-NEXT:    { // callseq 68, 0
1636; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1637; CHECK-NEXT:    st.param.v4.f32 [param0], {0f3F800000, 0f40000000, 0f40400000, 0f40800000};
1638; CHECK-NEXT:    call.uni
1639; CHECK-NEXT:    call_v4_f32,
1640; CHECK-NEXT:    (
1641; CHECK-NEXT:    param0
1642; CHECK-NEXT:    );
1643; CHECK-NEXT:    } // callseq 68
1644; CHECK-NEXT:    ret;
1645  call void @call_v4_f32(%struct.float4 { float 1.0, float 2.0, float 3.0, float 4.0 })
1646  ret void
1647}
1648define void @st_param_v4_f32_irrr(float %b, float %c, float %d) {
1649; CHECK-LABEL: st_param_v4_f32_irrr(
1650; CHECK:       {
1651; CHECK-NEXT:    .reg .f32 %f<4>;
1652; CHECK-EMPTY:
1653; CHECK-NEXT:  // %bb.0:
1654; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_irrr_param_0];
1655; CHECK-NEXT:    ld.param.f32 %f2, [st_param_v4_f32_irrr_param_1];
1656; CHECK-NEXT:    ld.param.f32 %f3, [st_param_v4_f32_irrr_param_2];
1657; CHECK-NEXT:    { // callseq 69, 0
1658; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1659; CHECK-NEXT:    st.param.v4.f32 [param0], {0f3F800000, %f1, %f2, %f3};
1660; CHECK-NEXT:    call.uni
1661; CHECK-NEXT:    call_v4_f32,
1662; CHECK-NEXT:    (
1663; CHECK-NEXT:    param0
1664; CHECK-NEXT:    );
1665; CHECK-NEXT:    } // callseq 69
1666; CHECK-NEXT:    ret;
1667  %struct.irrr0 = insertvalue %struct.float4 poison, float 1.0, 0
1668  %struct.irrr1 = insertvalue %struct.float4 %struct.irrr0, float %b, 1
1669  %struct.irrr2 = insertvalue %struct.float4 %struct.irrr1, float %c, 2
1670  %struct.irrr3 = insertvalue %struct.float4 %struct.irrr2, float %d, 3
1671  call void @call_v4_f32(%struct.float4 %struct.irrr3)
1672  ret void
1673}
1674define void @st_param_v4_f32_rirr(float %a, float %c, float %d) {
1675; CHECK-LABEL: st_param_v4_f32_rirr(
1676; CHECK:       {
1677; CHECK-NEXT:    .reg .f32 %f<4>;
1678; CHECK-EMPTY:
1679; CHECK-NEXT:  // %bb.0:
1680; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_rirr_param_0];
1681; CHECK-NEXT:    ld.param.f32 %f2, [st_param_v4_f32_rirr_param_1];
1682; CHECK-NEXT:    ld.param.f32 %f3, [st_param_v4_f32_rirr_param_2];
1683; CHECK-NEXT:    { // callseq 70, 0
1684; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1685; CHECK-NEXT:    st.param.v4.f32 [param0], {%f1, 0f40000000, %f2, %f3};
1686; CHECK-NEXT:    call.uni
1687; CHECK-NEXT:    call_v4_f32,
1688; CHECK-NEXT:    (
1689; CHECK-NEXT:    param0
1690; CHECK-NEXT:    );
1691; CHECK-NEXT:    } // callseq 70
1692; CHECK-NEXT:    ret;
1693  %struct.rirr0 = insertvalue %struct.float4 poison, float %a, 0
1694  %struct.rirr1 = insertvalue %struct.float4 %struct.rirr0, float 2.0, 1
1695  %struct.rirr2 = insertvalue %struct.float4 %struct.rirr1, float %c, 2
1696  %struct.rirr3 = insertvalue %struct.float4 %struct.rirr2, float %d, 3
1697  call void @call_v4_f32(%struct.float4 %struct.rirr3)
1698  ret void
1699}
1700define void @st_param_v4_f32_rrir(float %a, float %b, float %d) {
1701; CHECK-LABEL: st_param_v4_f32_rrir(
1702; CHECK:       {
1703; CHECK-NEXT:    .reg .f32 %f<4>;
1704; CHECK-EMPTY:
1705; CHECK-NEXT:  // %bb.0:
1706; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_rrir_param_0];
1707; CHECK-NEXT:    ld.param.f32 %f2, [st_param_v4_f32_rrir_param_1];
1708; CHECK-NEXT:    ld.param.f32 %f3, [st_param_v4_f32_rrir_param_2];
1709; CHECK-NEXT:    { // callseq 71, 0
1710; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1711; CHECK-NEXT:    st.param.v4.f32 [param0], {%f1, %f2, 0f40400000, %f3};
1712; CHECK-NEXT:    call.uni
1713; CHECK-NEXT:    call_v4_f32,
1714; CHECK-NEXT:    (
1715; CHECK-NEXT:    param0
1716; CHECK-NEXT:    );
1717; CHECK-NEXT:    } // callseq 71
1718; CHECK-NEXT:    ret;
1719  %struct.rrir0 = insertvalue %struct.float4 poison, float %a, 0
1720  %struct.rrir1 = insertvalue %struct.float4 %struct.rrir0, float %b, 1
1721  %struct.rrir2 = insertvalue %struct.float4 %struct.rrir1, float 3.0, 2
1722  %struct.rrir3 = insertvalue %struct.float4 %struct.rrir2, float %d, 3
1723  call void @call_v4_f32(%struct.float4 %struct.rrir3)
1724  ret void
1725}
1726define void @st_param_v4_f32_rrri(float %a, float %b, float %c) {
1727; CHECK-LABEL: st_param_v4_f32_rrri(
1728; CHECK:       {
1729; CHECK-NEXT:    .reg .f32 %f<4>;
1730; CHECK-EMPTY:
1731; CHECK-NEXT:  // %bb.0:
1732; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_rrri_param_0];
1733; CHECK-NEXT:    ld.param.f32 %f2, [st_param_v4_f32_rrri_param_1];
1734; CHECK-NEXT:    ld.param.f32 %f3, [st_param_v4_f32_rrri_param_2];
1735; CHECK-NEXT:    { // callseq 72, 0
1736; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1737; CHECK-NEXT:    st.param.v4.f32 [param0], {%f1, %f2, %f3, 0f40800000};
1738; CHECK-NEXT:    call.uni
1739; CHECK-NEXT:    call_v4_f32,
1740; CHECK-NEXT:    (
1741; CHECK-NEXT:    param0
1742; CHECK-NEXT:    );
1743; CHECK-NEXT:    } // callseq 72
1744; CHECK-NEXT:    ret;
1745  %struct.rrri0 = insertvalue %struct.float4 poison, float %a, 0
1746  %struct.rrri1 = insertvalue %struct.float4 %struct.rrri0, float %b, 1
1747  %struct.rrri2 = insertvalue %struct.float4 %struct.rrri1, float %c, 2
1748  %struct.rrri3 = insertvalue %struct.float4 %struct.rrri2, float 4.0, 3
1749  call void @call_v4_f32(%struct.float4 %struct.rrri3)
1750  ret void
1751}
1752define void @st_param_v4_f32_iirr(float %c, float %d) {
1753; CHECK-LABEL: st_param_v4_f32_iirr(
1754; CHECK:       {
1755; CHECK-NEXT:    .reg .f32 %f<3>;
1756; CHECK-EMPTY:
1757; CHECK-NEXT:  // %bb.0:
1758; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_iirr_param_0];
1759; CHECK-NEXT:    ld.param.f32 %f2, [st_param_v4_f32_iirr_param_1];
1760; CHECK-NEXT:    { // callseq 73, 0
1761; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1762; CHECK-NEXT:    st.param.v4.f32 [param0], {0f3F800000, 0f40000000, %f1, %f2};
1763; CHECK-NEXT:    call.uni
1764; CHECK-NEXT:    call_v4_f32,
1765; CHECK-NEXT:    (
1766; CHECK-NEXT:    param0
1767; CHECK-NEXT:    );
1768; CHECK-NEXT:    } // callseq 73
1769; CHECK-NEXT:    ret;
1770  %struct.iirr0 = insertvalue %struct.float4 poison, float 1.0, 0
1771  %struct.iirr1 = insertvalue %struct.float4 %struct.iirr0, float 2.0, 1
1772  %struct.iirr2 = insertvalue %struct.float4 %struct.iirr1, float %c, 2
1773  %struct.iirr3 = insertvalue %struct.float4 %struct.iirr2, float %d, 3
1774  call void @call_v4_f32(%struct.float4 %struct.iirr3)
1775  ret void
1776}
1777define void @st_param_v4_f32_irir(float %b, float %d) {
1778; CHECK-LABEL: st_param_v4_f32_irir(
1779; CHECK:       {
1780; CHECK-NEXT:    .reg .f32 %f<3>;
1781; CHECK-EMPTY:
1782; CHECK-NEXT:  // %bb.0:
1783; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_irir_param_0];
1784; CHECK-NEXT:    ld.param.f32 %f2, [st_param_v4_f32_irir_param_1];
1785; CHECK-NEXT:    { // callseq 74, 0
1786; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1787; CHECK-NEXT:    st.param.v4.f32 [param0], {0f3F800000, %f1, 0f40400000, %f2};
1788; CHECK-NEXT:    call.uni
1789; CHECK-NEXT:    call_v4_f32,
1790; CHECK-NEXT:    (
1791; CHECK-NEXT:    param0
1792; CHECK-NEXT:    );
1793; CHECK-NEXT:    } // callseq 74
1794; CHECK-NEXT:    ret;
1795  %struct.irir0 = insertvalue %struct.float4 poison, float 1.0, 0
1796  %struct.irir1 = insertvalue %struct.float4 %struct.irir0, float %b, 1
1797  %struct.irir2 = insertvalue %struct.float4 %struct.irir1, float 3.0, 2
1798  %struct.irir3 = insertvalue %struct.float4 %struct.irir2, float %d, 3
1799  call void @call_v4_f32(%struct.float4 %struct.irir3)
1800  ret void
1801}
1802define void @st_param_v4_f32_irri(float %b, float %c) {
1803; CHECK-LABEL: st_param_v4_f32_irri(
1804; CHECK:       {
1805; CHECK-NEXT:    .reg .f32 %f<3>;
1806; CHECK-EMPTY:
1807; CHECK-NEXT:  // %bb.0:
1808; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_irri_param_0];
1809; CHECK-NEXT:    ld.param.f32 %f2, [st_param_v4_f32_irri_param_1];
1810; CHECK-NEXT:    { // callseq 75, 0
1811; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1812; CHECK-NEXT:    st.param.v4.f32 [param0], {0f3F800000, %f1, %f2, 0f40800000};
1813; CHECK-NEXT:    call.uni
1814; CHECK-NEXT:    call_v4_f32,
1815; CHECK-NEXT:    (
1816; CHECK-NEXT:    param0
1817; CHECK-NEXT:    );
1818; CHECK-NEXT:    } // callseq 75
1819; CHECK-NEXT:    ret;
1820  %struct.irri0 = insertvalue %struct.float4 poison, float 1.0, 0
1821  %struct.irri1 = insertvalue %struct.float4 %struct.irri0, float %b, 1
1822  %struct.irri2 = insertvalue %struct.float4 %struct.irri1, float %c, 2
1823  %struct.irri3 = insertvalue %struct.float4 %struct.irri2, float 4.0, 3
1824  call void @call_v4_f32(%struct.float4 %struct.irri3)
1825  ret void
1826}
1827define void @st_param_v4_f32_riir(float %a, float %d) {
1828; CHECK-LABEL: st_param_v4_f32_riir(
1829; CHECK:       {
1830; CHECK-NEXT:    .reg .f32 %f<3>;
1831; CHECK-EMPTY:
1832; CHECK-NEXT:  // %bb.0:
1833; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_riir_param_0];
1834; CHECK-NEXT:    ld.param.f32 %f2, [st_param_v4_f32_riir_param_1];
1835; CHECK-NEXT:    { // callseq 76, 0
1836; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1837; CHECK-NEXT:    st.param.v4.f32 [param0], {%f1, 0f40000000, 0f40400000, %f2};
1838; CHECK-NEXT:    call.uni
1839; CHECK-NEXT:    call_v4_f32,
1840; CHECK-NEXT:    (
1841; CHECK-NEXT:    param0
1842; CHECK-NEXT:    );
1843; CHECK-NEXT:    } // callseq 76
1844; CHECK-NEXT:    ret;
1845  %struct.riir0 = insertvalue %struct.float4 poison, float %a, 0
1846  %struct.riir1 = insertvalue %struct.float4 %struct.riir0, float 2.0, 1
1847  %struct.riir2 = insertvalue %struct.float4 %struct.riir1, float 3.0, 2
1848  %struct.riir3 = insertvalue %struct.float4 %struct.riir2, float %d, 3
1849  call void @call_v4_f32(%struct.float4 %struct.riir3)
1850  ret void
1851}
1852define void @st_param_v4_f32_riri(float %a, float %c) {
1853; CHECK-LABEL: st_param_v4_f32_riri(
1854; CHECK:       {
1855; CHECK-NEXT:    .reg .f32 %f<3>;
1856; CHECK-EMPTY:
1857; CHECK-NEXT:  // %bb.0:
1858; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_riri_param_0];
1859; CHECK-NEXT:    ld.param.f32 %f2, [st_param_v4_f32_riri_param_1];
1860; CHECK-NEXT:    { // callseq 77, 0
1861; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1862; CHECK-NEXT:    st.param.v4.f32 [param0], {%f1, 0f40000000, %f2, 0f40800000};
1863; CHECK-NEXT:    call.uni
1864; CHECK-NEXT:    call_v4_f32,
1865; CHECK-NEXT:    (
1866; CHECK-NEXT:    param0
1867; CHECK-NEXT:    );
1868; CHECK-NEXT:    } // callseq 77
1869; CHECK-NEXT:    ret;
1870  %struct.riri0 = insertvalue %struct.float4 poison, float %a, 0
1871  %struct.riri1 = insertvalue %struct.float4 %struct.riri0, float 2.0, 1
1872  %struct.riri2 = insertvalue %struct.float4 %struct.riri1, float %c, 2
1873  %struct.riri3 = insertvalue %struct.float4 %struct.riri2, float 4.0, 3
1874  call void @call_v4_f32(%struct.float4 %struct.riri3)
1875  ret void
1876}
1877define void @st_param_v4_f32_rrii(float %a, float %b) {
1878; CHECK-LABEL: st_param_v4_f32_rrii(
1879; CHECK:       {
1880; CHECK-NEXT:    .reg .f32 %f<3>;
1881; CHECK-EMPTY:
1882; CHECK-NEXT:  // %bb.0:
1883; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_rrii_param_0];
1884; CHECK-NEXT:    ld.param.f32 %f2, [st_param_v4_f32_rrii_param_1];
1885; CHECK-NEXT:    { // callseq 78, 0
1886; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1887; CHECK-NEXT:    st.param.v4.f32 [param0], {%f1, %f2, 0f40400000, 0f40800000};
1888; CHECK-NEXT:    call.uni
1889; CHECK-NEXT:    call_v4_f32,
1890; CHECK-NEXT:    (
1891; CHECK-NEXT:    param0
1892; CHECK-NEXT:    );
1893; CHECK-NEXT:    } // callseq 78
1894; CHECK-NEXT:    ret;
1895  %struct.rrii0 = insertvalue %struct.float4 poison, float %a, 0
1896  %struct.rrii1 = insertvalue %struct.float4 %struct.rrii0, float %b, 1
1897  %struct.rrii2 = insertvalue %struct.float4 %struct.rrii1, float 3.0, 2
1898  %struct.rrii3 = insertvalue %struct.float4 %struct.rrii2, float 4.0, 3
1899  call void @call_v4_f32(%struct.float4 %struct.rrii3)
1900  ret void
1901}
1902define void @st_param_v4_f32_iiir(float %d) {
1903; CHECK-LABEL: st_param_v4_f32_iiir(
1904; CHECK:       {
1905; CHECK-NEXT:    .reg .f32 %f<2>;
1906; CHECK-EMPTY:
1907; CHECK-NEXT:  // %bb.0:
1908; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_iiir_param_0];
1909; CHECK-NEXT:    { // callseq 79, 0
1910; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1911; CHECK-NEXT:    st.param.v4.f32 [param0], {0f3F800000, 0f40000000, 0f40400000, %f1};
1912; CHECK-NEXT:    call.uni
1913; CHECK-NEXT:    call_v4_f32,
1914; CHECK-NEXT:    (
1915; CHECK-NEXT:    param0
1916; CHECK-NEXT:    );
1917; CHECK-NEXT:    } // callseq 79
1918; CHECK-NEXT:    ret;
1919  %struct.iiir0 = insertvalue %struct.float4 poison, float 1.0, 0
1920  %struct.iiir1 = insertvalue %struct.float4 %struct.iiir0, float 2.0, 1
1921  %struct.iiir2 = insertvalue %struct.float4 %struct.iiir1, float 3.0, 2
1922  %struct.iiir3 = insertvalue %struct.float4 %struct.iiir2, float %d, 3
1923  call void @call_v4_f32(%struct.float4 %struct.iiir3)
1924  ret void
1925}
1926define void @st_param_v4_f32_iiri(float %c) {
1927; CHECK-LABEL: st_param_v4_f32_iiri(
1928; CHECK:       {
1929; CHECK-NEXT:    .reg .f32 %f<2>;
1930; CHECK-EMPTY:
1931; CHECK-NEXT:  // %bb.0:
1932; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_iiri_param_0];
1933; CHECK-NEXT:    { // callseq 80, 0
1934; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1935; CHECK-NEXT:    st.param.v4.f32 [param0], {0f3F800000, 0f40000000, %f1, 0f40800000};
1936; CHECK-NEXT:    call.uni
1937; CHECK-NEXT:    call_v4_f32,
1938; CHECK-NEXT:    (
1939; CHECK-NEXT:    param0
1940; CHECK-NEXT:    );
1941; CHECK-NEXT:    } // callseq 80
1942; CHECK-NEXT:    ret;
1943  %struct.iiri0 = insertvalue %struct.float4 poison, float 1.0, 0
1944  %struct.iiri1 = insertvalue %struct.float4 %struct.iiri0, float 2.0, 1
1945  %struct.iiri2 = insertvalue %struct.float4 %struct.iiri1, float %c, 2
1946  %struct.iiri3 = insertvalue %struct.float4 %struct.iiri2, float 4.0, 3
1947  call void @call_v4_f32(%struct.float4 %struct.iiri3)
1948  ret void
1949}
1950define void @st_param_v4_f32_irii(float %b) {
1951; CHECK-LABEL: st_param_v4_f32_irii(
1952; CHECK:       {
1953; CHECK-NEXT:    .reg .f32 %f<2>;
1954; CHECK-EMPTY:
1955; CHECK-NEXT:  // %bb.0:
1956; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_irii_param_0];
1957; CHECK-NEXT:    { // callseq 81, 0
1958; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1959; CHECK-NEXT:    st.param.v4.f32 [param0], {0f3F800000, %f1, 0f40400000, 0f40800000};
1960; CHECK-NEXT:    call.uni
1961; CHECK-NEXT:    call_v4_f32,
1962; CHECK-NEXT:    (
1963; CHECK-NEXT:    param0
1964; CHECK-NEXT:    );
1965; CHECK-NEXT:    } // callseq 81
1966; CHECK-NEXT:    ret;
1967  %struct.irii0 = insertvalue %struct.float4 poison, float 1.0, 0
1968  %struct.irii1 = insertvalue %struct.float4 %struct.irii0, float %b, 1
1969  %struct.irii2 = insertvalue %struct.float4 %struct.irii1, float 3.0, 2
1970  %struct.irii3 = insertvalue %struct.float4 %struct.irii2, float 4.0, 3
1971  call void @call_v4_f32(%struct.float4 %struct.irii3)
1972  ret void
1973}
1974define void @st_param_v4_f32_riii(float %a) {
1975; CHECK-LABEL: st_param_v4_f32_riii(
1976; CHECK:       {
1977; CHECK-NEXT:    .reg .f32 %f<2>;
1978; CHECK-EMPTY:
1979; CHECK-NEXT:  // %bb.0:
1980; CHECK-NEXT:    ld.param.f32 %f1, [st_param_v4_f32_riii_param_0];
1981; CHECK-NEXT:    { // callseq 82, 0
1982; CHECK-NEXT:    .param .align 16 .b8 param0[16];
1983; CHECK-NEXT:    st.param.v4.f32 [param0], {%f1, 0f40000000, 0f40400000, 0f40800000};
1984; CHECK-NEXT:    call.uni
1985; CHECK-NEXT:    call_v4_f32,
1986; CHECK-NEXT:    (
1987; CHECK-NEXT:    param0
1988; CHECK-NEXT:    );
1989; CHECK-NEXT:    } // callseq 82
1990; CHECK-NEXT:    ret;
1991  %struct.riii0 = insertvalue %struct.float4 poison, float %a, 0
1992  %struct.riii1 = insertvalue %struct.float4 %struct.riii0, float 2.0, 1
1993  %struct.riii2 = insertvalue %struct.float4 %struct.riii1, float 3.0, 2
1994  %struct.riii3 = insertvalue %struct.float4 %struct.riii2, float 4.0, 3
1995  call void @call_v4_f32(%struct.float4 %struct.riii3)
1996  ret void
1997}
1998
1999declare void @call_v4_i8(%struct.char4 alignstack(4))
2000declare void @call_v4_i16(%struct.short4 alignstack(8))
2001declare void @call_v4_i32(%struct.int4 alignstack(16))
2002declare void @call_v4_f32(%struct.float4 alignstack(16))
2003