diff options
Diffstat (limited to 'llvm/test/CodeGen/NVPTX/st-param-imm.ll')
| -rw-r--r-- | llvm/test/CodeGen/NVPTX/st-param-imm.ll | 201 |
1 files changed, 126 insertions, 75 deletions
diff --git a/llvm/test/CodeGen/NVPTX/st-param-imm.ll b/llvm/test/CodeGen/NVPTX/st-param-imm.ll index 6aa1119..f90435a 100644 --- a/llvm/test/CodeGen/NVPTX/st-param-imm.ll +++ b/llvm/test/CodeGen/NVPTX/st-param-imm.ll @@ -26,8 +26,8 @@ define void @st_param_i8_i16() { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: { // callseq 0, 0 ; CHECK-NEXT: .param .align 2 .b8 param0[4]; -; CHECK-NEXT: st.param.b8 [param0], 1; ; CHECK-NEXT: st.param.b16 [param0+2], 2; +; CHECK-NEXT: st.param.b8 [param0], 1; ; CHECK-NEXT: call.uni call_i8_i16, (param0); ; CHECK-NEXT: } // callseq 0 ; CHECK-NEXT: ret; @@ -75,7 +75,7 @@ define void @st_param_f32() { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: { // callseq 3, 0 ; CHECK-NEXT: .param .b32 param0; -; CHECK-NEXT: st.param.b32 [param0], 0f40A00000; +; CHECK-NEXT: st.param.b32 [param0], 1084227584; ; CHECK-NEXT: call.uni call_f32, (param0); ; CHECK-NEXT: } // callseq 3 ; CHECK-NEXT: ret; @@ -91,7 +91,7 @@ define void @st_param_f64() { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: { // callseq 4, 0 ; CHECK-NEXT: .param .b64 param0; -; CHECK-NEXT: st.param.b64 [param0], 0d4018000000000000; +; CHECK-NEXT: st.param.b64 [param0], 4618441417868443648; ; CHECK-NEXT: call.uni call_f64, (param0); ; CHECK-NEXT: } // callseq 4 ; CHECK-NEXT: ret; @@ -165,7 +165,7 @@ define void @st_param_v2_i16_ii() { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: { // callseq 8, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v2.b16 [param0], {1, 2}; +; CHECK-NEXT: st.param.b32 [param0], 131073; ; CHECK-NEXT: call.uni call_v2_i16, (param0); ; CHECK-NEXT: } // callseq 8 ; CHECK-NEXT: ret; @@ -432,7 +432,7 @@ define void @st_param_v4_i8_iiii() { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: { // callseq 23, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {1, 2, 3, 4}; +; CHECK-NEXT: st.param.b32 [param0], 67305985; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 23 ; CHECK-NEXT: ret; @@ -442,15 +442,18 @@ define void @st_param_v4_i8_iiii() { define void @st_param_v4_i8_irrr(i8 %b, i8 %c, i8 %d) { ; CHECK-LABEL: st_param_v4_i8_irrr( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<4>; +; CHECK-NEXT: .reg .b32 %r<7>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_irrr_param_2]; -; CHECK-NEXT: ld.param.b8 %rs2, [st_param_v4_i8_irrr_param_1]; -; CHECK-NEXT: ld.param.b8 %rs3, [st_param_v4_i8_irrr_param_0]; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_irrr_param_2]; +; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_irrr_param_1]; +; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U; +; CHECK-NEXT: ld.param.b8 %r4, [st_param_v4_i8_irrr_param_0]; +; CHECK-NEXT: prmt.b32 %r5, 1, %r4, 0x3340U; +; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x5410U; ; CHECK-NEXT: { // callseq 24, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {1, %rs3, %rs2, %rs1}; +; CHECK-NEXT: st.param.b32 [param0], %r6; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 24 ; CHECK-NEXT: ret; @@ -464,15 +467,18 @@ define void @st_param_v4_i8_irrr(i8 %b, i8 %c, i8 %d) { define void @st_param_v4_i8_rirr(i8 %a, i8 %c, i8 %d) { ; CHECK-LABEL: st_param_v4_i8_rirr( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<4>; +; CHECK-NEXT: .reg .b32 %r<7>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_rirr_param_2]; -; CHECK-NEXT: ld.param.b8 %rs2, [st_param_v4_i8_rirr_param_1]; -; CHECK-NEXT: ld.param.b8 %rs3, [st_param_v4_i8_rirr_param_0]; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_rirr_param_2]; +; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_rirr_param_1]; +; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U; +; CHECK-NEXT: ld.param.b8 %r4, [st_param_v4_i8_rirr_param_0]; +; CHECK-NEXT: prmt.b32 %r5, %r4, 2, 0x3340U; +; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x5410U; ; CHECK-NEXT: { // callseq 25, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {%rs3, 2, %rs2, %rs1}; +; CHECK-NEXT: st.param.b32 [param0], %r6; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 25 ; CHECK-NEXT: ret; @@ -486,15 +492,18 @@ define void @st_param_v4_i8_rirr(i8 %a, i8 %c, i8 %d) { define void @st_param_v4_i8_rrir(i8 %a, i8 %b, i8 %d) { ; CHECK-LABEL: st_param_v4_i8_rrir( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<4>; +; CHECK-NEXT: .reg .b32 %r<7>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_rrir_param_2]; -; CHECK-NEXT: ld.param.b8 %rs2, [st_param_v4_i8_rrir_param_1]; -; CHECK-NEXT: ld.param.b8 %rs3, [st_param_v4_i8_rrir_param_0]; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_rrir_param_1]; +; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_rrir_param_0]; +; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U; +; CHECK-NEXT: ld.param.b8 %r4, [st_param_v4_i8_rrir_param_2]; +; CHECK-NEXT: prmt.b32 %r5, 3, %r4, 0x3340U; +; CHECK-NEXT: prmt.b32 %r6, %r3, %r5, 0x5410U; ; CHECK-NEXT: { // callseq 26, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {%rs3, %rs2, 3, %rs1}; +; CHECK-NEXT: st.param.b32 [param0], %r6; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 26 ; CHECK-NEXT: ret; @@ -508,15 +517,18 @@ define void @st_param_v4_i8_rrir(i8 %a, i8 %b, i8 %d) { define void @st_param_v4_i8_rrri(i8 %a, i8 %b, i8 %c) { ; CHECK-LABEL: st_param_v4_i8_rrri( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<4>; +; CHECK-NEXT: .reg .b32 %r<7>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_rrri_param_2]; -; CHECK-NEXT: ld.param.b8 %rs2, [st_param_v4_i8_rrri_param_1]; -; CHECK-NEXT: ld.param.b8 %rs3, [st_param_v4_i8_rrri_param_0]; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_rrri_param_1]; +; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_rrri_param_0]; +; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U; +; CHECK-NEXT: ld.param.b8 %r4, [st_param_v4_i8_rrri_param_2]; +; CHECK-NEXT: prmt.b32 %r5, %r4, 4, 0x3340U; +; CHECK-NEXT: prmt.b32 %r6, %r3, %r5, 0x5410U; ; CHECK-NEXT: { // callseq 27, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {%rs3, %rs2, %rs1, 4}; +; CHECK-NEXT: st.param.b32 [param0], %r6; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 27 ; CHECK-NEXT: ret; @@ -530,14 +542,16 @@ define void @st_param_v4_i8_rrri(i8 %a, i8 %b, i8 %c) { define void @st_param_v4_i8_iirr(i8 %c, i8 %d) { ; CHECK-LABEL: st_param_v4_i8_iirr( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<5>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_iirr_param_1]; -; CHECK-NEXT: ld.param.b8 %rs2, [st_param_v4_i8_iirr_param_0]; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_iirr_param_1]; +; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_iirr_param_0]; +; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U; +; CHECK-NEXT: prmt.b32 %r4, 513, %r3, 0x5410U; ; CHECK-NEXT: { // callseq 28, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {1, 2, %rs2, %rs1}; +; CHECK-NEXT: st.param.b32 [param0], %r4; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 28 ; CHECK-NEXT: ret; @@ -551,14 +565,17 @@ define void @st_param_v4_i8_iirr(i8 %c, i8 %d) { define void @st_param_v4_i8_irir(i8 %b, i8 %d) { ; CHECK-LABEL: st_param_v4_i8_irir( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<6>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_irir_param_1]; -; CHECK-NEXT: ld.param.b8 %rs2, [st_param_v4_i8_irir_param_0]; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_irir_param_1]; +; CHECK-NEXT: prmt.b32 %r2, 3, %r1, 0x3340U; +; CHECK-NEXT: ld.param.b8 %r3, [st_param_v4_i8_irir_param_0]; +; CHECK-NEXT: prmt.b32 %r4, 1, %r3, 0x3340U; +; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; ; CHECK-NEXT: { // callseq 29, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {1, %rs2, 3, %rs1}; +; CHECK-NEXT: st.param.b32 [param0], %r5; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 29 ; CHECK-NEXT: ret; @@ -572,14 +589,17 @@ define void @st_param_v4_i8_irir(i8 %b, i8 %d) { define void @st_param_v4_i8_irri(i8 %b, i8 %c) { ; CHECK-LABEL: st_param_v4_i8_irri( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<6>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_irri_param_1]; -; CHECK-NEXT: ld.param.b8 %rs2, [st_param_v4_i8_irri_param_0]; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_irri_param_1]; +; CHECK-NEXT: prmt.b32 %r2, %r1, 4, 0x3340U; +; CHECK-NEXT: ld.param.b8 %r3, [st_param_v4_i8_irri_param_0]; +; CHECK-NEXT: prmt.b32 %r4, 1, %r3, 0x3340U; +; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; ; CHECK-NEXT: { // callseq 30, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {1, %rs2, %rs1, 4}; +; CHECK-NEXT: st.param.b32 [param0], %r5; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 30 ; CHECK-NEXT: ret; @@ -593,14 +613,17 @@ define void @st_param_v4_i8_irri(i8 %b, i8 %c) { define void @st_param_v4_i8_riir(i8 %a, i8 %d) { ; CHECK-LABEL: st_param_v4_i8_riir( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<6>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_riir_param_1]; -; CHECK-NEXT: ld.param.b8 %rs2, [st_param_v4_i8_riir_param_0]; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_riir_param_1]; +; CHECK-NEXT: prmt.b32 %r2, 3, %r1, 0x3340U; +; CHECK-NEXT: ld.param.b8 %r3, [st_param_v4_i8_riir_param_0]; +; CHECK-NEXT: prmt.b32 %r4, %r3, 2, 0x3340U; +; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; ; CHECK-NEXT: { // callseq 31, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {%rs2, 2, 3, %rs1}; +; CHECK-NEXT: st.param.b32 [param0], %r5; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 31 ; CHECK-NEXT: ret; @@ -614,14 +637,17 @@ define void @st_param_v4_i8_riir(i8 %a, i8 %d) { define void @st_param_v4_i8_riri(i8 %a, i8 %c) { ; CHECK-LABEL: st_param_v4_i8_riri( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<6>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_riri_param_1]; -; CHECK-NEXT: ld.param.b8 %rs2, [st_param_v4_i8_riri_param_0]; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_riri_param_1]; +; CHECK-NEXT: prmt.b32 %r2, %r1, 4, 0x3340U; +; CHECK-NEXT: ld.param.b8 %r3, [st_param_v4_i8_riri_param_0]; +; CHECK-NEXT: prmt.b32 %r4, %r3, 2, 0x3340U; +; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; ; CHECK-NEXT: { // callseq 32, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {%rs2, 2, %rs1, 4}; +; CHECK-NEXT: st.param.b32 [param0], %r5; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 32 ; CHECK-NEXT: ret; @@ -635,14 +661,16 @@ define void @st_param_v4_i8_riri(i8 %a, i8 %c) { define void @st_param_v4_i8_rrii(i8 %a, i8 %b) { ; CHECK-LABEL: st_param_v4_i8_rrii( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<5>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_rrii_param_1]; -; CHECK-NEXT: ld.param.b8 %rs2, [st_param_v4_i8_rrii_param_0]; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_rrii_param_1]; +; CHECK-NEXT: ld.param.b8 %r2, [st_param_v4_i8_rrii_param_0]; +; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U; +; CHECK-NEXT: prmt.b32 %r4, %r3, 1027, 0x5410U; ; CHECK-NEXT: { // callseq 33, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {%rs2, %rs1, 3, 4}; +; CHECK-NEXT: st.param.b32 [param0], %r4; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 33 ; CHECK-NEXT: ret; @@ -656,13 +684,15 @@ define void @st_param_v4_i8_rrii(i8 %a, i8 %b) { define void @st_param_v4_i8_iiir(i8 %d) { ; CHECK-LABEL: st_param_v4_i8_iiir( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_iiir_param_0]; ; CHECK-NEXT: { // callseq 34, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {1, 2, 3, %rs1}; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_iiir_param_0]; +; CHECK-NEXT: prmt.b32 %r2, 3, %r1, 0x3340U; +; CHECK-NEXT: prmt.b32 %r3, 513, %r2, 0x5410U; +; CHECK-NEXT: st.param.b32 [param0], %r3; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 34 ; CHECK-NEXT: ret; @@ -676,13 +706,15 @@ define void @st_param_v4_i8_iiir(i8 %d) { define void @st_param_v4_i8_iiri(i8 %c) { ; CHECK-LABEL: st_param_v4_i8_iiri( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_iiri_param_0]; ; CHECK-NEXT: { // callseq 35, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {1, 2, %rs1, 4}; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_iiri_param_0]; +; CHECK-NEXT: prmt.b32 %r2, %r1, 4, 0x3340U; +; CHECK-NEXT: prmt.b32 %r3, 513, %r2, 0x5410U; +; CHECK-NEXT: st.param.b32 [param0], %r3; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 35 ; CHECK-NEXT: ret; @@ -696,13 +728,15 @@ define void @st_param_v4_i8_iiri(i8 %c) { define void @st_param_v4_i8_irii(i8 %b) { ; CHECK-LABEL: st_param_v4_i8_irii( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_irii_param_0]; ; CHECK-NEXT: { // callseq 36, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {1, %rs1, 3, 4}; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_irii_param_0]; +; CHECK-NEXT: prmt.b32 %r2, 1, %r1, 0x3340U; +; CHECK-NEXT: prmt.b32 %r3, %r2, 1027, 0x5410U; +; CHECK-NEXT: st.param.b32 [param0], %r3; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 36 ; CHECK-NEXT: ret; @@ -716,13 +750,15 @@ define void @st_param_v4_i8_irii(i8 %b) { define void @st_param_v4_i8_riii(i8 %a) { ; CHECK-LABEL: st_param_v4_i8_riii( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b32 %r<4>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [st_param_v4_i8_riii_param_0]; ; CHECK-NEXT: { // callseq 37, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.v4.b8 [param0], {%rs1, 2, 3, 4}; +; CHECK-NEXT: ld.param.b8 %r1, [st_param_v4_i8_riii_param_0]; +; CHECK-NEXT: prmt.b32 %r2, %r1, 2, 0x3340U; +; CHECK-NEXT: prmt.b32 %r3, %r2, 1027, 0x5410U; +; CHECK-NEXT: st.param.b32 [param0], %r3; ; CHECK-NEXT: call.uni call_v4_i8, (param0); ; CHECK-NEXT: } // callseq 37 ; CHECK-NEXT: ret; @@ -742,7 +778,7 @@ define void @st_param_v4_i16_iiii() { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: { // callseq 38, 0 ; CHECK-NEXT: .param .align 8 .b8 param0[8]; -; CHECK-NEXT: st.param.v4.b16 [param0], {1, 2, 3, 4}; +; CHECK-NEXT: st.param.v2.b32 [param0], {131073, 262147}; ; CHECK-NEXT: call.uni call_v4_i16, (param0); ; CHECK-NEXT: } // callseq 38 ; CHECK-NEXT: ret; @@ -841,13 +877,15 @@ define void @st_param_v4_i16_iirr(i16 %c, i16 %d) { ; CHECK-LABEL: st_param_v4_i16_iirr( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<2>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_iirr_param_0]; ; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_iirr_param_1]; +; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2}; ; CHECK-NEXT: { // callseq 43, 0 ; CHECK-NEXT: .param .align 8 .b8 param0[8]; -; CHECK-NEXT: st.param.v4.b16 [param0], {1, 2, %rs1, %rs2}; +; CHECK-NEXT: st.param.v2.b32 [param0], {131073, %r1}; ; CHECK-NEXT: call.uni call_v4_i16, (param0); ; CHECK-NEXT: } // callseq 43 ; CHECK-NEXT: ret; @@ -946,13 +984,15 @@ define void @st_param_v4_i16_rrii(i16 %a, i16 %b) { ; CHECK-LABEL: st_param_v4_i16_rrii( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<2>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_rrii_param_0]; ; CHECK-NEXT: ld.param.b16 %rs2, [st_param_v4_i16_rrii_param_1]; +; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2}; ; CHECK-NEXT: { // callseq 48, 0 ; CHECK-NEXT: .param .align 8 .b8 param0[8]; -; CHECK-NEXT: st.param.v4.b16 [param0], {%rs1, %rs2, 3, 4}; +; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, 262147}; ; CHECK-NEXT: call.uni call_v4_i16, (param0); ; CHECK-NEXT: } // callseq 48 ; CHECK-NEXT: ret; @@ -966,13 +1006,16 @@ define void @st_param_v4_i16_rrii(i16 %a, i16 %b) { define void @st_param_v4_i16_iiir(i16 %d) { ; CHECK-LABEL: st_param_v4_i16_iiir( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<2>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_iiir_param_0]; +; CHECK-NEXT: mov.b16 %rs2, 3; +; CHECK-NEXT: mov.b32 %r1, {%rs2, %rs1}; ; CHECK-NEXT: { // callseq 49, 0 ; CHECK-NEXT: .param .align 8 .b8 param0[8]; -; CHECK-NEXT: st.param.v4.b16 [param0], {1, 2, 3, %rs1}; +; CHECK-NEXT: st.param.v2.b32 [param0], {131073, %r1}; ; CHECK-NEXT: call.uni call_v4_i16, (param0); ; CHECK-NEXT: } // callseq 49 ; CHECK-NEXT: ret; @@ -986,13 +1029,16 @@ define void @st_param_v4_i16_iiir(i16 %d) { define void @st_param_v4_i16_iiri(i16 %c) { ; CHECK-LABEL: st_param_v4_i16_iiri( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<2>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_iiri_param_0]; +; CHECK-NEXT: mov.b16 %rs2, 4; +; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2}; ; CHECK-NEXT: { // callseq 50, 0 ; CHECK-NEXT: .param .align 8 .b8 param0[8]; -; CHECK-NEXT: st.param.v4.b16 [param0], {1, 2, %rs1, 4}; +; CHECK-NEXT: st.param.v2.b32 [param0], {131073, %r1}; ; CHECK-NEXT: call.uni call_v4_i16, (param0); ; CHECK-NEXT: } // callseq 50 ; CHECK-NEXT: ret; @@ -1006,13 +1052,16 @@ define void @st_param_v4_i16_iiri(i16 %c) { define void @st_param_v4_i16_irii(i16 %b) { ; CHECK-LABEL: st_param_v4_i16_irii( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<2>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_irii_param_0]; +; CHECK-NEXT: mov.b16 %rs2, 1; +; CHECK-NEXT: mov.b32 %r1, {%rs2, %rs1}; ; CHECK-NEXT: { // callseq 51, 0 ; CHECK-NEXT: .param .align 8 .b8 param0[8]; -; CHECK-NEXT: st.param.v4.b16 [param0], {1, %rs1, 3, 4}; +; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, 262147}; ; CHECK-NEXT: call.uni call_v4_i16, (param0); ; CHECK-NEXT: } // callseq 51 ; CHECK-NEXT: ret; @@ -1026,13 +1075,16 @@ define void @st_param_v4_i16_irii(i16 %b) { define void @st_param_v4_i16_riii(i16 %a) { ; CHECK-LABEL: st_param_v4_i16_riii( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b32 %r<2>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b16 %rs1, [st_param_v4_i16_riii_param_0]; +; CHECK-NEXT: mov.b16 %rs2, 2; +; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2}; ; CHECK-NEXT: { // callseq 52, 0 ; CHECK-NEXT: .param .align 8 .b8 param0[8]; -; CHECK-NEXT: st.param.v4.b16 [param0], {%rs1, 2, 3, 4}; +; CHECK-NEXT: st.param.v2.b32 [param0], {%r1, 262147}; ; CHECK-NEXT: call.uni call_v4_i16, (param0); ; CHECK-NEXT: } // callseq 52 ; CHECK-NEXT: ret; @@ -1672,13 +1724,12 @@ declare void @call_v4_f32(%struct.float4 alignstack(16)) define void @st_param_bfloat() { ; CHECK-LABEL: st_param_bfloat( ; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; +; CHECK-EMPTY: ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: mov.b16 %rs1, 0x4100; ; CHECK-NEXT: { // callseq 83, 0 ; CHECK-NEXT: .param .align 2 .b8 param0[2]; -; CHECK-NEXT: st.param.b16 [param0], %rs1; +; CHECK-NEXT: st.param.b16 [param0], 0x4100; ; CHECK-NEXT: call.uni call_bfloat, (param0); ; CHECK-NEXT: } // callseq 83 ; CHECK-NEXT: ret; |