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

[JOE1994]

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

[llvmbot]

@llvm/pr-subscribers-debuginfo

Author: Youngsuk Kim (JOE1994)

Changes

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

---

Full diff: https://github.com/llvm/llvm-project/pull/89004.diff

5 Files Affected:

• (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+1-2)
• (modified) llvm/test/CodeGen/NVPTX/dynamic_stackalloc.ll (-2)
• (modified) llvm/test/CodeGen/NVPTX/i8x4-instructions.ll (-3)
• (modified) llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll (-1)
• (modified) llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected (-1)

diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 4daf47cfd4df8a..cd8546005c0289 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/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),
diff --git a/llvm/test/CodeGen/NVPTX/dynamic_stackalloc.ll b/llvm/test/CodeGen/NVPTX/dynamic_stackalloc.ll
index 09297fb819ce55..ce81957f2a3934 100644
--- a/llvm/test/CodeGen/NVPTX/dynamic_stackalloc.ll
+++ b/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]];
 
diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
index 6895699a1dfea1..96a4359d0ec43e 100644
--- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
+++ b/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];
diff --git a/llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll b/llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll
index 8d7608ead38edb..de367dfa4acb4c 100644
--- a/llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll
+++ b/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
diff --git a/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected b/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected
index a65f140b460159..3ac63d070933dd 100644
--- a/llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected
+++ b/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};

[Artem-B]

Good catch. LGTM.