[clang] [llvm] [NVPTX] Fix v2i8 call lowering, use generic ld/st nodes for call params (PR #146930)

[Princeton Ferro via llvm-commits]

================
@@ -442,15 +442,19 @@ 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<8>;
 ; 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:    mov.b32 %r5, 1;
+; CHECK-NEXT:    prmt.b32 %r6, %r5, %r4, 0x3340U;
+; CHECK-NEXT:    prmt.b32 %r7, %r6, %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], %r7;
----------------
Prince781 wrote:

What is causing this pattern here? Can this be improved?

https://github.com/llvm/llvm-project/pull/146930