1742 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 42; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s3; RUN: llc < %s -mtriple=nvptx | FileCheck %s4; RUN: %if ptxas-ptr32 %{ 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, 028; CHECK-NEXT: .param .align 2 .b8 param0[4];29; CHECK-NEXT: st.param.b16 [param0+2], 2;30; CHECK-NEXT: st.param.b8 [param0], 1;31; CHECK-NEXT: call.uni call_i8_i16, (param0);32; CHECK-NEXT: } // callseq 033; CHECK-NEXT: ret;34 call void @call_i8_i16(%struct.A { i8 1, i16 2 })35 ret void36}37 38define void @st_param_i32() {39; CHECK-LABEL: st_param_i32(40; CHECK: {41; CHECK-EMPTY:42; CHECK-EMPTY:43; CHECK-NEXT: // %bb.0:44; CHECK-NEXT: { // callseq 1, 045; CHECK-NEXT: .param .b32 param0;46; CHECK-NEXT: st.param.b32 [param0], 3;47; CHECK-NEXT: call.uni call_i32, (param0);48; CHECK-NEXT: } // callseq 149; CHECK-NEXT: ret;50 call void @call_i32(i32 3)51 ret void52}53 54define void @st_param_i64() {55; CHECK-LABEL: st_param_i64(56; CHECK: {57; CHECK-EMPTY:58; CHECK-EMPTY:59; CHECK-NEXT: // %bb.0:60; CHECK-NEXT: { // callseq 2, 061; CHECK-NEXT: .param .b64 param0;62; CHECK-NEXT: st.param.b64 [param0], 4;63; CHECK-NEXT: call.uni call_i64, (param0);64; CHECK-NEXT: } // callseq 265; CHECK-NEXT: ret;66 call void @call_i64(i64 4)67 ret void68}69 70define void @st_param_f32() {71; CHECK-LABEL: st_param_f32(72; CHECK: {73; CHECK-EMPTY:74; CHECK-EMPTY:75; CHECK-NEXT: // %bb.0:76; CHECK-NEXT: { // callseq 3, 077; CHECK-NEXT: .param .b32 param0;78; CHECK-NEXT: st.param.b32 [param0], 1084227584;79; CHECK-NEXT: call.uni call_f32, (param0);80; CHECK-NEXT: } // callseq 381; CHECK-NEXT: ret;82 call void @call_f32(float 5.0)83 ret void84}85 86define void @st_param_f64() {87; CHECK-LABEL: st_param_f64(88; CHECK: {89; CHECK-EMPTY:90; CHECK-EMPTY:91; CHECK-NEXT: // %bb.0:92; CHECK-NEXT: { // callseq 4, 093; CHECK-NEXT: .param .b64 param0;94; CHECK-NEXT: st.param.b64 [param0], 4618441417868443648;95; CHECK-NEXT: call.uni call_f64, (param0);96; CHECK-NEXT: } // callseq 497; CHECK-NEXT: ret;98 call void @call_f64(double 6.0)99 ret void100}101 102declare void @call_i8_i16(%struct.A)103declare void @call_i32(i32)104declare void @call_i64(i64)105declare void @call_f32(float)106declare void @call_f64(double)107 108define void @st_param_v2_i8_ii() {109; CHECK-LABEL: st_param_v2_i8_ii(110; CHECK: {111; CHECK-EMPTY:112; CHECK-EMPTY:113; CHECK-NEXT: // %bb.0:114; CHECK-NEXT: { // callseq 5, 0115; CHECK-NEXT: .param .align 2 .b8 param0[2];116; CHECK-NEXT: st.param.v2.b8 [param0], {1, 2};117; CHECK-NEXT: call.uni call_v2_i8, (param0);118; CHECK-NEXT: } // callseq 5119; CHECK-NEXT: ret;120 call void @call_v2_i8(%struct.char2 { i8 1, i8 2 })121 ret void122}123define void @st_param_v2_i8_ir(i8 %val) {124; CHECK-LABEL: st_param_v2_i8_ir(125; CHECK: {126; CHECK-NEXT: .reg .b16 %rs<2>;127; CHECK-EMPTY:128; CHECK-NEXT: // %bb.0:129; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v2_i8_ir_param_0];130; CHECK-NEXT: { // callseq 6, 0131; CHECK-NEXT: .param .align 2 .b8 param0[2];132; CHECK-NEXT: st.param.v2.b8 [param0], {1, %rs1};133; CHECK-NEXT: call.uni call_v2_i8, (param0);134; CHECK-NEXT: } // callseq 6135; CHECK-NEXT: ret;136 %struct.ir0 = insertvalue %struct.char2 poison, i8 1, 0137 %struct.ir1 = insertvalue %struct.char2 %struct.ir0, i8 %val, 1138 call void @call_v2_i8(%struct.char2 %struct.ir1)139 ret void140}141define void @st_param_v2_i8_ri(i8 %val) {142; CHECK-LABEL: st_param_v2_i8_ri(143; CHECK: {144; CHECK-NEXT: .reg .b16 %rs<2>;145; CHECK-EMPTY:146; CHECK-NEXT: // %bb.0:147; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v2_i8_ri_param_0];148; CHECK-NEXT: { // callseq 7, 0149; CHECK-NEXT: .param .align 2 .b8 param0[2];150; CHECK-NEXT: st.param.v2.b8 [param0], {%rs1, 2};151; CHECK-NEXT: call.uni call_v2_i8, (param0);152; CHECK-NEXT: } // callseq 7153; CHECK-NEXT: ret;154 %struct.ri0 = insertvalue %struct.char2 poison, i8 %val, 0155 %struct.ri1 = insertvalue %struct.char2 %struct.ri0, i8 2, 1156 call void @call_v2_i8(%struct.char2 %struct.ri1)157 ret void158}159 160define void @st_param_v2_i16_ii() {161; CHECK-LABEL: st_param_v2_i16_ii(162; CHECK: {163; CHECK-EMPTY:164; CHECK-EMPTY:165; CHECK-NEXT: // %bb.0:166; CHECK-NEXT: { // callseq 8, 0167; CHECK-NEXT: .param .align 4 .b8 param0[4];168; CHECK-NEXT: st.param.b32 [param0], 131073;169; CHECK-NEXT: call.uni call_v2_i16, (param0);170; CHECK-NEXT: } // callseq 8171; CHECK-NEXT: ret;172 call void @call_v2_i16(%struct.short2 { i16 1, i16 2 })173 ret void174}175define void @st_param_v2_i16_ir(i16 %val) {176; CHECK-LABEL: st_param_v2_i16_ir(177; CHECK: {178; CHECK-NEXT: .reg .b16 %rs<2>;179; CHECK-EMPTY:180; CHECK-NEXT: // %bb.0:181; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v2_i16_ir_param_0];182; CHECK-NEXT: { // callseq 9, 0183; CHECK-NEXT: .param .align 4 .b8 param0[4];184; CHECK-NEXT: st.param.v2.b16 [param0], {1, %rs1};185; CHECK-NEXT: call.uni call_v2_i16, (param0);186; CHECK-NEXT: } // callseq 9187; CHECK-NEXT: ret;188 %struct.ir0 = insertvalue %struct.short2 poison, i16 1, 0189 %struct.ir1 = insertvalue %struct.short2 %struct.ir0, i16 %val, 1190 call void @call_v2_i16(%struct.short2 %struct.ir1)191 ret void192}193define void @st_param_v2_i16_ri(i16 %val) {194; CHECK-LABEL: st_param_v2_i16_ri(195; CHECK: {196; CHECK-NEXT: .reg .b16 %rs<2>;197; CHECK-EMPTY:198; CHECK-NEXT: // %bb.0:199; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v2_i16_ri_param_0];200; CHECK-NEXT: { // callseq 10, 0201; CHECK-NEXT: .param .align 4 .b8 param0[4];202; CHECK-NEXT: st.param.v2.b16 [param0], {%rs1, 2};203; CHECK-NEXT: call.uni call_v2_i16, (param0);204; CHECK-NEXT: } // callseq 10205; CHECK-NEXT: ret;206 %struct.ri0 = insertvalue %struct.short2 poison, i16 %val, 0207 %struct.ri1 = insertvalue %struct.short2 %struct.ri0, i16 2, 1208 call void @call_v2_i16(%struct.short2 %struct.ri1)209 ret void210}211 212define void @st_param_v2_i32_ii() {213; CHECK-LABEL: st_param_v2_i32_ii(214; CHECK: {215; CHECK-EMPTY:216; CHECK-EMPTY:217; CHECK-NEXT: // %bb.0:218; CHECK-NEXT: { // callseq 11, 0219; CHECK-NEXT: .param .align 8 .b8 param0[8];220; CHECK-NEXT: st.param.v2.b32 [param0], {1, 2};221; CHECK-NEXT: call.uni call_v2_i32, (param0);222; CHECK-NEXT: } // callseq 11223; CHECK-NEXT: ret;224 call void @call_v2_i32(%struct.int2 { i32 1, i32 2 })225 ret void226}227define void @st_param_v2_i32_ir(i32 %val) {228; CHECK-LABEL: st_param_v2_i32_ir(229; CHECK: {230; CHECK-NEXT: .reg .b32 %r<2>;231; CHECK-EMPTY:232; CHECK-NEXT: // %bb.0:233; CHECK-NEXT: ld.param.b32 %r1, [st_param_v2_i32_ir_param_0];234; CHECK-NEXT: { // callseq 12, 0235; CHECK-NEXT: .param .align 8 .b8 param0[8];236; CHECK-NEXT: st.param.v2.b32 [param0], {1, %r1};237; CHECK-NEXT: call.uni call_v2_i32, (param0);238; CHECK-NEXT: } // callseq 12239; CHECK-NEXT: ret;240 %struct.ir0 = insertvalue %struct.int2 poison, i32 1, 0241 %struct.ir1 = insertvalue %struct.int2 %struct.ir0, i32 %val, 1242 call void @call_v2_i32(%struct.int2 %struct.ir1)243 ret void244}245define void @st_param_v2_i32_ri(i32 %val) {246; CHECK-LABEL: st_param_v2_i32_ri(247; CHECK: {248; CHECK-NEXT: .reg .b32 %r<2>;249; CHECK-EMPTY:250; CHECK-NEXT: // %bb.0:251; CHECK-NEXT: ld.param.b32 %r1, [st_param_v2_i32_ri_param_0];252; CHECK-NEXT: { // callseq 13, 0253; CHECK-NEXT: .param .align 8 .b8 param0[8];254; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, 2};255; CHECK-NEXT: call.uni call_v2_i32, (param0);256; CHECK-NEXT: } // callseq 13257; CHECK-NEXT: ret;258 %struct.ri0 = insertvalue %struct.int2 poison, i32 %val, 0259 %struct.ri1 = insertvalue %struct.int2 %struct.ri0, i32 2, 1260 call void @call_v2_i32(%struct.int2 %struct.ri1)261 ret void262}263 264define void @st_param_v2_i64_ii() {265; CHECK-LABEL: st_param_v2_i64_ii(266; CHECK: {267; CHECK-EMPTY:268; CHECK-EMPTY:269; CHECK-NEXT: // %bb.0:270; CHECK-NEXT: { // callseq 14, 0271; CHECK-NEXT: .param .align 16 .b8 param0[16];272; CHECK-NEXT: st.param.v2.b64 [param0], {1, 2};273; CHECK-NEXT: call.uni call_v2_i64, (param0);274; CHECK-NEXT: } // callseq 14275; CHECK-NEXT: ret;276 call void @call_v2_i64(%struct.longlong2 { i64 1, i64 2 })277 ret void278}279define void @st_param_v2_i64_ir(i64 %val) {280; CHECK-LABEL: st_param_v2_i64_ir(281; CHECK: {282; CHECK-NEXT: .reg .b64 %rd<2>;283; CHECK-EMPTY:284; CHECK-NEXT: // %bb.0:285; CHECK-NEXT: ld.param.b64 %rd1, [st_param_v2_i64_ir_param_0];286; CHECK-NEXT: { // callseq 15, 0287; CHECK-NEXT: .param .align 16 .b8 param0[16];288; CHECK-NEXT: st.param.v2.b64 [param0], {1, %rd1};289; CHECK-NEXT: call.uni call_v2_i64, (param0);290; CHECK-NEXT: } // callseq 15291; CHECK-NEXT: ret;292 %struct.ir0 = insertvalue %struct.longlong2 poison, i64 1, 0293 %struct.ir1 = insertvalue %struct.longlong2 %struct.ir0, i64 %val, 1294 call void @call_v2_i64(%struct.longlong2 %struct.ir1)295 ret void296}297define void @st_param_v2_i64_ri(i64 %val) {298; CHECK-LABEL: st_param_v2_i64_ri(299; CHECK: {300; CHECK-NEXT: .reg .b64 %rd<2>;301; CHECK-EMPTY:302; CHECK-NEXT: // %bb.0:303; CHECK-NEXT: ld.param.b64 %rd1, [st_param_v2_i64_ri_param_0];304; CHECK-NEXT: { // callseq 16, 0305; CHECK-NEXT: .param .align 16 .b8 param0[16];306; CHECK-NEXT: st.param.v2.b64 [param0], {%rd1, 2};307; CHECK-NEXT: call.uni call_v2_i64, (param0);308; CHECK-NEXT: } // callseq 16309; CHECK-NEXT: ret;310 %struct.ri0 = insertvalue %struct.longlong2 poison, i64 %val, 0311 %struct.ri1 = insertvalue %struct.longlong2 %struct.ri0, i64 2, 1312 call void @call_v2_i64(%struct.longlong2 %struct.ri1)313 ret void314}315 316define void @st_param_v2_f32_ii(float %val) {317; CHECK-LABEL: st_param_v2_f32_ii(318; CHECK: {319; CHECK-EMPTY:320; CHECK-EMPTY:321; CHECK-NEXT: // %bb.0:322; CHECK-NEXT: { // callseq 17, 0323; CHECK-NEXT: .param .align 8 .b8 param0[8];324; CHECK-NEXT: st.param.v2.b32 [param0], {0f3F800000, 0f40000000};325; CHECK-NEXT: call.uni call_v2_f32, (param0);326; CHECK-NEXT: } // callseq 17327; CHECK-NEXT: ret;328 call void @call_v2_f32(%struct.float2 { float 1.0, float 2.0 })329 ret void330}331define void @st_param_v2_f32_ir(float %val) {332; CHECK-LABEL: st_param_v2_f32_ir(333; CHECK: {334; CHECK-NEXT: .reg .b32 %r<2>;335; CHECK-EMPTY:336; CHECK-NEXT: // %bb.0:337; CHECK-NEXT: ld.param.b32 %r1, [st_param_v2_f32_ir_param_0];338; CHECK-NEXT: { // callseq 18, 0339; CHECK-NEXT: .param .align 8 .b8 param0[8];340; CHECK-NEXT: st.param.v2.b32 [param0], {0f3F800000, %r1};341; CHECK-NEXT: call.uni call_v2_f32, (param0);342; CHECK-NEXT: } // callseq 18343; CHECK-NEXT: ret;344 %struct.ir0 = insertvalue %struct.float2 poison, float 1.0, 0345 %struct.ir1 = insertvalue %struct.float2 %struct.ir0, float %val, 1346 call void @call_v2_f32(%struct.float2 %struct.ir1)347 ret void348}349define void @st_param_v2_f32_ri(float %val) {350; CHECK-LABEL: st_param_v2_f32_ri(351; CHECK: {352; CHECK-NEXT: .reg .b32 %r<2>;353; CHECK-EMPTY:354; CHECK-NEXT: // %bb.0:355; CHECK-NEXT: ld.param.b32 %r1, [st_param_v2_f32_ri_param_0];356; CHECK-NEXT: { // callseq 19, 0357; CHECK-NEXT: .param .align 8 .b8 param0[8];358; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, 0f40000000};359; CHECK-NEXT: call.uni call_v2_f32, (param0);360; CHECK-NEXT: } // callseq 19361; CHECK-NEXT: ret;362 %struct.ri0 = insertvalue %struct.float2 poison, float %val, 0363 %struct.ri1 = insertvalue %struct.float2 %struct.ri0, float 2.0, 1364 call void @call_v2_f32(%struct.float2 %struct.ri1)365 ret void366}367 368define void @st_param_v2_f64_ii(double %val) {369; CHECK-LABEL: st_param_v2_f64_ii(370; CHECK: {371; CHECK-EMPTY:372; CHECK-EMPTY:373; CHECK-NEXT: // %bb.0:374; CHECK-NEXT: { // callseq 20, 0375; CHECK-NEXT: .param .align 16 .b8 param0[16];376; CHECK-NEXT: st.param.v2.b64 [param0], {0d3FF0000000000000, 0d4000000000000000};377; CHECK-NEXT: call.uni call_v2_f64, (param0);378; CHECK-NEXT: } // callseq 20379; CHECK-NEXT: ret;380 call void @call_v2_f64(%struct.double2 { double 1.0, double 2.0 })381 ret void382}383define void @st_param_v2_f64_ir(double %val) {384; CHECK-LABEL: st_param_v2_f64_ir(385; CHECK: {386; CHECK-NEXT: .reg .b64 %rd<2>;387; CHECK-EMPTY:388; CHECK-NEXT: // %bb.0:389; CHECK-NEXT: ld.param.b64 %rd1, [st_param_v2_f64_ir_param_0];390; CHECK-NEXT: { // callseq 21, 0391; CHECK-NEXT: .param .align 16 .b8 param0[16];392; CHECK-NEXT: st.param.v2.b64 [param0], {0d3FF0000000000000, %rd1};393; CHECK-NEXT: call.uni call_v2_f64, (param0);394; CHECK-NEXT: } // callseq 21395; CHECK-NEXT: ret;396 %struct.ir0 = insertvalue %struct.double2 poison, double 1.0, 0397 %struct.ir1 = insertvalue %struct.double2 %struct.ir0, double %val, 1398 call void @call_v2_f64(%struct.double2 %struct.ir1)399 ret void400}401define void @st_param_v2_f64_ri(double %val) {402; CHECK-LABEL: st_param_v2_f64_ri(403; CHECK: {404; CHECK-NEXT: .reg .b64 %rd<2>;405; CHECK-EMPTY:406; CHECK-NEXT: // %bb.0:407; CHECK-NEXT: ld.param.b64 %rd1, [st_param_v2_f64_ri_param_0];408; CHECK-NEXT: { // callseq 22, 0409; CHECK-NEXT: .param .align 16 .b8 param0[16];410; CHECK-NEXT: st.param.v2.b64 [param0], {%rd1, 0d4000000000000000};411; CHECK-NEXT: call.uni call_v2_f64, (param0);412; CHECK-NEXT: } // callseq 22413; CHECK-NEXT: ret;414 %struct.ri0 = insertvalue %struct.double2 poison, double %val, 0415 %struct.ri1 = insertvalue %struct.double2 %struct.ri0, double 2.0, 1416 call void @call_v2_f64(%struct.double2 %struct.ri1)417 ret void418}419 420declare void @call_v2_i8(%struct.char2 alignstack(2))421declare void @call_v2_i16(%struct.short2 alignstack(4))422declare void @call_v2_i32(%struct.int2 alignstack(8))423declare void @call_v2_i64(%struct.longlong2 alignstack(16))424declare void @call_v2_f32(%struct.float2 alignstack(8))425declare void @call_v2_f64(%struct.double2 alignstack(16))426 427define void @st_param_v4_i8_iiii() {428; CHECK-LABEL: st_param_v4_i8_iiii(429; CHECK: {430; CHECK-EMPTY:431; CHECK-EMPTY:432; CHECK-NEXT: // %bb.0:433; CHECK-NEXT: { // callseq 23, 0434; CHECK-NEXT: .param .align 4 .b8 param0[4];435; CHECK-NEXT: st.param.b32 [param0], 67305985;436; CHECK-NEXT: call.uni call_v4_i8, (param0);437; CHECK-NEXT: } // callseq 23438; CHECK-NEXT: ret;439 call void @call_v4_i8(%struct.char4 { i8 1, i8 2, i8 3, i8 4 })440 ret void441}442define void @st_param_v4_i8_irrr(i8 %b, i8 %c, i8 %d) {443; CHECK-LABEL: st_param_v4_i8_irrr(444; CHECK: {445; CHECK-NEXT: .reg .b32 %r<7>;446; CHECK-EMPTY:447; CHECK-NEXT: // %bb.0:448; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_irrr_param_2];449; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_irrr_param_1];450; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U;451; CHECK-NEXT: ld.param.b8 %r4, [st_param_v4_i8_irrr_param_0];452; CHECK-NEXT: prmt.b32 %r5, 1, %r4, 0x3340U;453; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x5410U;454; CHECK-NEXT: { // callseq 24, 0455; CHECK-NEXT: .param .align 4 .b8 param0[4];456; CHECK-NEXT: st.param.b32 [param0], %r6;457; CHECK-NEXT: call.uni call_v4_i8, (param0);458; CHECK-NEXT: } // callseq 24459; CHECK-NEXT: ret;460 %struct.irrr0 = insertvalue %struct.char4 poison, i8 1, 0461 %struct.irrr1 = insertvalue %struct.char4 %struct.irrr0, i8 %b, 1462 %struct.irrr2 = insertvalue %struct.char4 %struct.irrr1, i8 %c, 2463 %struct.irrr3 = insertvalue %struct.char4 %struct.irrr2, i8 %d, 3464 call void @call_v4_i8(%struct.char4 %struct.irrr3)465 ret void466}467define void @st_param_v4_i8_rirr(i8 %a, i8 %c, i8 %d) {468; CHECK-LABEL: st_param_v4_i8_rirr(469; CHECK: {470; CHECK-NEXT: .reg .b32 %r<7>;471; CHECK-EMPTY:472; CHECK-NEXT: // %bb.0:473; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_rirr_param_2];474; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_rirr_param_1];475; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U;476; CHECK-NEXT: ld.param.b8 %r4, [st_param_v4_i8_rirr_param_0];477; CHECK-NEXT: prmt.b32 %r5, %r4, 2, 0x3340U;478; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x5410U;479; CHECK-NEXT: { // callseq 25, 0480; CHECK-NEXT: .param .align 4 .b8 param0[4];481; CHECK-NEXT: st.param.b32 [param0], %r6;482; CHECK-NEXT: call.uni call_v4_i8, (param0);483; CHECK-NEXT: } // callseq 25484; CHECK-NEXT: ret;485 %struct.rirr0 = insertvalue %struct.char4 poison, i8 %a, 0486 %struct.rirr1 = insertvalue %struct.char4 %struct.rirr0, i8 2, 1487 %struct.rirr2 = insertvalue %struct.char4 %struct.rirr1, i8 %c, 2488 %struct.rirr3 = insertvalue %struct.char4 %struct.rirr2, i8 %d, 3489 call void @call_v4_i8(%struct.char4 %struct.rirr3)490 ret void491}492define void @st_param_v4_i8_rrir(i8 %a, i8 %b, i8 %d) {493; CHECK-LABEL: st_param_v4_i8_rrir(494; CHECK: {495; CHECK-NEXT: .reg .b32 %r<7>;496; CHECK-EMPTY:497; CHECK-NEXT: // %bb.0:498; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_rrir_param_1];499; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_rrir_param_0];500; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U;501; CHECK-NEXT: ld.param.b8 %r4, [st_param_v4_i8_rrir_param_2];502; CHECK-NEXT: prmt.b32 %r5, 3, %r4, 0x3340U;503; CHECK-NEXT: prmt.b32 %r6, %r3, %r5, 0x5410U;504; CHECK-NEXT: { // callseq 26, 0505; CHECK-NEXT: .param .align 4 .b8 param0[4];506; CHECK-NEXT: st.param.b32 [param0], %r6;507; CHECK-NEXT: call.uni call_v4_i8, (param0);508; CHECK-NEXT: } // callseq 26509; CHECK-NEXT: ret;510 %struct.rrir0 = insertvalue %struct.char4 poison, i8 %a, 0511 %struct.rrir1 = insertvalue %struct.char4 %struct.rrir0, i8 %b, 1512 %struct.rrir2 = insertvalue %struct.char4 %struct.rrir1, i8 3, 2513 %struct.rrir3 = insertvalue %struct.char4 %struct.rrir2, i8 %d, 3514 call void @call_v4_i8(%struct.char4 %struct.rrir3)515 ret void516}517define void @st_param_v4_i8_rrri(i8 %a, i8 %b, i8 %c) {518; CHECK-LABEL: st_param_v4_i8_rrri(519; CHECK: {520; CHECK-NEXT: .reg .b32 %r<7>;521; CHECK-EMPTY:522; CHECK-NEXT: // %bb.0:523; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_rrri_param_1];524; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_rrri_param_0];525; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U;526; CHECK-NEXT: ld.param.b8 %r4, [st_param_v4_i8_rrri_param_2];527; CHECK-NEXT: prmt.b32 %r5, %r4, 4, 0x3340U;528; CHECK-NEXT: prmt.b32 %r6, %r3, %r5, 0x5410U;529; CHECK-NEXT: { // callseq 27, 0530; CHECK-NEXT: .param .align 4 .b8 param0[4];531; CHECK-NEXT: st.param.b32 [param0], %r6;532; CHECK-NEXT: call.uni call_v4_i8, (param0);533; CHECK-NEXT: } // callseq 27534; CHECK-NEXT: ret;535 %struct.rrri0 = insertvalue %struct.char4 poison, i8 %a, 0536 %struct.rrri1 = insertvalue %struct.char4 %struct.rrri0, i8 %b, 1537 %struct.rrri2 = insertvalue %struct.char4 %struct.rrri1, i8 %c, 2538 %struct.rrri3 = insertvalue %struct.char4 %struct.rrri2, i8 4, 3539 call void @call_v4_i8(%struct.char4 %struct.rrri3)540 ret void541}542define void @st_param_v4_i8_iirr(i8 %c, i8 %d) {543; CHECK-LABEL: st_param_v4_i8_iirr(544; CHECK: {545; CHECK-NEXT: .reg .b32 %r<5>;546; CHECK-EMPTY:547; CHECK-NEXT: // %bb.0:548; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_iirr_param_1];549; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_iirr_param_0];550; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U;551; CHECK-NEXT: prmt.b32 %r4, 513, %r3, 0x5410U;552; CHECK-NEXT: { // callseq 28, 0553; CHECK-NEXT: .param .align 4 .b8 param0[4];554; CHECK-NEXT: st.param.b32 [param0], %r4;555; CHECK-NEXT: call.uni call_v4_i8, (param0);556; CHECK-NEXT: } // callseq 28557; CHECK-NEXT: ret;558 %struct.iirr0 = insertvalue %struct.char4 poison, i8 1, 0559 %struct.iirr1 = insertvalue %struct.char4 %struct.iirr0, i8 2, 1560 %struct.iirr2 = insertvalue %struct.char4 %struct.iirr1, i8 %c, 2561 %struct.iirr3 = insertvalue %struct.char4 %struct.iirr2, i8 %d, 3562 call void @call_v4_i8(%struct.char4 %struct.iirr3)563 ret void564}565define void @st_param_v4_i8_irir(i8 %b, i8 %d) {566; CHECK-LABEL: st_param_v4_i8_irir(567; CHECK: {568; CHECK-NEXT: .reg .b32 %r<6>;569; CHECK-EMPTY:570; CHECK-NEXT: // %bb.0:571; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_irir_param_1];572; CHECK-NEXT: prmt.b32 %r2, 3, %r1, 0x3340U;573; CHECK-NEXT: ld.param.b8 %r3, [st_param_v4_i8_irir_param_0];574; CHECK-NEXT: prmt.b32 %r4, 1, %r3, 0x3340U;575; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;576; CHECK-NEXT: { // callseq 29, 0577; CHECK-NEXT: .param .align 4 .b8 param0[4];578; CHECK-NEXT: st.param.b32 [param0], %r5;579; CHECK-NEXT: call.uni call_v4_i8, (param0);580; CHECK-NEXT: } // callseq 29581; CHECK-NEXT: ret;582 %struct.irir0 = insertvalue %struct.char4 poison, i8 1, 0583 %struct.irir1 = insertvalue %struct.char4 %struct.irir0, i8 %b, 1584 %struct.irir2 = insertvalue %struct.char4 %struct.irir1, i8 3, 2585 %struct.irir3 = insertvalue %struct.char4 %struct.irir2, i8 %d, 3586 call void @call_v4_i8(%struct.char4 %struct.irir3)587 ret void588}589define void @st_param_v4_i8_irri(i8 %b, i8 %c) {590; CHECK-LABEL: st_param_v4_i8_irri(591; CHECK: {592; CHECK-NEXT: .reg .b32 %r<6>;593; CHECK-EMPTY:594; CHECK-NEXT: // %bb.0:595; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_irri_param_1];596; CHECK-NEXT: prmt.b32 %r2, %r1, 4, 0x3340U;597; CHECK-NEXT: ld.param.b8 %r3, [st_param_v4_i8_irri_param_0];598; CHECK-NEXT: prmt.b32 %r4, 1, %r3, 0x3340U;599; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;600; CHECK-NEXT: { // callseq 30, 0601; CHECK-NEXT: .param .align 4 .b8 param0[4];602; CHECK-NEXT: st.param.b32 [param0], %r5;603; CHECK-NEXT: call.uni call_v4_i8, (param0);604; CHECK-NEXT: } // callseq 30605; CHECK-NEXT: ret;606 %struct.irri0 = insertvalue %struct.char4 poison, i8 1, 0607 %struct.irri1 = insertvalue %struct.char4 %struct.irri0, i8 %b, 1608 %struct.irri2 = insertvalue %struct.char4 %struct.irri1, i8 %c, 2609 %struct.irri3 = insertvalue %struct.char4 %struct.irri2, i8 4, 3610 call void @call_v4_i8(%struct.char4 %struct.irri3)611 ret void612}613define void @st_param_v4_i8_riir(i8 %a, i8 %d) {614; CHECK-LABEL: st_param_v4_i8_riir(615; CHECK: {616; CHECK-NEXT: .reg .b32 %r<6>;617; CHECK-EMPTY:618; CHECK-NEXT: // %bb.0:619; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_riir_param_1];620; CHECK-NEXT: prmt.b32 %r2, 3, %r1, 0x3340U;621; CHECK-NEXT: ld.param.b8 %r3, [st_param_v4_i8_riir_param_0];622; CHECK-NEXT: prmt.b32 %r4, %r3, 2, 0x3340U;623; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;624; CHECK-NEXT: { // callseq 31, 0625; CHECK-NEXT: .param .align 4 .b8 param0[4];626; CHECK-NEXT: st.param.b32 [param0], %r5;627; CHECK-NEXT: call.uni call_v4_i8, (param0);628; CHECK-NEXT: } // callseq 31629; CHECK-NEXT: ret;630 %struct.riir0 = insertvalue %struct.char4 poison, i8 %a, 0631 %struct.riir1 = insertvalue %struct.char4 %struct.riir0, i8 2, 1632 %struct.riir2 = insertvalue %struct.char4 %struct.riir1, i8 3, 2633 %struct.riir3 = insertvalue %struct.char4 %struct.riir2, i8 %d, 3634 call void @call_v4_i8(%struct.char4 %struct.riir3)635 ret void636}637define void @st_param_v4_i8_riri(i8 %a, i8 %c) {638; CHECK-LABEL: st_param_v4_i8_riri(639; CHECK: {640; CHECK-NEXT: .reg .b32 %r<6>;641; CHECK-EMPTY:642; CHECK-NEXT: // %bb.0:643; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_riri_param_1];644; CHECK-NEXT: prmt.b32 %r2, %r1, 4, 0x3340U;645; CHECK-NEXT: ld.param.b8 %r3, [st_param_v4_i8_riri_param_0];646; CHECK-NEXT: prmt.b32 %r4, %r3, 2, 0x3340U;647; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;648; CHECK-NEXT: { // callseq 32, 0649; CHECK-NEXT: .param .align 4 .b8 param0[4];650; CHECK-NEXT: st.param.b32 [param0], %r5;651; CHECK-NEXT: call.uni call_v4_i8, (param0);652; CHECK-NEXT: } // callseq 32653; CHECK-NEXT: ret;654 %struct.riri0 = insertvalue %struct.char4 poison, i8 %a, 0655 %struct.riri1 = insertvalue %struct.char4 %struct.riri0, i8 2, 1656 %struct.riri2 = insertvalue %struct.char4 %struct.riri1, i8 %c, 2657 %struct.riri3 = insertvalue %struct.char4 %struct.riri2, i8 4, 3658 call void @call_v4_i8(%struct.char4 %struct.riri3)659 ret void660}661define void @st_param_v4_i8_rrii(i8 %a, i8 %b) {662; CHECK-LABEL: st_param_v4_i8_rrii(663; CHECK: {664; CHECK-NEXT: .reg .b32 %r<5>;665; CHECK-EMPTY:666; CHECK-NEXT: // %bb.0:667; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_rrii_param_1];668; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_rrii_param_0];669; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U;670; CHECK-NEXT: prmt.b32 %r4, %r3, 1027, 0x5410U;671; CHECK-NEXT: { // callseq 33, 0672; CHECK-NEXT: .param .align 4 .b8 param0[4];673; CHECK-NEXT: st.param.b32 [param0], %r4;674; CHECK-NEXT: call.uni call_v4_i8, (param0);675; CHECK-NEXT: } // callseq 33676; CHECK-NEXT: ret;677 %struct.rrii0 = insertvalue %struct.char4 poison, i8 %a, 0678 %struct.rrii1 = insertvalue %struct.char4 %struct.rrii0, i8 %b, 1679 %struct.rrii2 = insertvalue %struct.char4 %struct.rrii1, i8 3, 2680 %struct.rrii3 = insertvalue %struct.char4 %struct.rrii2, i8 4, 3681 call void @call_v4_i8(%struct.char4 %struct.rrii3)682 ret void683}684define void @st_param_v4_i8_iiir(i8 %d) {685; CHECK-LABEL: st_param_v4_i8_iiir(686; CHECK: {687; CHECK-NEXT: .reg .b32 %r<4>;688; CHECK-EMPTY:689; CHECK-NEXT: // %bb.0:690; CHECK-NEXT: { // callseq 34, 0691; CHECK-NEXT: .param .align 4 .b8 param0[4];692; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_iiir_param_0];693; CHECK-NEXT: prmt.b32 %r2, 3, %r1, 0x3340U;694; CHECK-NEXT: prmt.b32 %r3, 513, %r2, 0x5410U;695; CHECK-NEXT: st.param.b32 [param0], %r3;696; CHECK-NEXT: call.uni call_v4_i8, (param0);697; CHECK-NEXT: } // callseq 34698; CHECK-NEXT: ret;699 %struct.iiir0 = insertvalue %struct.char4 poison, i8 1, 0700 %struct.iiir1 = insertvalue %struct.char4 %struct.iiir0, i8 2, 1701 %struct.iiir2 = insertvalue %struct.char4 %struct.iiir1, i8 3, 2702 %struct.iiir3 = insertvalue %struct.char4 %struct.iiir2, i8 %d, 3703 call void @call_v4_i8(%struct.char4 %struct.iiir3)704 ret void705}706define void @st_param_v4_i8_iiri(i8 %c) {707; CHECK-LABEL: st_param_v4_i8_iiri(708; CHECK: {709; CHECK-NEXT: .reg .b32 %r<4>;710; CHECK-EMPTY:711; CHECK-NEXT: // %bb.0:712; CHECK-NEXT: { // callseq 35, 0713; CHECK-NEXT: .param .align 4 .b8 param0[4];714; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_iiri_param_0];715; CHECK-NEXT: prmt.b32 %r2, %r1, 4, 0x3340U;716; CHECK-NEXT: prmt.b32 %r3, 513, %r2, 0x5410U;717; CHECK-NEXT: st.param.b32 [param0], %r3;718; CHECK-NEXT: call.uni call_v4_i8, (param0);719; CHECK-NEXT: } // callseq 35720; CHECK-NEXT: ret;721 %struct.iiri0 = insertvalue %struct.char4 poison, i8 1, 0722 %struct.iiri1 = insertvalue %struct.char4 %struct.iiri0, i8 2, 1723 %struct.iiri2 = insertvalue %struct.char4 %struct.iiri1, i8 %c, 2724 %struct.iiri3 = insertvalue %struct.char4 %struct.iiri2, i8 4, 3725 call void @call_v4_i8(%struct.char4 %struct.iiri3)726 ret void727}728define void @st_param_v4_i8_irii(i8 %b) {729; CHECK-LABEL: st_param_v4_i8_irii(730; CHECK: {731; CHECK-NEXT: .reg .b32 %r<4>;732; CHECK-EMPTY:733; CHECK-NEXT: // %bb.0:734; CHECK-NEXT: { // callseq 36, 0735; CHECK-NEXT: .param .align 4 .b8 param0[4];736; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_irii_param_0];737; CHECK-NEXT: prmt.b32 %r2, 1, %r1, 0x3340U;738; CHECK-NEXT: prmt.b32 %r3, %r2, 1027, 0x5410U;739; CHECK-NEXT: st.param.b32 [param0], %r3;740; CHECK-NEXT: call.uni call_v4_i8, (param0);741; CHECK-NEXT: } // callseq 36742; CHECK-NEXT: ret;743 %struct.irii0 = insertvalue %struct.char4 poison, i8 1, 0744 %struct.irii1 = insertvalue %struct.char4 %struct.irii0, i8 %b, 1745 %struct.irii2 = insertvalue %struct.char4 %struct.irii1, i8 3, 2746 %struct.irii3 = insertvalue %struct.char4 %struct.irii2, i8 4, 3747 call void @call_v4_i8(%struct.char4 %struct.irii3)748 ret void749}750define void @st_param_v4_i8_riii(i8 %a) {751; CHECK-LABEL: st_param_v4_i8_riii(752; CHECK: {753; CHECK-NEXT: .reg .b32 %r<4>;754; CHECK-EMPTY:755; CHECK-NEXT: // %bb.0:756; CHECK-NEXT: { // callseq 37, 0757; CHECK-NEXT: .param .align 4 .b8 param0[4];758; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_riii_param_0];759; CHECK-NEXT: prmt.b32 %r2, %r1, 2, 0x3340U;760; CHECK-NEXT: prmt.b32 %r3, %r2, 1027, 0x5410U;761; CHECK-NEXT: st.param.b32 [param0], %r3;762; CHECK-NEXT: call.uni call_v4_i8, (param0);763; CHECK-NEXT: } // callseq 37764; CHECK-NEXT: ret;765 %struct.riii0 = insertvalue %struct.char4 poison, i8 %a, 0766 %struct.riii1 = insertvalue %struct.char4 %struct.riii0, i8 2, 1767 %struct.riii2 = insertvalue %struct.char4 %struct.riii1, i8 3, 2768 %struct.riii3 = insertvalue %struct.char4 %struct.riii2, i8 4, 3769 call void @call_v4_i8(%struct.char4 %struct.riii3)770 ret void771}772 773define void @st_param_v4_i16_iiii() {774; CHECK-LABEL: st_param_v4_i16_iiii(775; CHECK: {776; CHECK-EMPTY:777; CHECK-EMPTY:778; CHECK-NEXT: // %bb.0:779; CHECK-NEXT: { // callseq 38, 0780; CHECK-NEXT: .param .align 8 .b8 param0[8];781; CHECK-NEXT: st.param.v2.b32 [param0], {131073, 262147};782; CHECK-NEXT: call.uni call_v4_i16, (param0);783; CHECK-NEXT: } // callseq 38784; CHECK-NEXT: ret;785 call void @call_v4_i16(%struct.short4 { i16 1, i16 2, i16 3, i16 4 })786 ret void787}788define void @st_param_v4_i16_irrr(i16 %b, i16 %c, i16 %d) {789; CHECK-LABEL: st_param_v4_i16_irrr(790; CHECK: {791; CHECK-NEXT: .reg .b16 %rs<4>;792; CHECK-EMPTY:793; CHECK-NEXT: // %bb.0:794; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_irrr_param_0];795; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_irrr_param_1];796; CHECK-NEXT: ld.param.b16 %rs3, [st_param_v4_i16_irrr_param_2];797; CHECK-NEXT: { // callseq 39, 0798; CHECK-NEXT: .param .align 8 .b8 param0[8];799; CHECK-NEXT: st.param.v4.b16 [param0], {1, %rs1, %rs2, %rs3};800; CHECK-NEXT: call.uni call_v4_i16, (param0);801; CHECK-NEXT: } // callseq 39802; CHECK-NEXT: ret;803 %struct.irrr0 = insertvalue %struct.short4 poison, i16 1, 0804 %struct.irrr1 = insertvalue %struct.short4 %struct.irrr0, i16 %b, 1805 %struct.irrr2 = insertvalue %struct.short4 %struct.irrr1, i16 %c, 2806 %struct.irrr3 = insertvalue %struct.short4 %struct.irrr2, i16 %d, 3807 call void @call_v4_i16(%struct.short4 %struct.irrr3)808 ret void809}810define void @st_param_v4_i16_rirr(i16 %a, i16 %c, i16 %d) {811; CHECK-LABEL: st_param_v4_i16_rirr(812; CHECK: {813; CHECK-NEXT: .reg .b16 %rs<4>;814; CHECK-EMPTY:815; CHECK-NEXT: // %bb.0:816; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_rirr_param_0];817; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_rirr_param_1];818; CHECK-NEXT: ld.param.b16 %rs3, [st_param_v4_i16_rirr_param_2];819; CHECK-NEXT: { // callseq 40, 0820; CHECK-NEXT: .param .align 8 .b8 param0[8];821; CHECK-NEXT: st.param.v4.b16 [param0], {%rs1, 2, %rs2, %rs3};822; CHECK-NEXT: call.uni call_v4_i16, (param0);823; CHECK-NEXT: } // callseq 40824; CHECK-NEXT: ret;825 %struct.rirr0 = insertvalue %struct.short4 poison, i16 %a, 0826 %struct.rirr1 = insertvalue %struct.short4 %struct.rirr0, i16 2, 1827 %struct.rirr2 = insertvalue %struct.short4 %struct.rirr1, i16 %c, 2828 %struct.rirr3 = insertvalue %struct.short4 %struct.rirr2, i16 %d, 3829 call void @call_v4_i16(%struct.short4 %struct.rirr3)830 ret void831}832define void @st_param_v4_i16_rrir(i16 %a, i16 %b, i16 %d) {833; CHECK-LABEL: st_param_v4_i16_rrir(834; CHECK: {835; CHECK-NEXT: .reg .b16 %rs<4>;836; CHECK-EMPTY:837; CHECK-NEXT: // %bb.0:838; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_rrir_param_0];839; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_rrir_param_1];840; CHECK-NEXT: ld.param.b16 %rs3, [st_param_v4_i16_rrir_param_2];841; CHECK-NEXT: { // callseq 41, 0842; CHECK-NEXT: .param .align 8 .b8 param0[8];843; CHECK-NEXT: st.param.v4.b16 [param0], {%rs1, %rs2, 3, %rs3};844; CHECK-NEXT: call.uni call_v4_i16, (param0);845; CHECK-NEXT: } // callseq 41846; CHECK-NEXT: ret;847 %struct.rrir0 = insertvalue %struct.short4 poison, i16 %a, 0848 %struct.rrir1 = insertvalue %struct.short4 %struct.rrir0, i16 %b, 1849 %struct.rrir2 = insertvalue %struct.short4 %struct.rrir1, i16 3, 2850 %struct.rrir3 = insertvalue %struct.short4 %struct.rrir2, i16 %d, 3851 call void @call_v4_i16(%struct.short4 %struct.rrir3)852 ret void853}854define void @st_param_v4_i16_rrri(i16 %a, i16 %b, i16 %c) {855; CHECK-LABEL: st_param_v4_i16_rrri(856; CHECK: {857; CHECK-NEXT: .reg .b16 %rs<4>;858; CHECK-EMPTY:859; CHECK-NEXT: // %bb.0:860; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_rrri_param_0];861; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_rrri_param_1];862; CHECK-NEXT: ld.param.b16 %rs3, [st_param_v4_i16_rrri_param_2];863; CHECK-NEXT: { // callseq 42, 0864; CHECK-NEXT: .param .align 8 .b8 param0[8];865; CHECK-NEXT: st.param.v4.b16 [param0], {%rs1, %rs2, %rs3, 4};866; CHECK-NEXT: call.uni call_v4_i16, (param0);867; CHECK-NEXT: } // callseq 42868; CHECK-NEXT: ret;869 %struct.rrri0 = insertvalue %struct.short4 poison, i16 %a, 0870 %struct.rrri1 = insertvalue %struct.short4 %struct.rrri0, i16 %b, 1871 %struct.rrri2 = insertvalue %struct.short4 %struct.rrri1, i16 %c, 2872 %struct.rrri3 = insertvalue %struct.short4 %struct.rrri2, i16 4, 3873 call void @call_v4_i16(%struct.short4 %struct.rrri3)874 ret void875}876define void @st_param_v4_i16_iirr(i16 %c, i16 %d) {877; CHECK-LABEL: st_param_v4_i16_iirr(878; CHECK: {879; CHECK-NEXT: .reg .b16 %rs<3>;880; CHECK-NEXT: .reg .b32 %r<2>;881; CHECK-EMPTY:882; CHECK-NEXT: // %bb.0:883; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_iirr_param_0];884; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_iirr_param_1];885; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};886; CHECK-NEXT: { // callseq 43, 0887; CHECK-NEXT: .param .align 8 .b8 param0[8];888; CHECK-NEXT: st.param.v2.b32 [param0], {131073, %r1};889; CHECK-NEXT: call.uni call_v4_i16, (param0);890; CHECK-NEXT: } // callseq 43891; CHECK-NEXT: ret;892 %struct.iirr0 = insertvalue %struct.short4 poison, i16 1, 0893 %struct.iirr1 = insertvalue %struct.short4 %struct.iirr0, i16 2, 1894 %struct.iirr2 = insertvalue %struct.short4 %struct.iirr1, i16 %c, 2895 %struct.iirr3 = insertvalue %struct.short4 %struct.iirr2, i16 %d, 3896 call void @call_v4_i16(%struct.short4 %struct.iirr3)897 ret void898}899define void @st_param_v4_i16_irir(i16 %b, i16 %d) {900; CHECK-LABEL: st_param_v4_i16_irir(901; CHECK: {902; CHECK-NEXT: .reg .b16 %rs<3>;903; CHECK-EMPTY:904; CHECK-NEXT: // %bb.0:905; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_irir_param_0];906; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_irir_param_1];907; CHECK-NEXT: { // callseq 44, 0908; CHECK-NEXT: .param .align 8 .b8 param0[8];909; CHECK-NEXT: st.param.v4.b16 [param0], {1, %rs1, 3, %rs2};910; CHECK-NEXT: call.uni call_v4_i16, (param0);911; CHECK-NEXT: } // callseq 44912; CHECK-NEXT: ret;913 %struct.irir0 = insertvalue %struct.short4 poison, i16 1, 0914 %struct.irir1 = insertvalue %struct.short4 %struct.irir0, i16 %b, 1915 %struct.irir2 = insertvalue %struct.short4 %struct.irir1, i16 3, 2916 %struct.irir3 = insertvalue %struct.short4 %struct.irir2, i16 %d, 3917 call void @call_v4_i16(%struct.short4 %struct.irir3)918 ret void919}920define void @st_param_v4_i16_irri(i16 %b, i16 %c) {921; CHECK-LABEL: st_param_v4_i16_irri(922; CHECK: {923; CHECK-NEXT: .reg .b16 %rs<3>;924; CHECK-EMPTY:925; CHECK-NEXT: // %bb.0:926; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_irri_param_0];927; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_irri_param_1];928; CHECK-NEXT: { // callseq 45, 0929; CHECK-NEXT: .param .align 8 .b8 param0[8];930; CHECK-NEXT: st.param.v4.b16 [param0], {1, %rs1, %rs2, 4};931; CHECK-NEXT: call.uni call_v4_i16, (param0);932; CHECK-NEXT: } // callseq 45933; CHECK-NEXT: ret;934 %struct.irri0 = insertvalue %struct.short4 poison, i16 1, 0935 %struct.irri1 = insertvalue %struct.short4 %struct.irri0, i16 %b, 1936 %struct.irri2 = insertvalue %struct.short4 %struct.irri1, i16 %c, 2937 %struct.irri3 = insertvalue %struct.short4 %struct.irri2, i16 4, 3938 call void @call_v4_i16(%struct.short4 %struct.irri3)939 ret void940}941define void @st_param_v4_i16_riir(i16 %a, i16 %d) {942; CHECK-LABEL: st_param_v4_i16_riir(943; CHECK: {944; CHECK-NEXT: .reg .b16 %rs<3>;945; CHECK-EMPTY:946; CHECK-NEXT: // %bb.0:947; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_riir_param_0];948; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_riir_param_1];949; CHECK-NEXT: { // callseq 46, 0950; CHECK-NEXT: .param .align 8 .b8 param0[8];951; CHECK-NEXT: st.param.v4.b16 [param0], {%rs1, 2, 3, %rs2};952; CHECK-NEXT: call.uni call_v4_i16, (param0);953; CHECK-NEXT: } // callseq 46954; CHECK-NEXT: ret;955 %struct.riir0 = insertvalue %struct.short4 poison, i16 %a, 0956 %struct.riir1 = insertvalue %struct.short4 %struct.riir0, i16 2, 1957 %struct.riir2 = insertvalue %struct.short4 %struct.riir1, i16 3, 2958 %struct.riir3 = insertvalue %struct.short4 %struct.riir2, i16 %d, 3959 call void @call_v4_i16(%struct.short4 %struct.riir3)960 ret void961}962define void @st_param_v4_i16_riri(i16 %a, i16 %c) {963; CHECK-LABEL: st_param_v4_i16_riri(964; CHECK: {965; CHECK-NEXT: .reg .b16 %rs<3>;966; CHECK-EMPTY:967; CHECK-NEXT: // %bb.0:968; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_riri_param_0];969; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_riri_param_1];970; CHECK-NEXT: { // callseq 47, 0971; CHECK-NEXT: .param .align 8 .b8 param0[8];972; CHECK-NEXT: st.param.v4.b16 [param0], {%rs1, 2, %rs2, 4};973; CHECK-NEXT: call.uni call_v4_i16, (param0);974; CHECK-NEXT: } // callseq 47975; CHECK-NEXT: ret;976 %struct.riri0 = insertvalue %struct.short4 poison, i16 %a, 0977 %struct.riri1 = insertvalue %struct.short4 %struct.riri0, i16 2, 1978 %struct.riri2 = insertvalue %struct.short4 %struct.riri1, i16 %c, 2979 %struct.riri3 = insertvalue %struct.short4 %struct.riri2, i16 4, 3980 call void @call_v4_i16(%struct.short4 %struct.riri3)981 ret void982}983define void @st_param_v4_i16_rrii(i16 %a, i16 %b) {984; CHECK-LABEL: st_param_v4_i16_rrii(985; CHECK: {986; CHECK-NEXT: .reg .b16 %rs<3>;987; CHECK-NEXT: .reg .b32 %r<2>;988; CHECK-EMPTY:989; CHECK-NEXT: // %bb.0:990; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_rrii_param_0];991; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_rrii_param_1];992; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};993; CHECK-NEXT: { // callseq 48, 0994; CHECK-NEXT: .param .align 8 .b8 param0[8];995; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, 262147};996; CHECK-NEXT: call.uni call_v4_i16, (param0);997; CHECK-NEXT: } // callseq 48998; CHECK-NEXT: ret;999 %struct.rrii0 = insertvalue %struct.short4 poison, i16 %a, 01000 %struct.rrii1 = insertvalue %struct.short4 %struct.rrii0, i16 %b, 11001 %struct.rrii2 = insertvalue %struct.short4 %struct.rrii1, i16 3, 21002 %struct.rrii3 = insertvalue %struct.short4 %struct.rrii2, i16 4, 31003 call void @call_v4_i16(%struct.short4 %struct.rrii3)1004 ret void1005}1006define void @st_param_v4_i16_iiir(i16 %d) {1007; CHECK-LABEL: st_param_v4_i16_iiir(1008; CHECK: {1009; CHECK-NEXT: .reg .b16 %rs<3>;1010; CHECK-NEXT: .reg .b32 %r<2>;1011; CHECK-EMPTY:1012; CHECK-NEXT: // %bb.0:1013; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_iiir_param_0];1014; CHECK-NEXT: mov.b16 %rs2, 3;1015; CHECK-NEXT: mov.b32 %r1, {%rs2, %rs1};1016; CHECK-NEXT: { // callseq 49, 01017; CHECK-NEXT: .param .align 8 .b8 param0[8];1018; CHECK-NEXT: st.param.v2.b32 [param0], {131073, %r1};1019; CHECK-NEXT: call.uni call_v4_i16, (param0);1020; CHECK-NEXT: } // callseq 491021; CHECK-NEXT: ret;1022 %struct.iiir0 = insertvalue %struct.short4 poison, i16 1, 01023 %struct.iiir1 = insertvalue %struct.short4 %struct.iiir0, i16 2, 11024 %struct.iiir2 = insertvalue %struct.short4 %struct.iiir1, i16 3, 21025 %struct.iiir3 = insertvalue %struct.short4 %struct.iiir2, i16 %d, 31026 call void @call_v4_i16(%struct.short4 %struct.iiir3)1027 ret void1028}1029define void @st_param_v4_i16_iiri(i16 %c) {1030; CHECK-LABEL: st_param_v4_i16_iiri(1031; CHECK: {1032; CHECK-NEXT: .reg .b16 %rs<3>;1033; CHECK-NEXT: .reg .b32 %r<2>;1034; CHECK-EMPTY:1035; CHECK-NEXT: // %bb.0:1036; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_iiri_param_0];1037; CHECK-NEXT: mov.b16 %rs2, 4;1038; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};1039; CHECK-NEXT: { // callseq 50, 01040; CHECK-NEXT: .param .align 8 .b8 param0[8];1041; CHECK-NEXT: st.param.v2.b32 [param0], {131073, %r1};1042; CHECK-NEXT: call.uni call_v4_i16, (param0);1043; CHECK-NEXT: } // callseq 501044; CHECK-NEXT: ret;1045 %struct.iiri0 = insertvalue %struct.short4 poison, i16 1, 01046 %struct.iiri1 = insertvalue %struct.short4 %struct.iiri0, i16 2, 11047 %struct.iiri2 = insertvalue %struct.short4 %struct.iiri1, i16 %c, 21048 %struct.iiri3 = insertvalue %struct.short4 %struct.iiri2, i16 4, 31049 call void @call_v4_i16(%struct.short4 %struct.iiri3)1050 ret void1051}1052define void @st_param_v4_i16_irii(i16 %b) {1053; CHECK-LABEL: st_param_v4_i16_irii(1054; CHECK: {1055; CHECK-NEXT: .reg .b16 %rs<3>;1056; CHECK-NEXT: .reg .b32 %r<2>;1057; CHECK-EMPTY:1058; CHECK-NEXT: // %bb.0:1059; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_irii_param_0];1060; CHECK-NEXT: mov.b16 %rs2, 1;1061; CHECK-NEXT: mov.b32 %r1, {%rs2, %rs1};1062; CHECK-NEXT: { // callseq 51, 01063; CHECK-NEXT: .param .align 8 .b8 param0[8];1064; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, 262147};1065; CHECK-NEXT: call.uni call_v4_i16, (param0);1066; CHECK-NEXT: } // callseq 511067; CHECK-NEXT: ret;1068 %struct.irii0 = insertvalue %struct.short4 poison, i16 1, 01069 %struct.irii1 = insertvalue %struct.short4 %struct.irii0, i16 %b, 11070 %struct.irii2 = insertvalue %struct.short4 %struct.irii1, i16 3, 21071 %struct.irii3 = insertvalue %struct.short4 %struct.irii2, i16 4, 31072 call void @call_v4_i16(%struct.short4 %struct.irii3)1073 ret void1074}1075define void @st_param_v4_i16_riii(i16 %a) {1076; CHECK-LABEL: st_param_v4_i16_riii(1077; CHECK: {1078; CHECK-NEXT: .reg .b16 %rs<3>;1079; CHECK-NEXT: .reg .b32 %r<2>;1080; CHECK-EMPTY:1081; CHECK-NEXT: // %bb.0:1082; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_riii_param_0];1083; CHECK-NEXT: mov.b16 %rs2, 2;1084; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};1085; CHECK-NEXT: { // callseq 52, 01086; CHECK-NEXT: .param .align 8 .b8 param0[8];1087; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, 262147};1088; CHECK-NEXT: call.uni call_v4_i16, (param0);1089; CHECK-NEXT: } // callseq 521090; CHECK-NEXT: ret;1091 %struct.riii0 = insertvalue %struct.short4 poison, i16 %a, 01092 %struct.riii1 = insertvalue %struct.short4 %struct.riii0, i16 2, 11093 %struct.riii2 = insertvalue %struct.short4 %struct.riii1, i16 3, 21094 %struct.riii3 = insertvalue %struct.short4 %struct.riii2, i16 4, 31095 call void @call_v4_i16(%struct.short4 %struct.riii3)1096 ret void1097}1098 1099define void @st_param_v4_i32_iiii() {1100; CHECK-LABEL: st_param_v4_i32_iiii(1101; CHECK: {1102; CHECK-EMPTY:1103; CHECK-EMPTY:1104; CHECK-NEXT: // %bb.0:1105; CHECK-NEXT: { // callseq 53, 01106; CHECK-NEXT: .param .align 16 .b8 param0[16];1107; CHECK-NEXT: st.param.v4.b32 [param0], {1, 2, 3, 4};1108; CHECK-NEXT: call.uni call_v4_i32, (param0);1109; CHECK-NEXT: } // callseq 531110; CHECK-NEXT: ret;1111 call void @call_v4_i32(%struct.int4 { i32 1, i32 2, i32 3, i32 4 })1112 ret void1113}1114define void @st_param_v4_i32_irrr(i32 %b, i32 %c, i32 %d) {1115; CHECK-LABEL: st_param_v4_i32_irrr(1116; CHECK: {1117; CHECK-NEXT: .reg .b32 %r<4>;1118; CHECK-EMPTY:1119; CHECK-NEXT: // %bb.0:1120; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_irrr_param_0];1121; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_i32_irrr_param_1];1122; CHECK-NEXT: ld.param.b32 %r3, [st_param_v4_i32_irrr_param_2];1123; CHECK-NEXT: { // callseq 54, 01124; CHECK-NEXT: .param .align 16 .b8 param0[16];1125; CHECK-NEXT: st.param.v4.b32 [param0], {1, %r1, %r2, %r3};1126; CHECK-NEXT: call.uni call_v4_i32, (param0);1127; CHECK-NEXT: } // callseq 541128; CHECK-NEXT: ret;1129 %struct.irrr0 = insertvalue %struct.int4 poison, i32 1, 01130 %struct.irrr1 = insertvalue %struct.int4 %struct.irrr0, i32 %b, 11131 %struct.irrr2 = insertvalue %struct.int4 %struct.irrr1, i32 %c, 21132 %struct.irrr3 = insertvalue %struct.int4 %struct.irrr2, i32 %d, 31133 call void @call_v4_i32(%struct.int4 %struct.irrr3)1134 ret void1135}1136define void @st_param_v4_i32_rirr(i32 %a, i32 %c, i32 %d) {1137; CHECK-LABEL: st_param_v4_i32_rirr(1138; CHECK: {1139; CHECK-NEXT: .reg .b32 %r<4>;1140; CHECK-EMPTY:1141; CHECK-NEXT: // %bb.0:1142; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_rirr_param_0];1143; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_i32_rirr_param_1];1144; CHECK-NEXT: ld.param.b32 %r3, [st_param_v4_i32_rirr_param_2];1145; CHECK-NEXT: { // callseq 55, 01146; CHECK-NEXT: .param .align 16 .b8 param0[16];1147; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, 2, %r2, %r3};1148; CHECK-NEXT: call.uni call_v4_i32, (param0);1149; CHECK-NEXT: } // callseq 551150; CHECK-NEXT: ret;1151 %struct.rirr0 = insertvalue %struct.int4 poison, i32 %a, 01152 %struct.rirr1 = insertvalue %struct.int4 %struct.rirr0, i32 2, 11153 %struct.rirr2 = insertvalue %struct.int4 %struct.rirr1, i32 %c, 21154 %struct.rirr3 = insertvalue %struct.int4 %struct.rirr2, i32 %d, 31155 call void @call_v4_i32(%struct.int4 %struct.rirr3)1156 ret void1157}1158define void @st_param_v4_i32_rrir(i32 %a, i32 %b, i32 %d) {1159; CHECK-LABEL: st_param_v4_i32_rrir(1160; CHECK: {1161; CHECK-NEXT: .reg .b32 %r<4>;1162; CHECK-EMPTY:1163; CHECK-NEXT: // %bb.0:1164; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_rrir_param_0];1165; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_i32_rrir_param_1];1166; CHECK-NEXT: ld.param.b32 %r3, [st_param_v4_i32_rrir_param_2];1167; CHECK-NEXT: { // callseq 56, 01168; CHECK-NEXT: .param .align 16 .b8 param0[16];1169; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, %r2, 3, %r3};1170; CHECK-NEXT: call.uni call_v4_i32, (param0);1171; CHECK-NEXT: } // callseq 561172; CHECK-NEXT: ret;1173 %struct.rrir0 = insertvalue %struct.int4 poison, i32 %a, 01174 %struct.rrir1 = insertvalue %struct.int4 %struct.rrir0, i32 %b, 11175 %struct.rrir2 = insertvalue %struct.int4 %struct.rrir1, i32 3, 21176 %struct.rrir3 = insertvalue %struct.int4 %struct.rrir2, i32 %d, 31177 call void @call_v4_i32(%struct.int4 %struct.rrir3)1178 ret void1179}1180define void @st_param_v4_i32_rrri(i32 %a, i32 %b, i32 %c) {1181; CHECK-LABEL: st_param_v4_i32_rrri(1182; CHECK: {1183; CHECK-NEXT: .reg .b32 %r<4>;1184; CHECK-EMPTY:1185; CHECK-NEXT: // %bb.0:1186; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_rrri_param_0];1187; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_i32_rrri_param_1];1188; CHECK-NEXT: ld.param.b32 %r3, [st_param_v4_i32_rrri_param_2];1189; CHECK-NEXT: { // callseq 57, 01190; CHECK-NEXT: .param .align 16 .b8 param0[16];1191; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, %r2, %r3, 4};1192; CHECK-NEXT: call.uni call_v4_i32, (param0);1193; CHECK-NEXT: } // callseq 571194; CHECK-NEXT: ret;1195 %struct.rrri0 = insertvalue %struct.int4 poison, i32 %a, 01196 %struct.rrri1 = insertvalue %struct.int4 %struct.rrri0, i32 %b, 11197 %struct.rrri2 = insertvalue %struct.int4 %struct.rrri1, i32 %c, 21198 %struct.rrri3 = insertvalue %struct.int4 %struct.rrri2, i32 4, 31199 call void @call_v4_i32(%struct.int4 %struct.rrri3)1200 ret void1201}1202define void @st_param_v4_i32_iirr(i32 %c, i32 %d) {1203; CHECK-LABEL: st_param_v4_i32_iirr(1204; CHECK: {1205; CHECK-NEXT: .reg .b32 %r<3>;1206; CHECK-EMPTY:1207; CHECK-NEXT: // %bb.0:1208; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_iirr_param_0];1209; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_i32_iirr_param_1];1210; CHECK-NEXT: { // callseq 58, 01211; CHECK-NEXT: .param .align 16 .b8 param0[16];1212; CHECK-NEXT: st.param.v4.b32 [param0], {1, 2, %r1, %r2};1213; CHECK-NEXT: call.uni call_v4_i32, (param0);1214; CHECK-NEXT: } // callseq 581215; CHECK-NEXT: ret;1216 %struct.iirr0 = insertvalue %struct.int4 poison, i32 1, 01217 %struct.iirr1 = insertvalue %struct.int4 %struct.iirr0, i32 2, 11218 %struct.iirr2 = insertvalue %struct.int4 %struct.iirr1, i32 %c, 21219 %struct.iirr3 = insertvalue %struct.int4 %struct.iirr2, i32 %d, 31220 call void @call_v4_i32(%struct.int4 %struct.iirr3)1221 ret void1222}1223define void @st_param_v4_i32_irir(i32 %b, i32 %d) {1224; CHECK-LABEL: st_param_v4_i32_irir(1225; CHECK: {1226; CHECK-NEXT: .reg .b32 %r<3>;1227; CHECK-EMPTY:1228; CHECK-NEXT: // %bb.0:1229; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_irir_param_0];1230; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_i32_irir_param_1];1231; CHECK-NEXT: { // callseq 59, 01232; CHECK-NEXT: .param .align 16 .b8 param0[16];1233; CHECK-NEXT: st.param.v4.b32 [param0], {1, %r1, 3, %r2};1234; CHECK-NEXT: call.uni call_v4_i32, (param0);1235; CHECK-NEXT: } // callseq 591236; CHECK-NEXT: ret;1237 %struct.irir0 = insertvalue %struct.int4 poison, i32 1, 01238 %struct.irir1 = insertvalue %struct.int4 %struct.irir0, i32 %b, 11239 %struct.irir2 = insertvalue %struct.int4 %struct.irir1, i32 3, 21240 %struct.irir3 = insertvalue %struct.int4 %struct.irir2, i32 %d, 31241 call void @call_v4_i32(%struct.int4 %struct.irir3)1242 ret void1243}1244define void @st_param_v4_i32_irri(i32 %b, i32 %c) {1245; CHECK-LABEL: st_param_v4_i32_irri(1246; CHECK: {1247; CHECK-NEXT: .reg .b32 %r<3>;1248; CHECK-EMPTY:1249; CHECK-NEXT: // %bb.0:1250; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_irri_param_0];1251; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_i32_irri_param_1];1252; CHECK-NEXT: { // callseq 60, 01253; CHECK-NEXT: .param .align 16 .b8 param0[16];1254; CHECK-NEXT: st.param.v4.b32 [param0], {1, %r1, %r2, 4};1255; CHECK-NEXT: call.uni call_v4_i32, (param0);1256; CHECK-NEXT: } // callseq 601257; CHECK-NEXT: ret;1258 %struct.irri0 = insertvalue %struct.int4 poison, i32 1, 01259 %struct.irri1 = insertvalue %struct.int4 %struct.irri0, i32 %b, 11260 %struct.irri2 = insertvalue %struct.int4 %struct.irri1, i32 %c, 21261 %struct.irri3 = insertvalue %struct.int4 %struct.irri2, i32 4, 31262 call void @call_v4_i32(%struct.int4 %struct.irri3)1263 ret void1264}1265define void @st_param_v4_i32_riir(i32 %a, i32 %d) {1266; CHECK-LABEL: st_param_v4_i32_riir(1267; CHECK: {1268; CHECK-NEXT: .reg .b32 %r<3>;1269; CHECK-EMPTY:1270; CHECK-NEXT: // %bb.0:1271; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_riir_param_0];1272; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_i32_riir_param_1];1273; CHECK-NEXT: { // callseq 61, 01274; CHECK-NEXT: .param .align 16 .b8 param0[16];1275; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, 2, 3, %r2};1276; CHECK-NEXT: call.uni call_v4_i32, (param0);1277; CHECK-NEXT: } // callseq 611278; CHECK-NEXT: ret;1279 %struct.riir0 = insertvalue %struct.int4 poison, i32 %a, 01280 %struct.riir1 = insertvalue %struct.int4 %struct.riir0, i32 2, 11281 %struct.riir2 = insertvalue %struct.int4 %struct.riir1, i32 3, 21282 %struct.riir3 = insertvalue %struct.int4 %struct.riir2, i32 %d, 31283 call void @call_v4_i32(%struct.int4 %struct.riir3)1284 ret void1285}1286define void @st_param_v4_i32_riri(i32 %a, i32 %c) {1287; CHECK-LABEL: st_param_v4_i32_riri(1288; CHECK: {1289; CHECK-NEXT: .reg .b32 %r<3>;1290; CHECK-EMPTY:1291; CHECK-NEXT: // %bb.0:1292; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_riri_param_0];1293; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_i32_riri_param_1];1294; CHECK-NEXT: { // callseq 62, 01295; CHECK-NEXT: .param .align 16 .b8 param0[16];1296; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, 2, %r2, 4};1297; CHECK-NEXT: call.uni call_v4_i32, (param0);1298; CHECK-NEXT: } // callseq 621299; CHECK-NEXT: ret;1300 %struct.riri0 = insertvalue %struct.int4 poison, i32 %a, 01301 %struct.riri1 = insertvalue %struct.int4 %struct.riri0, i32 2, 11302 %struct.riri2 = insertvalue %struct.int4 %struct.riri1, i32 %c, 21303 %struct.riri3 = insertvalue %struct.int4 %struct.riri2, i32 4, 31304 call void @call_v4_i32(%struct.int4 %struct.riri3)1305 ret void1306}1307define void @st_param_v4_i32_rrii(i32 %a, i32 %b) {1308; CHECK-LABEL: st_param_v4_i32_rrii(1309; CHECK: {1310; CHECK-NEXT: .reg .b32 %r<3>;1311; CHECK-EMPTY:1312; CHECK-NEXT: // %bb.0:1313; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_rrii_param_0];1314; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_i32_rrii_param_1];1315; CHECK-NEXT: { // callseq 63, 01316; CHECK-NEXT: .param .align 16 .b8 param0[16];1317; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, %r2, 3, 4};1318; CHECK-NEXT: call.uni call_v4_i32, (param0);1319; CHECK-NEXT: } // callseq 631320; CHECK-NEXT: ret;1321 %struct.rrii0 = insertvalue %struct.int4 poison, i32 %a, 01322 %struct.rrii1 = insertvalue %struct.int4 %struct.rrii0, i32 %b, 11323 %struct.rrii2 = insertvalue %struct.int4 %struct.rrii1, i32 3, 21324 %struct.rrii3 = insertvalue %struct.int4 %struct.rrii2, i32 4, 31325 call void @call_v4_i32(%struct.int4 %struct.rrii3)1326 ret void1327}1328define void @st_param_v4_i32_iiir(i32 %d) {1329; CHECK-LABEL: st_param_v4_i32_iiir(1330; CHECK: {1331; CHECK-NEXT: .reg .b32 %r<2>;1332; CHECK-EMPTY:1333; CHECK-NEXT: // %bb.0:1334; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_iiir_param_0];1335; CHECK-NEXT: { // callseq 64, 01336; CHECK-NEXT: .param .align 16 .b8 param0[16];1337; CHECK-NEXT: st.param.v4.b32 [param0], {1, 2, 3, %r1};1338; CHECK-NEXT: call.uni call_v4_i32, (param0);1339; CHECK-NEXT: } // callseq 641340; CHECK-NEXT: ret;1341 %struct.iiir0 = insertvalue %struct.int4 poison, i32 1, 01342 %struct.iiir1 = insertvalue %struct.int4 %struct.iiir0, i32 2, 11343 %struct.iiir2 = insertvalue %struct.int4 %struct.iiir1, i32 3, 21344 %struct.iiir3 = insertvalue %struct.int4 %struct.iiir2, i32 %d, 31345 call void @call_v4_i32(%struct.int4 %struct.iiir3)1346 ret void1347}1348define void @st_param_v4_i32_iiri(i32 %c) {1349; CHECK-LABEL: st_param_v4_i32_iiri(1350; CHECK: {1351; CHECK-NEXT: .reg .b32 %r<2>;1352; CHECK-EMPTY:1353; CHECK-NEXT: // %bb.0:1354; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_iiri_param_0];1355; CHECK-NEXT: { // callseq 65, 01356; CHECK-NEXT: .param .align 16 .b8 param0[16];1357; CHECK-NEXT: st.param.v4.b32 [param0], {1, 2, %r1, 4};1358; CHECK-NEXT: call.uni call_v4_i32, (param0);1359; CHECK-NEXT: } // callseq 651360; CHECK-NEXT: ret;1361 %struct.iiri0 = insertvalue %struct.int4 poison, i32 1, 01362 %struct.iiri1 = insertvalue %struct.int4 %struct.iiri0, i32 2, 11363 %struct.iiri2 = insertvalue %struct.int4 %struct.iiri1, i32 %c, 21364 %struct.iiri3 = insertvalue %struct.int4 %struct.iiri2, i32 4, 31365 call void @call_v4_i32(%struct.int4 %struct.iiri3)1366 ret void1367}1368define void @st_param_v4_i32_irii(i32 %b) {1369; CHECK-LABEL: st_param_v4_i32_irii(1370; CHECK: {1371; CHECK-NEXT: .reg .b32 %r<2>;1372; CHECK-EMPTY:1373; CHECK-NEXT: // %bb.0:1374; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_irii_param_0];1375; CHECK-NEXT: { // callseq 66, 01376; CHECK-NEXT: .param .align 16 .b8 param0[16];1377; CHECK-NEXT: st.param.v4.b32 [param0], {1, %r1, 3, 4};1378; CHECK-NEXT: call.uni call_v4_i32, (param0);1379; CHECK-NEXT: } // callseq 661380; CHECK-NEXT: ret;1381 %struct.irii0 = insertvalue %struct.int4 poison, i32 1, 01382 %struct.irii1 = insertvalue %struct.int4 %struct.irii0, i32 %b, 11383 %struct.irii2 = insertvalue %struct.int4 %struct.irii1, i32 3, 21384 %struct.irii3 = insertvalue %struct.int4 %struct.irii2, i32 4, 31385 call void @call_v4_i32(%struct.int4 %struct.irii3)1386 ret void1387}1388define void @st_param_v4_i32_riii(i32 %a) {1389; CHECK-LABEL: st_param_v4_i32_riii(1390; CHECK: {1391; CHECK-NEXT: .reg .b32 %r<2>;1392; CHECK-EMPTY:1393; CHECK-NEXT: // %bb.0:1394; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_i32_riii_param_0];1395; CHECK-NEXT: { // callseq 67, 01396; CHECK-NEXT: .param .align 16 .b8 param0[16];1397; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, 2, 3, 4};1398; CHECK-NEXT: call.uni call_v4_i32, (param0);1399; CHECK-NEXT: } // callseq 671400; CHECK-NEXT: ret;1401 %struct.riii0 = insertvalue %struct.int4 poison, i32 %a, 01402 %struct.riii1 = insertvalue %struct.int4 %struct.riii0, i32 2, 11403 %struct.riii2 = insertvalue %struct.int4 %struct.riii1, i32 3, 21404 %struct.riii3 = insertvalue %struct.int4 %struct.riii2, i32 4, 31405 call void @call_v4_i32(%struct.int4 %struct.riii3)1406 ret void1407}1408 1409define void @st_param_v4_f32_iiii() {1410; CHECK-LABEL: st_param_v4_f32_iiii(1411; CHECK: {1412; CHECK-EMPTY:1413; CHECK-EMPTY:1414; CHECK-NEXT: // %bb.0:1415; CHECK-NEXT: { // callseq 68, 01416; CHECK-NEXT: .param .align 16 .b8 param0[16];1417; CHECK-NEXT: st.param.v4.b32 [param0], {0f3F800000, 0f40000000, 0f40400000, 0f40800000};1418; CHECK-NEXT: call.uni call_v4_f32, (param0);1419; CHECK-NEXT: } // callseq 681420; CHECK-NEXT: ret;1421 call void @call_v4_f32(%struct.float4 { float 1.0, float 2.0, float 3.0, float 4.0 })1422 ret void1423}1424define void @st_param_v4_f32_irrr(float %b, float %c, float %d) {1425; CHECK-LABEL: st_param_v4_f32_irrr(1426; CHECK: {1427; CHECK-NEXT: .reg .b32 %r<4>;1428; CHECK-EMPTY:1429; CHECK-NEXT: // %bb.0:1430; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_irrr_param_0];1431; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_f32_irrr_param_1];1432; CHECK-NEXT: ld.param.b32 %r3, [st_param_v4_f32_irrr_param_2];1433; CHECK-NEXT: { // callseq 69, 01434; CHECK-NEXT: .param .align 16 .b8 param0[16];1435; CHECK-NEXT: st.param.v4.b32 [param0], {0f3F800000, %r1, %r2, %r3};1436; CHECK-NEXT: call.uni call_v4_f32, (param0);1437; CHECK-NEXT: } // callseq 691438; CHECK-NEXT: ret;1439 %struct.irrr0 = insertvalue %struct.float4 poison, float 1.0, 01440 %struct.irrr1 = insertvalue %struct.float4 %struct.irrr0, float %b, 11441 %struct.irrr2 = insertvalue %struct.float4 %struct.irrr1, float %c, 21442 %struct.irrr3 = insertvalue %struct.float4 %struct.irrr2, float %d, 31443 call void @call_v4_f32(%struct.float4 %struct.irrr3)1444 ret void1445}1446define void @st_param_v4_f32_rirr(float %a, float %c, float %d) {1447; CHECK-LABEL: st_param_v4_f32_rirr(1448; CHECK: {1449; CHECK-NEXT: .reg .b32 %r<4>;1450; CHECK-EMPTY:1451; CHECK-NEXT: // %bb.0:1452; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_rirr_param_0];1453; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_f32_rirr_param_1];1454; CHECK-NEXT: ld.param.b32 %r3, [st_param_v4_f32_rirr_param_2];1455; CHECK-NEXT: { // callseq 70, 01456; CHECK-NEXT: .param .align 16 .b8 param0[16];1457; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, 0f40000000, %r2, %r3};1458; CHECK-NEXT: call.uni call_v4_f32, (param0);1459; CHECK-NEXT: } // callseq 701460; CHECK-NEXT: ret;1461 %struct.rirr0 = insertvalue %struct.float4 poison, float %a, 01462 %struct.rirr1 = insertvalue %struct.float4 %struct.rirr0, float 2.0, 11463 %struct.rirr2 = insertvalue %struct.float4 %struct.rirr1, float %c, 21464 %struct.rirr3 = insertvalue %struct.float4 %struct.rirr2, float %d, 31465 call void @call_v4_f32(%struct.float4 %struct.rirr3)1466 ret void1467}1468define void @st_param_v4_f32_rrir(float %a, float %b, float %d) {1469; CHECK-LABEL: st_param_v4_f32_rrir(1470; CHECK: {1471; CHECK-NEXT: .reg .b32 %r<4>;1472; CHECK-EMPTY:1473; CHECK-NEXT: // %bb.0:1474; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_rrir_param_0];1475; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_f32_rrir_param_1];1476; CHECK-NEXT: ld.param.b32 %r3, [st_param_v4_f32_rrir_param_2];1477; CHECK-NEXT: { // callseq 71, 01478; CHECK-NEXT: .param .align 16 .b8 param0[16];1479; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, %r2, 0f40400000, %r3};1480; CHECK-NEXT: call.uni call_v4_f32, (param0);1481; CHECK-NEXT: } // callseq 711482; CHECK-NEXT: ret;1483 %struct.rrir0 = insertvalue %struct.float4 poison, float %a, 01484 %struct.rrir1 = insertvalue %struct.float4 %struct.rrir0, float %b, 11485 %struct.rrir2 = insertvalue %struct.float4 %struct.rrir1, float 3.0, 21486 %struct.rrir3 = insertvalue %struct.float4 %struct.rrir2, float %d, 31487 call void @call_v4_f32(%struct.float4 %struct.rrir3)1488 ret void1489}1490define void @st_param_v4_f32_rrri(float %a, float %b, float %c) {1491; CHECK-LABEL: st_param_v4_f32_rrri(1492; CHECK: {1493; CHECK-NEXT: .reg .b32 %r<4>;1494; CHECK-EMPTY:1495; CHECK-NEXT: // %bb.0:1496; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_rrri_param_0];1497; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_f32_rrri_param_1];1498; CHECK-NEXT: ld.param.b32 %r3, [st_param_v4_f32_rrri_param_2];1499; CHECK-NEXT: { // callseq 72, 01500; CHECK-NEXT: .param .align 16 .b8 param0[16];1501; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, %r2, %r3, 0f40800000};1502; CHECK-NEXT: call.uni call_v4_f32, (param0);1503; CHECK-NEXT: } // callseq 721504; CHECK-NEXT: ret;1505 %struct.rrri0 = insertvalue %struct.float4 poison, float %a, 01506 %struct.rrri1 = insertvalue %struct.float4 %struct.rrri0, float %b, 11507 %struct.rrri2 = insertvalue %struct.float4 %struct.rrri1, float %c, 21508 %struct.rrri3 = insertvalue %struct.float4 %struct.rrri2, float 4.0, 31509 call void @call_v4_f32(%struct.float4 %struct.rrri3)1510 ret void1511}1512define void @st_param_v4_f32_iirr(float %c, float %d) {1513; CHECK-LABEL: st_param_v4_f32_iirr(1514; CHECK: {1515; CHECK-NEXT: .reg .b32 %r<3>;1516; CHECK-EMPTY:1517; CHECK-NEXT: // %bb.0:1518; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_iirr_param_0];1519; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_f32_iirr_param_1];1520; CHECK-NEXT: { // callseq 73, 01521; CHECK-NEXT: .param .align 16 .b8 param0[16];1522; CHECK-NEXT: st.param.v4.b32 [param0], {0f3F800000, 0f40000000, %r1, %r2};1523; CHECK-NEXT: call.uni call_v4_f32, (param0);1524; CHECK-NEXT: } // callseq 731525; CHECK-NEXT: ret;1526 %struct.iirr0 = insertvalue %struct.float4 poison, float 1.0, 01527 %struct.iirr1 = insertvalue %struct.float4 %struct.iirr0, float 2.0, 11528 %struct.iirr2 = insertvalue %struct.float4 %struct.iirr1, float %c, 21529 %struct.iirr3 = insertvalue %struct.float4 %struct.iirr2, float %d, 31530 call void @call_v4_f32(%struct.float4 %struct.iirr3)1531 ret void1532}1533define void @st_param_v4_f32_irir(float %b, float %d) {1534; CHECK-LABEL: st_param_v4_f32_irir(1535; CHECK: {1536; CHECK-NEXT: .reg .b32 %r<3>;1537; CHECK-EMPTY:1538; CHECK-NEXT: // %bb.0:1539; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_irir_param_0];1540; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_f32_irir_param_1];1541; CHECK-NEXT: { // callseq 74, 01542; CHECK-NEXT: .param .align 16 .b8 param0[16];1543; CHECK-NEXT: st.param.v4.b32 [param0], {0f3F800000, %r1, 0f40400000, %r2};1544; CHECK-NEXT: call.uni call_v4_f32, (param0);1545; CHECK-NEXT: } // callseq 741546; CHECK-NEXT: ret;1547 %struct.irir0 = insertvalue %struct.float4 poison, float 1.0, 01548 %struct.irir1 = insertvalue %struct.float4 %struct.irir0, float %b, 11549 %struct.irir2 = insertvalue %struct.float4 %struct.irir1, float 3.0, 21550 %struct.irir3 = insertvalue %struct.float4 %struct.irir2, float %d, 31551 call void @call_v4_f32(%struct.float4 %struct.irir3)1552 ret void1553}1554define void @st_param_v4_f32_irri(float %b, float %c) {1555; CHECK-LABEL: st_param_v4_f32_irri(1556; CHECK: {1557; CHECK-NEXT: .reg .b32 %r<3>;1558; CHECK-EMPTY:1559; CHECK-NEXT: // %bb.0:1560; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_irri_param_0];1561; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_f32_irri_param_1];1562; CHECK-NEXT: { // callseq 75, 01563; CHECK-NEXT: .param .align 16 .b8 param0[16];1564; CHECK-NEXT: st.param.v4.b32 [param0], {0f3F800000, %r1, %r2, 0f40800000};1565; CHECK-NEXT: call.uni call_v4_f32, (param0);1566; CHECK-NEXT: } // callseq 751567; CHECK-NEXT: ret;1568 %struct.irri0 = insertvalue %struct.float4 poison, float 1.0, 01569 %struct.irri1 = insertvalue %struct.float4 %struct.irri0, float %b, 11570 %struct.irri2 = insertvalue %struct.float4 %struct.irri1, float %c, 21571 %struct.irri3 = insertvalue %struct.float4 %struct.irri2, float 4.0, 31572 call void @call_v4_f32(%struct.float4 %struct.irri3)1573 ret void1574}1575define void @st_param_v4_f32_riir(float %a, float %d) {1576; CHECK-LABEL: st_param_v4_f32_riir(1577; CHECK: {1578; CHECK-NEXT: .reg .b32 %r<3>;1579; CHECK-EMPTY:1580; CHECK-NEXT: // %bb.0:1581; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_riir_param_0];1582; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_f32_riir_param_1];1583; CHECK-NEXT: { // callseq 76, 01584; CHECK-NEXT: .param .align 16 .b8 param0[16];1585; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, 0f40000000, 0f40400000, %r2};1586; CHECK-NEXT: call.uni call_v4_f32, (param0);1587; CHECK-NEXT: } // callseq 761588; CHECK-NEXT: ret;1589 %struct.riir0 = insertvalue %struct.float4 poison, float %a, 01590 %struct.riir1 = insertvalue %struct.float4 %struct.riir0, float 2.0, 11591 %struct.riir2 = insertvalue %struct.float4 %struct.riir1, float 3.0, 21592 %struct.riir3 = insertvalue %struct.float4 %struct.riir2, float %d, 31593 call void @call_v4_f32(%struct.float4 %struct.riir3)1594 ret void1595}1596define void @st_param_v4_f32_riri(float %a, float %c) {1597; CHECK-LABEL: st_param_v4_f32_riri(1598; CHECK: {1599; CHECK-NEXT: .reg .b32 %r<3>;1600; CHECK-EMPTY:1601; CHECK-NEXT: // %bb.0:1602; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_riri_param_0];1603; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_f32_riri_param_1];1604; CHECK-NEXT: { // callseq 77, 01605; CHECK-NEXT: .param .align 16 .b8 param0[16];1606; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, 0f40000000, %r2, 0f40800000};1607; CHECK-NEXT: call.uni call_v4_f32, (param0);1608; CHECK-NEXT: } // callseq 771609; CHECK-NEXT: ret;1610 %struct.riri0 = insertvalue %struct.float4 poison, float %a, 01611 %struct.riri1 = insertvalue %struct.float4 %struct.riri0, float 2.0, 11612 %struct.riri2 = insertvalue %struct.float4 %struct.riri1, float %c, 21613 %struct.riri3 = insertvalue %struct.float4 %struct.riri2, float 4.0, 31614 call void @call_v4_f32(%struct.float4 %struct.riri3)1615 ret void1616}1617define void @st_param_v4_f32_rrii(float %a, float %b) {1618; CHECK-LABEL: st_param_v4_f32_rrii(1619; CHECK: {1620; CHECK-NEXT: .reg .b32 %r<3>;1621; CHECK-EMPTY:1622; CHECK-NEXT: // %bb.0:1623; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_rrii_param_0];1624; CHECK-NEXT: ld.param.b32 %r2, [st_param_v4_f32_rrii_param_1];1625; CHECK-NEXT: { // callseq 78, 01626; CHECK-NEXT: .param .align 16 .b8 param0[16];1627; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, %r2, 0f40400000, 0f40800000};1628; CHECK-NEXT: call.uni call_v4_f32, (param0);1629; CHECK-NEXT: } // callseq 781630; CHECK-NEXT: ret;1631 %struct.rrii0 = insertvalue %struct.float4 poison, float %a, 01632 %struct.rrii1 = insertvalue %struct.float4 %struct.rrii0, float %b, 11633 %struct.rrii2 = insertvalue %struct.float4 %struct.rrii1, float 3.0, 21634 %struct.rrii3 = insertvalue %struct.float4 %struct.rrii2, float 4.0, 31635 call void @call_v4_f32(%struct.float4 %struct.rrii3)1636 ret void1637}1638define void @st_param_v4_f32_iiir(float %d) {1639; CHECK-LABEL: st_param_v4_f32_iiir(1640; CHECK: {1641; CHECK-NEXT: .reg .b32 %r<2>;1642; CHECK-EMPTY:1643; CHECK-NEXT: // %bb.0:1644; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_iiir_param_0];1645; CHECK-NEXT: { // callseq 79, 01646; CHECK-NEXT: .param .align 16 .b8 param0[16];1647; CHECK-NEXT: st.param.v4.b32 [param0], {0f3F800000, 0f40000000, 0f40400000, %r1};1648; CHECK-NEXT: call.uni call_v4_f32, (param0);1649; CHECK-NEXT: } // callseq 791650; CHECK-NEXT: ret;1651 %struct.iiir0 = insertvalue %struct.float4 poison, float 1.0, 01652 %struct.iiir1 = insertvalue %struct.float4 %struct.iiir0, float 2.0, 11653 %struct.iiir2 = insertvalue %struct.float4 %struct.iiir1, float 3.0, 21654 %struct.iiir3 = insertvalue %struct.float4 %struct.iiir2, float %d, 31655 call void @call_v4_f32(%struct.float4 %struct.iiir3)1656 ret void1657}1658define void @st_param_v4_f32_iiri(float %c) {1659; CHECK-LABEL: st_param_v4_f32_iiri(1660; CHECK: {1661; CHECK-NEXT: .reg .b32 %r<2>;1662; CHECK-EMPTY:1663; CHECK-NEXT: // %bb.0:1664; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_iiri_param_0];1665; CHECK-NEXT: { // callseq 80, 01666; CHECK-NEXT: .param .align 16 .b8 param0[16];1667; CHECK-NEXT: st.param.v4.b32 [param0], {0f3F800000, 0f40000000, %r1, 0f40800000};1668; CHECK-NEXT: call.uni call_v4_f32, (param0);1669; CHECK-NEXT: } // callseq 801670; CHECK-NEXT: ret;1671 %struct.iiri0 = insertvalue %struct.float4 poison, float 1.0, 01672 %struct.iiri1 = insertvalue %struct.float4 %struct.iiri0, float 2.0, 11673 %struct.iiri2 = insertvalue %struct.float4 %struct.iiri1, float %c, 21674 %struct.iiri3 = insertvalue %struct.float4 %struct.iiri2, float 4.0, 31675 call void @call_v4_f32(%struct.float4 %struct.iiri3)1676 ret void1677}1678define void @st_param_v4_f32_irii(float %b) {1679; CHECK-LABEL: st_param_v4_f32_irii(1680; CHECK: {1681; CHECK-NEXT: .reg .b32 %r<2>;1682; CHECK-EMPTY:1683; CHECK-NEXT: // %bb.0:1684; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_irii_param_0];1685; CHECK-NEXT: { // callseq 81, 01686; CHECK-NEXT: .param .align 16 .b8 param0[16];1687; CHECK-NEXT: st.param.v4.b32 [param0], {0f3F800000, %r1, 0f40400000, 0f40800000};1688; CHECK-NEXT: call.uni call_v4_f32, (param0);1689; CHECK-NEXT: } // callseq 811690; CHECK-NEXT: ret;1691 %struct.irii0 = insertvalue %struct.float4 poison, float 1.0, 01692 %struct.irii1 = insertvalue %struct.float4 %struct.irii0, float %b, 11693 %struct.irii2 = insertvalue %struct.float4 %struct.irii1, float 3.0, 21694 %struct.irii3 = insertvalue %struct.float4 %struct.irii2, float 4.0, 31695 call void @call_v4_f32(%struct.float4 %struct.irii3)1696 ret void1697}1698define void @st_param_v4_f32_riii(float %a) {1699; CHECK-LABEL: st_param_v4_f32_riii(1700; CHECK: {1701; CHECK-NEXT: .reg .b32 %r<2>;1702; CHECK-EMPTY:1703; CHECK-NEXT: // %bb.0:1704; CHECK-NEXT: ld.param.b32 %r1, [st_param_v4_f32_riii_param_0];1705; CHECK-NEXT: { // callseq 82, 01706; CHECK-NEXT: .param .align 16 .b8 param0[16];1707; CHECK-NEXT: st.param.v4.b32 [param0], {%r1, 0f40000000, 0f40400000, 0f40800000};1708; CHECK-NEXT: call.uni call_v4_f32, (param0);1709; CHECK-NEXT: } // callseq 821710; CHECK-NEXT: ret;1711 %struct.riii0 = insertvalue %struct.float4 poison, float %a, 01712 %struct.riii1 = insertvalue %struct.float4 %struct.riii0, float 2.0, 11713 %struct.riii2 = insertvalue %struct.float4 %struct.riii1, float 3.0, 21714 %struct.riii3 = insertvalue %struct.float4 %struct.riii2, float 4.0, 31715 call void @call_v4_f32(%struct.float4 %struct.riii3)1716 ret void1717}1718 1719declare void @call_v4_i8(%struct.char4 alignstack(4))1720declare void @call_v4_i16(%struct.short4 alignstack(8))1721declare void @call_v4_i32(%struct.int4 alignstack(16))1722declare void @call_v4_f32(%struct.float4 alignstack(16))1723 1724define void @st_param_bfloat() {1725; CHECK-LABEL: st_param_bfloat(1726; CHECK: {1727; CHECK-EMPTY:1728; CHECK-EMPTY:1729; CHECK-NEXT: // %bb.0:1730; CHECK-NEXT: { // callseq 83, 01731; CHECK-NEXT: .param .align 2 .b8 param0[2];1732; CHECK-NEXT: st.param.b16 [param0], 0x4100;1733; CHECK-NEXT: call.uni call_bfloat, (param0);1734; CHECK-NEXT: } // callseq 831735; CHECK-NEXT: ret;1736 %five = bitcast i16 16640 to bfloat1737 call void @call_bfloat(bfloat %five)1738 ret void1739}1740 1741declare void @call_bfloat(bfloat)1742