[llvm][NVPTX] Don't emit unused var 'temp_param_reg' (NFC) (#89004)

Don't emit unused variable 'temp_param_reg' which has been around since
ae556d3 .

Author: JOE1994

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

@@ -3785,8 +3785,7 @@ class Pseudo<dag outs, dag ins, string asmstr, list<dag> pattern>
 
 def Callseq_Start :
   NVPTXInst<(outs), (ins i32imm:$amt1, i32imm:$amt2),
-            "\\{ // callseq $amt1, $amt2\n"
-            "\t.reg .b32 temp_param_reg;",
+            "\\{ // callseq $amt1, $amt2",
             [(callseq_start timm:$amt1, timm:$amt2)]>;
 def Callseq_End :
   NVPTXInst<(outs), (ins i32imm:$amt1, i32imm:$amt2),

llvm/test/CodeGen/NVPTX/dynamic_stackalloc.ll

@@ -17,7 +17,6 @@
 ; CHECK-32-NEXT:  alloca.u32  %r[[ALLOCA:[0-9]]], %r[[SIZE3]], 16;
 ; CHECK-32-NEXT:  cvta.local.u32  %r[[ALLOCA]], %r[[ALLOCA]];
 ; CHECK-32-NEXT:  { // callseq 0, 0
-; CHECK-32-NEXT:  .reg .b32 temp_param_reg;
 ; CHECK-32-NEXT:  .param .b32 param0;
 ; CHECK-32-NEXT:  st.param.b32  [param0+0], %r[[ALLOCA]];
 
@@ -27,7 +26,6 @@
 ; CHECK-64-NEXT:  alloca.u64  %rd[[ALLOCA:[0-9]]], %rd[[SIZE3]], 16;
 ; CHECK-64-NEXT:  cvta.local.u64  %rd[[ALLOCA]], %rd[[ALLOCA]];
 ; CHECK-64-NEXT:  { // callseq 0, 0
-; CHECK-64-NEXT:  .reg .b32 temp_param_reg;
 ; CHECK-64-NEXT:  .param .b64 param0;
 ; CHECK-64-NEXT:  st.param.b64  [param0+0], %rd[[ALLOCA]];
 

llvm/test/CodeGen/NVPTX/i8x4-instructions.ll

@@ -827,7 +827,6 @@ define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 {
 ; CHECK-NEXT:    ld.param.u32 %r2, [test_call_param_1];
 ; CHECK-NEXT:    ld.param.u32 %r1, [test_call_param_0];
 ; CHECK-NEXT:    { // callseq 0, 0
-; CHECK-NEXT:    .reg .b32 temp_param_reg;
 ; CHECK-NEXT:    .param .align 4 .b8 param0[4];
 ; CHECK-NEXT:    st.param.b32 [param0+0], %r1;
 ; CHECK-NEXT:    .param .align 4 .b8 param1[4];
@@ -856,7 +855,6 @@ define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
 ; CHECK-NEXT:    ld.param.u32 %r2, [test_call_flipped_param_1];
 ; CHECK-NEXT:    ld.param.u32 %r1, [test_call_flipped_param_0];
 ; CHECK-NEXT:    { // callseq 1, 0
-; CHECK-NEXT:    .reg .b32 temp_param_reg;
 ; CHECK-NEXT:    .param .align 4 .b8 param0[4];
 ; CHECK-NEXT:    st.param.b32 [param0+0], %r2;
 ; CHECK-NEXT:    .param .align 4 .b8 param1[4];
@@ -885,7 +883,6 @@ define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
 ; CHECK-NEXT:    ld.param.u32 %r2, [test_tailcall_flipped_param_1];
 ; CHECK-NEXT:    ld.param.u32 %r1, [test_tailcall_flipped_param_0];
 ; CHECK-NEXT:    { // callseq 2, 0
-; CHECK-NEXT:    .reg .b32 temp_param_reg;
 ; CHECK-NEXT:    .param .align 4 .b8 param0[4];
 ; CHECK-NEXT:    st.param.b32 [param0+0], %r2;
 ; CHECK-NEXT:    .param .align 4 .b8 param1[4];

llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll

@@ -9,7 +9,6 @@
 ; CHECK: add.u64 %rd1, %SP, 0;
 ; CHECK: .loc 1 5 3                   // t.c:5:3
 ; CHECK: { // callseq 0, 0
-; CHECK: .reg .b32 temp_param_reg;
 ; CHECK: .param .b64 param0;
 ; CHECK: st.param.b64 [param0+0], %rd1;
 ; CHECK: call.uni

llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected

@@ -30,7 +30,6 @@ define dso_local void @caller_St8x4(ptr nocapture noundef readonly byval(%struct
 ; CHECK-NEXT:    ld.u64 %rd7, [%SP+24];
 ; CHECK-NEXT:    ld.u64 %rd8, [%SP+16];
 ; CHECK-NEXT:    { // callseq 0, 0
-; CHECK-NEXT:    .reg .b32 temp_param_reg;
 ; CHECK-NEXT:    .param .align 16 .b8 param0[32];
 ; CHECK-NEXT:    st.param.v2.b64 [param0+0], {%rd6, %rd5};
 ; CHECK-NEXT:    st.param.v2.b64 [param0+16], {%rd8, %rd7};