[clang] [llvm] [NVPTX] Fix for device function pointer having the same name as generated callprototype (PR #205639)
Daniel Donenfeld via cfe-commits
cfe-commits at lists.llvm.org
Wed Sep 9 14:49:03 PDT 2026
https://github.com/daniel-donenfeld updated https://github.com/llvm/llvm-project/pull/205639
>From ea6a8d8078433a60c6e53e17a54ed2e8909434cc Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Tue, 2 Jun 2026 20:15:39 +0000
Subject: [PATCH 01/13] Bug fix for device function pointer having same name as
a generated callprototype
---
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 2 +-
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 2 +-
llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll | 2 +-
llvm/test/CodeGen/NVPTX/callchain.ll | 2 +-
.../NVPTX/callprototype-local-label.ll | 17 ++++++++++
.../CodeGen/NVPTX/convert-call-to-indirect.ll | 16 +++++-----
llvm/test/CodeGen/NVPTX/indirect_byval.ll | 8 ++---
.../CodeGen/NVPTX/lower-args-gridconstant.ll | 32 +++++++++----------
llvm/test/CodeGen/NVPTX/noreturn.ll | 4 +--
llvm/test/CodeGen/NVPTX/param-align.ll | 10 +++---
10 files changed, 56 insertions(+), 39 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/callprototype-local-label.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index e2b55b3be4dc0..6a51abcc23521 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -737,7 +737,7 @@ void NVPTXAsmPrinter::emitCallPrototype(const CallBase &CB,
const auto PtrVT = TLI->getPointerTy(DL);
Type *RetTy = CB.getFunctionType()->getReturnType();
- O << "prototype_" << UniqueCallSite << " : .callprototype ";
+ O << "$L__prototype_" << UniqueCallSite << " : .callprototype ";
if (RetTy->isVoidTy() || RetTy->isEmptyTy()) {
O << "()";
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 846bcc0557464..87cc6f41d60d0 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -1985,7 +1985,7 @@ foreach is_convergent = [0, 1] in {
NVPTXInst<(outs),
(ins ADDR_base:$addr, CallOperand:$rets, CallOperand:$params,
i32imm:$proto),
- "call${rets:RetList} $addr, (${params:ParamList}), prototype_$proto;">;
+ "call${rets:RetList} $addr, (${params:ParamList}), $$L__prototype_$proto;">;
def CALL_UNI # convergent_suffix :
NVPTXInst<(outs),
diff --git a/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll b/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
index 2a36c4162ba02..3cd64105b65c3 100644
--- a/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
@@ -13,7 +13,7 @@ target triple = "nvptx64-nvidia-cuda"
%"class.sycl::_V1::detail::half_impl::half" = type { half }
%complex_half = type { half, half }
-; CHECK: prototype_0 : .callprototype (.param .align 2 .b8 _[4]) _ (.param .b32 _, .param .b32 _, .param .align 2 .b8 _[4]);
+; CHECK: $L__prototype_0 : .callprototype (.param .align 2 .b8 _[4]) _ (.param .b32 _, .param .b32 _, .param .align 2 .b8 _[4]);
; CHECK: .param .align 2 .b8 param2[4];
; CHECK: .param .align 2 .b8 retval0[4];
; CHECK-DAG: st.param.b16 [param2], %rs{{[0-9]+}};
diff --git a/llvm/test/CodeGen/NVPTX/callchain.ll b/llvm/test/CodeGen/NVPTX/callchain.ll
index 847d8b80ad38a..be61a9fbe3291 100644
--- a/llvm/test/CodeGen/NVPTX/callchain.ll
+++ b/llvm/test/CodeGen/NVPTX/callchain.ll
@@ -4,7 +4,7 @@
target triple = "nvptx"
define void @foo(ptr %ptr) {
-; CHECK: prototype_0 : .callprototype ()_ ()
+; CHECK: $L__prototype_0 : .callprototype ()_ ()
tail call void %ptr()
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/callprototype-local-label.ll b/llvm/test/CodeGen/NVPTX/callprototype-local-label.ll
new file mode 100644
index 0000000000000..371282a28745c
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/callprototype-local-label.ll
@@ -0,0 +1,17 @@
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_75 | FileCheck %s
+; RUN: %if ptxas-sm_75 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_75 | %ptxas-verify --compile-only -arch=sm_75 %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+ at prototype_0 = addrspace(1) global ptr null, align 8
+
+define i32 @call_via_prototype_0(i32 %a, i32 %b, i32 %c, i32 %d) {
+; CHECK-LABEL: call_via_prototype_0(
+; CHECK: ld.global.{{u|b}}64 {{%rd[0-9]+}}, [prototype_0];
+; CHECK: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _, .param .b32 _, .param .b32 _);
+; CHECK-NEXT: call (retval0), %rd{{[0-9]+}}, (param0, param1, param2, param3), $L__prototype_0;
+; CHECK-NOT: prototype_0 : .callprototype
+ %fp = load ptr, ptr addrspace(1) @prototype_0, align 8
+ %ret = call i32 %fp(i32 %a, i32 %b, i32 %c, i32 %d)
+ ret i32 %ret
+}
diff --git a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
index 79be7ed2f18f1..e04c7fdef81df 100644
--- a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
@@ -10,7 +10,7 @@ define %struct.64 @test_return_type_mismatch(ptr %p) {
; CHECK-LABEL: test_return_type_mismatch(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<32>;
-; CHECK-NEXT: prototype_0 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _);
+; CHECK-NEXT: $L__prototype_0 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_return_type_mismatch_param_0];
; CHECK-NEXT: { // callseq 0, 0
@@ -18,7 +18,7 @@ define %struct.64 @test_return_type_mismatch(ptr %p) {
; CHECK-NEXT: .param .align 1 .b8 retval0[8];
; CHECK-NEXT: st.param.b64 [param0], %rd1;
; CHECK-NEXT: mov.b64 %rd2, callee;
-; CHECK-NEXT: call (retval0), %rd2, (param0), prototype_0;
+; CHECK-NEXT: call (retval0), %rd2, (param0), $L__prototype_0;
; CHECK-NEXT: ld.param.b8 %rd3, [retval0+7];
; CHECK-NEXT: ld.param.b8 %rd4, [retval0+6];
; CHECK-NEXT: ld.param.b8 %rd5, [retval0+5];
@@ -66,14 +66,14 @@ define i64 @test_param_type_mismatch(ptr %p) {
; CHECK-LABEL: test_param_type_mismatch(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-NEXT: prototype_1 : .callprototype (.param .b64 _) _ (.param .b64 _);
+; CHECK-NEXT: $L__prototype_1 : .callprototype (.param .b64 _) _ (.param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: { // callseq 1, 0
; CHECK-NEXT: .param .b64 param0;
; CHECK-NEXT: .param .b64 retval0;
; CHECK-NEXT: st.param.b64 [param0], 7;
; CHECK-NEXT: mov.b64 %rd1, callee;
-; CHECK-NEXT: call (retval0), %rd1, (param0), prototype_1;
+; CHECK-NEXT: call (retval0), %rd1, (param0), $L__prototype_1;
; CHECK-NEXT: ld.param.b64 %rd2, [retval0];
; CHECK-NEXT: } // callseq 1
; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
@@ -86,7 +86,7 @@ define i64 @test_param_count_mismatch(ptr %p) {
; CHECK-LABEL: test_param_count_mismatch(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<4>;
-; CHECK-NEXT: prototype_2 : .callprototype (.param .b64 _) _ (.param .b64 _, .param .b64 _);
+; CHECK-NEXT: $L__prototype_2 : .callprototype (.param .b64 _) _ (.param .b64 _, .param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_param_count_mismatch_param_0];
; CHECK-NEXT: { // callseq 2, 0
@@ -96,7 +96,7 @@ define i64 @test_param_count_mismatch(ptr %p) {
; CHECK-NEXT: st.param.b64 [param0], %rd1;
; CHECK-NEXT: st.param.b64 [param1], 7;
; CHECK-NEXT: mov.b64 %rd2, callee;
-; CHECK-NEXT: call (retval0), %rd2, (param0, param1), prototype_2;
+; CHECK-NEXT: call (retval0), %rd2, (param0, param1), $L__prototype_2;
; CHECK-NEXT: ld.param.b64 %rd3, [retval0];
; CHECK-NEXT: } // callseq 2
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
@@ -112,7 +112,7 @@ define %struct.64 @test_return_type_mismatch_variadic(ptr %p) {
; CHECK-NEXT: .reg .b64 %SP;
; CHECK-NEXT: .reg .b64 %SPL;
; CHECK-NEXT: .reg .b64 %rd<33>;
-; CHECK-NEXT: prototype_3 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _, .param .b64 _);
+; CHECK-NEXT: $L__prototype_3 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _, .param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: mov.b64 %SPL, __local_depot3;
; CHECK-NEXT: ld.param.b64 %rd1, [test_return_type_mismatch_variadic_param_0];
@@ -124,7 +124,7 @@ define %struct.64 @test_return_type_mismatch_variadic(ptr %p) {
; CHECK-NEXT: st.param.b64 [param1], %rd2;
; CHECK-NEXT: st.param.b64 [param0], %rd1;
; CHECK-NEXT: mov.b64 %rd3, callee_variadic;
-; CHECK-NEXT: call (retval0), %rd3, (param0, param1), prototype_3;
+; CHECK-NEXT: call (retval0), %rd3, (param0, param1), $L__prototype_3;
; CHECK-NEXT: ld.param.b8 %rd4, [retval0+7];
; CHECK-NEXT: ld.param.b8 %rd5, [retval0+6];
; CHECK-NEXT: ld.param.b8 %rd6, [retval0+5];
diff --git a/llvm/test/CodeGen/NVPTX/indirect_byval.ll b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
index 7a13bcd1f6674..a5a394386dc76 100644
--- a/llvm/test/CodeGen/NVPTX/indirect_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
@@ -18,7 +18,7 @@ define internal i32 @foo() {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-NEXT: .reg .b64 %rd<4>;
-; CHECK-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .align 1 .b8 _[1], .param .b64 _);
+; CHECK-NEXT: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .align 1 .b8 _[1], .param .b64 _);
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: mov.b64 %SPL, __local_depot0;
; CHECK-NEXT: add.u64 %rd1, %SPL, 0;
@@ -31,7 +31,7 @@ define internal i32 @foo() {
; CHECK-NEXT: st.param.b64 [param1], %rd2;
; CHECK-NEXT: ld.local.b8 %rs1, [%SPL+1];
; CHECK-NEXT: st.param.b8 [param0], %rs1;
-; CHECK-NEXT: call (retval0), %rd3, (param0, param1), prototype_0;
+; CHECK-NEXT: call (retval0), %rd3, (param0, param1), $L__prototype_0;
; CHECK-NEXT: ld.param.b32 %r1, [retval0];
; CHECK-NEXT: } // callseq 0
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
@@ -53,7 +53,7 @@ define internal i32 @bar() {
; CHECK-NEXT: .reg .b64 %SPL;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-NEXT: .reg .b64 %rd<5>;
-; CHECK-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .align 8 .b8 _[8], .param .b64 _);
+; CHECK-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .align 8 .b8 _[8], .param .b64 _);
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: mov.b64 %SPL, __local_depot1;
; CHECK-NEXT: add.u64 %rd1, %SPL, 0;
@@ -66,7 +66,7 @@ define internal i32 @bar() {
; CHECK-NEXT: st.param.b64 [param1], %rd2;
; CHECK-NEXT: ld.local.b64 %rd4, [%SPL+8];
; CHECK-NEXT: st.param.b64 [param0], %rd4;
-; CHECK-NEXT: call (retval0), %rd3, (param0, param1), prototype_1;
+; CHECK-NEXT: call (retval0), %rd3, (param0, param1), $L__prototype_1;
; CHECK-NEXT: ld.param.b32 %r1, [retval0];
; CHECK-NEXT: } // callseq 1
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
diff --git a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
index 5b23f009fea6c..90e1e29410027 100644
--- a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
@@ -154,7 +154,7 @@ define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.gri
; PTX-DEFAULT-LABEL: grid_const_escape(
; PTX-DEFAULT: {
; PTX-DEFAULT-NEXT: .reg .b64 %rd<4>;
-; PTX-DEFAULT-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-DEFAULT-NEXT: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-DEFAULT-NEXT: // %bb.0:
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_escape_param_0;
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd2, %rd1;
@@ -163,7 +163,7 @@ define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.gri
; PTX-DEFAULT-NEXT: .param .b32 retval0;
; PTX-DEFAULT-NEXT: st.param.b64 [param0], %rd2;
; PTX-DEFAULT-NEXT: mov.b64 %rd3, escape;
-; PTX-DEFAULT-NEXT: call (retval0), %rd3, (param0), prototype_0;
+; PTX-DEFAULT-NEXT: call (retval0), %rd3, (param0), $L__prototype_0;
; PTX-DEFAULT-NEXT: } // callseq 0
; PTX-DEFAULT-NEXT: ret;
;
@@ -171,7 +171,7 @@ define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.gri
; PTX-SHORT-PTR: {
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<2>;
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<4>;
-; PTX-SHORT-PTR-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-SHORT-PTR-NEXT: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-SHORT-PTR-NEXT: // %bb.0:
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_escape_param_0;
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd1, %r1;
@@ -181,7 +181,7 @@ define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.gri
; PTX-SHORT-PTR-NEXT: .param .b32 retval0;
; PTX-SHORT-PTR-NEXT: st.param.b64 [param0], %rd2;
; PTX-SHORT-PTR-NEXT: mov.b64 %rd3, escape;
-; PTX-SHORT-PTR-NEXT: call (retval0), %rd3, (param0), prototype_0;
+; PTX-SHORT-PTR-NEXT: call (retval0), %rd3, (param0), $L__prototype_0;
; PTX-SHORT-PTR-NEXT: } // callseq 0
; PTX-SHORT-PTR-NEXT: ret;
; OPT-LABEL: define ptx_kernel void @grid_const_escape(
@@ -201,7 +201,7 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-DEFAULT-NEXT: .reg .b64 %SPL;
; PTX-DEFAULT-NEXT: .reg .b32 %r<2>;
; PTX-DEFAULT-NEXT: .reg .b64 %rd<8>;
-; PTX-DEFAULT-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
+; PTX-DEFAULT-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
; PTX-DEFAULT-NEXT: // %bb.0:
; PTX-DEFAULT-NEXT: mov.b64 %SPL, __local_depot4;
; PTX-DEFAULT-NEXT: mov.b64 %rd1, multiple_grid_const_escape_param_0;
@@ -221,7 +221,7 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-DEFAULT-NEXT: st.param.b64 [param1], %rd6;
; PTX-DEFAULT-NEXT: st.param.b64 [param0], %rd4;
; PTX-DEFAULT-NEXT: mov.b64 %rd7, escape3;
-; PTX-DEFAULT-NEXT: call (retval0), %rd7, (param0, param1, param2), prototype_1;
+; PTX-DEFAULT-NEXT: call (retval0), %rd7, (param0, param1, param2), $L__prototype_1;
; PTX-DEFAULT-NEXT: } // callseq 1
; PTX-DEFAULT-NEXT: ret;
;
@@ -232,7 +232,7 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-SHORT-PTR-NEXT: .reg .b32 %SPL;
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<5>;
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<8>;
-; PTX-SHORT-PTR-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
+; PTX-SHORT-PTR-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
; PTX-SHORT-PTR-NEXT: // %bb.0:
; PTX-SHORT-PTR-NEXT: mov.b32 %SPL, __local_depot4;
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, multiple_grid_const_escape_param_0;
@@ -255,7 +255,7 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-SHORT-PTR-NEXT: st.param.b64 [param1], %rd6;
; PTX-SHORT-PTR-NEXT: st.param.b64 [param0], %rd4;
; PTX-SHORT-PTR-NEXT: mov.b64 %rd7, escape3;
-; PTX-SHORT-PTR-NEXT: call (retval0), %rd7, (param0, param1, param2), prototype_1;
+; PTX-SHORT-PTR-NEXT: call (retval0), %rd7, (param0, param1, param2), $L__prototype_1;
; PTX-SHORT-PTR-NEXT: } // callseq 1
; PTX-SHORT-PTR-NEXT: ret;
; OPT-LABEL: define ptx_kernel void @multiple_grid_const_escape(
@@ -362,7 +362,7 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g
; PTX-DEFAULT: {
; PTX-DEFAULT-NEXT: .reg .b32 %r<3>;
; PTX-DEFAULT-NEXT: .reg .b64 %rd<6>;
-; PTX-DEFAULT-NEXT: prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-DEFAULT-NEXT: $L__prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-DEFAULT-NEXT: // %bb.0:
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_partial_escape_param_0;
; PTX-DEFAULT-NEXT: ld.param.b64 %rd2, [grid_const_partial_escape_param_1];
@@ -376,7 +376,7 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g
; PTX-DEFAULT-NEXT: .param .b32 retval0;
; PTX-DEFAULT-NEXT: st.param.b64 [param0], %rd4;
; PTX-DEFAULT-NEXT: mov.b64 %rd5, escape;
-; PTX-DEFAULT-NEXT: call (retval0), %rd5, (param0), prototype_2;
+; PTX-DEFAULT-NEXT: call (retval0), %rd5, (param0), $L__prototype_2;
; PTX-DEFAULT-NEXT: } // callseq 2
; PTX-DEFAULT-NEXT: ret;
;
@@ -384,7 +384,7 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g
; PTX-SHORT-PTR: {
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<4>;
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<6>;
-; PTX-SHORT-PTR-NEXT: prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-SHORT-PTR-NEXT: $L__prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-SHORT-PTR-NEXT: // %bb.0:
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_partial_escape_param_0;
; PTX-SHORT-PTR-NEXT: ld.param.b64 %rd1, [grid_const_partial_escape_param_1];
@@ -399,7 +399,7 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g
; PTX-SHORT-PTR-NEXT: .param .b32 retval0;
; PTX-SHORT-PTR-NEXT: st.param.b64 [param0], %rd4;
; PTX-SHORT-PTR-NEXT: mov.b64 %rd5, escape;
-; PTX-SHORT-PTR-NEXT: call (retval0), %rd5, (param0), prototype_2;
+; PTX-SHORT-PTR-NEXT: call (retval0), %rd5, (param0), $L__prototype_2;
; PTX-SHORT-PTR-NEXT: } // callseq 2
; PTX-SHORT-PTR-NEXT: ret;
; OPT-LABEL: define ptx_kernel void @grid_const_partial_escape(
@@ -422,7 +422,7 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4
; PTX-DEFAULT: {
; PTX-DEFAULT-NEXT: .reg .b32 %r<4>;
; PTX-DEFAULT-NEXT: .reg .b64 %rd<6>;
-; PTX-DEFAULT-NEXT: prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-DEFAULT-NEXT: $L__prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-DEFAULT-NEXT: // %bb.0:
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_partial_escapemem_param_0;
; PTX-DEFAULT-NEXT: ld.param.b64 %rd2, [grid_const_partial_escapemem_param_1];
@@ -437,7 +437,7 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4
; PTX-DEFAULT-NEXT: .param .b32 retval0;
; PTX-DEFAULT-NEXT: st.param.b64 [param0], %rd4;
; PTX-DEFAULT-NEXT: mov.b64 %rd5, escape;
-; PTX-DEFAULT-NEXT: call (retval0), %rd5, (param0), prototype_3;
+; PTX-DEFAULT-NEXT: call (retval0), %rd5, (param0), $L__prototype_3;
; PTX-DEFAULT-NEXT: } // callseq 3
; PTX-DEFAULT-NEXT: st.param.b32 [func_retval0], %r3;
; PTX-DEFAULT-NEXT: ret;
@@ -446,7 +446,7 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4
; PTX-SHORT-PTR: {
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<5>;
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<6>;
-; PTX-SHORT-PTR-NEXT: prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-SHORT-PTR-NEXT: $L__prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-SHORT-PTR-NEXT: // %bb.0:
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_partial_escapemem_param_0;
; PTX-SHORT-PTR-NEXT: ld.param.b64 %rd1, [grid_const_partial_escapemem_param_1];
@@ -462,7 +462,7 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4
; PTX-SHORT-PTR-NEXT: .param .b32 retval0;
; PTX-SHORT-PTR-NEXT: st.param.b64 [param0], %rd4;
; PTX-SHORT-PTR-NEXT: mov.b64 %rd5, escape;
-; PTX-SHORT-PTR-NEXT: call (retval0), %rd5, (param0), prototype_3;
+; PTX-SHORT-PTR-NEXT: call (retval0), %rd5, (param0), $L__prototype_3;
; PTX-SHORT-PTR-NEXT: } // callseq 3
; PTX-SHORT-PTR-NEXT: st.param.b32 [func_retval0], %r4;
; PTX-SHORT-PTR-NEXT: ret;
diff --git a/llvm/test/CodeGen/NVPTX/noreturn.ll b/llvm/test/CodeGen/NVPTX/noreturn.ll
index 0062e62756d36..416418c7ffde7 100644
--- a/llvm/test/CodeGen/NVPTX/noreturn.ll
+++ b/llvm/test/CodeGen/NVPTX/noreturn.ll
@@ -32,8 +32,8 @@ define ptx_kernel void @ignore_kernel_noreturn() #0 {
}
; CHECK-LABEL: .entry callprototype_noreturn(
-; CHECK: prototype_{{[0-9]+}} : .callprototype ()_ (.param .b32 _) .noreturn;
-; CHECK: prototype_{{[0-9]+}} : .callprototype (.param .b32 _) _ (.param .b32 _);
+; CHECK: {{[$]}}L__prototype_{{[0-9]+}} : .callprototype ()_ (.param .b32 _) .noreturn;
+; CHECK: {{[$]}}L__prototype_{{[0-9]+}} : .callprototype (.param .b32 _) _ (.param .b32 _);
define ptx_kernel void @callprototype_noreturn(i32) {
%fn = load ptr, ptr addrspace(1) @function_pointer
diff --git a/llvm/test/CodeGen/NVPTX/param-align.ll b/llvm/test/CodeGen/NVPTX/param-align.ll
index c85080fdf295a..7f0a1646962ae 100644
--- a/llvm/test/CodeGen/NVPTX/param-align.ll
+++ b/llvm/test/CodeGen/NVPTX/param-align.ll
@@ -55,17 +55,17 @@ declare ptr @getfp(i32 %n)
define ptx_device void @t6() {
; CHECK: .func t6
%fp = call ptr @getfp(i32 0)
-; CHECK: prototype_2 : .callprototype ()_ (.param .align 8 .b8 _[8]);
+; CHECK: $L__prototype_2 : .callprototype ()_ (.param .align 8 .b8 _[8]);
call void %fp(ptr byval(double) null);
%fp2 = call ptr @getfp(i32 1)
-; NOALIGN4: prototype_4 : .callprototype ()_ (.param .align 2 .b8 _[4]);
-; ALIGN4: prototype_4 : .callprototype ()_ (.param .align 4 .b8 _[4]);
+; NOALIGN4: $L__prototype_4 : .callprototype ()_ (.param .align 2 .b8 _[4]);
+; ALIGN4: $L__prototype_4 : .callprototype ()_ (.param .align 4 .b8 _[4]);
call void %fp(ptr byval(%struct.half2) null);
%fp3 = call ptr @getfp(i32 2)
-; NOALIGN4: prototype_6 : .callprototype ()_ (.param .align 1 .b8 _[1]);
-; ALIGN4: prototype_6 : .callprototype ()_ (.param .align 4 .b8 _[1]);
+; NOALIGN4: $L__prototype_6 : .callprototype ()_ (.param .align 1 .b8 _[1]);
+; ALIGN4: $L__prototype_6 : .callprototype ()_ (.param .align 4 .b8 _[1]);
call void %fp(ptr byval(i8) null);
ret void
}
>From 193c47ea7a4d799fd3af33a197e27ad2c74bebd7 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Wed, 24 Jun 2026 21:34:42 +0000
Subject: [PATCH 02/13] Fix missed test
---
llvm/test/CodeGen/NVPTX/vaargs.ll | 8 ++++----
1 file changed, 4 insertions(+), 4 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/vaargs.ll b/llvm/test/CodeGen/NVPTX/vaargs.ll
index 8ff429a16b74c..34ff8c18e2ca0 100644
--- a/llvm/test/CodeGen/NVPTX/vaargs.ll
+++ b/llvm/test/CodeGen/NVPTX/vaargs.ll
@@ -216,7 +216,7 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK32-NEXT: .reg .b32 %SPL;
; CHECK32-NEXT: .reg .b32 %r<9>;
; CHECK32-NEXT: .reg .b64 %rd<3>;
-; CHECK32-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _);
+; CHECK32-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _);
; CHECK32-NEXT: // %bb.0: // %entry
; CHECK32-NEXT: mov.b32 %SPL, __local_depot1;
; CHECK32-NEXT: ld.param.b32 %r2, [test_foo_param_3];
@@ -238,7 +238,7 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK32-NEXT: .param .b32 retval0;
; CHECK32-NEXT: st.param.b32 [param1], %r3;
; CHECK32-NEXT: st.param.b32 [param0], 4;
-; CHECK32-NEXT: call (retval0), %r7, (param0, param1), prototype_1;
+; CHECK32-NEXT: call (retval0), %r7, (param0, param1), $L__prototype_1;
; CHECK32-NEXT: ld.param.b32 %r8, [retval0];
; CHECK32-NEXT: } // callseq 1
; CHECK32-NEXT: st.param.b32 [func_retval0], %r8;
@@ -251,7 +251,7 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK64-NEXT: .reg .b64 %SPL;
; CHECK64-NEXT: .reg .b32 %r<3>;
; CHECK64-NEXT: .reg .b64 %rd<9>;
-; CHECK64-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b64 _);
+; CHECK64-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b64 _);
; CHECK64-NEXT: // %bb.0: // %entry
; CHECK64-NEXT: mov.b64 %SPL, __local_depot1;
; CHECK64-NEXT: ld.param.b64 %rd3, [test_foo_param_3];
@@ -273,7 +273,7 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK64-NEXT: .param .b32 retval0;
; CHECK64-NEXT: st.param.b64 [param1], %rd4;
; CHECK64-NEXT: st.param.b32 [param0], 4;
-; CHECK64-NEXT: call (retval0), %rd8, (param0, param1), prototype_1;
+; CHECK64-NEXT: call (retval0), %rd8, (param0, param1), $L__prototype_1;
; CHECK64-NEXT: ld.param.b32 %r2, [retval0];
; CHECK64-NEXT: } // callseq 1
; CHECK64-NEXT: st.param.b32 [func_retval0], %r2;
>From 72f3d46c4ba9bec3448a5e75fee07eecc1cb6a80 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Mon, 29 Jun 2026 19:54:26 +0000
Subject: [PATCH 03/13] Add cuda C++ lit test for prototype collision
---
.../CodeGenCUDA/callprototype-local-label.cu | 29 +++++++++++++++++++
1 file changed, 29 insertions(+)
create mode 100644 clang/test/CodeGenCUDA/callprototype-local-label.cu
diff --git a/clang/test/CodeGenCUDA/callprototype-local-label.cu b/clang/test/CodeGenCUDA/callprototype-local-label.cu
new file mode 100644
index 0000000000000..6f4bb2de0b02a
--- /dev/null
+++ b/clang/test/CodeGenCUDA/callprototype-local-label.cu
@@ -0,0 +1,29 @@
+// REQUIRES: nvptx-registered-target
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_75 \
+// RUN: -fcuda-is-device -S -o - -x cuda %s \
+// RUN: | FileCheck %s
+
+// Test that a global named 'prototype_0' does not cause a callprototype label
+// collision: the label must be '$L__prototype_0', not 'prototype_0'.
+// extern "C" is used to keep the PTX global name unmangled.
+
+#define __device__ __attribute__((device))
+
+extern "C" {
+
+__device__ int simple_func() { return 42; }
+
+__device__ int (*prototype_0)(int, int, int, int) = nullptr;
+__device__ int call_via_prototype_0(int a, int b, int c, int d) {
+ if (prototype_0 != nullptr)
+ return prototype_0(a, b, c, d);
+ return a + b + c + d;
+}
+
+} // extern "C"
+
+// CHECK: .visible .global .align 8 .u64 prototype_0;
+// CHECK-LABEL: .visible .func (.param .b32 func_retval0) call_via_prototype_0(
+// CHECK: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _, .param .b32 _, .param .b32 _);
+// CHECK-NEXT: call (retval0), %rd{{[0-9]+}}, (param0, param1, param2, param3), $L__prototype_0;
+// CHECK-NOT: prototype_0 : .callprototype
>From 18ad8e90a058e9530f6822b40ac58a8d25fd32f3 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Mon, 29 Jun 2026 20:31:48 +0000
Subject: [PATCH 04/13] Fix tests after rebase
---
clang/test/CodeGenCUDA/callprototype-local-label.cu | 2 +-
llvm/test/CodeGen/NVPTX/callprototype-local-label.ll | 4 ++--
2 files changed, 3 insertions(+), 3 deletions(-)
diff --git a/clang/test/CodeGenCUDA/callprototype-local-label.cu b/clang/test/CodeGenCUDA/callprototype-local-label.cu
index 6f4bb2de0b02a..9e5f3d868a6d5 100644
--- a/clang/test/CodeGenCUDA/callprototype-local-label.cu
+++ b/clang/test/CodeGenCUDA/callprototype-local-label.cu
@@ -25,5 +25,5 @@ __device__ int call_via_prototype_0(int a, int b, int c, int d) {
// CHECK: .visible .global .align 8 .u64 prototype_0;
// CHECK-LABEL: .visible .func (.param .b32 func_retval0) call_via_prototype_0(
// CHECK: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _, .param .b32 _, .param .b32 _);
-// CHECK-NEXT: call (retval0), %rd{{[0-9]+}}, (param0, param1, param2, param3), $L__prototype_0;
+// CHECK: call (retval0), %rd{{[0-9]+}}, (param0, param1, param2, param3), $L__prototype_0;
// CHECK-NOT: prototype_0 : .callprototype
diff --git a/llvm/test/CodeGen/NVPTX/callprototype-local-label.ll b/llvm/test/CodeGen/NVPTX/callprototype-local-label.ll
index 371282a28745c..894a5dc97d11b 100644
--- a/llvm/test/CodeGen/NVPTX/callprototype-local-label.ll
+++ b/llvm/test/CodeGen/NVPTX/callprototype-local-label.ll
@@ -7,9 +7,9 @@ target triple = "nvptx64-nvidia-cuda"
define i32 @call_via_prototype_0(i32 %a, i32 %b, i32 %c, i32 %d) {
; CHECK-LABEL: call_via_prototype_0(
-; CHECK: ld.global.{{u|b}}64 {{%rd[0-9]+}}, [prototype_0];
; CHECK: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _, .param .b32 _, .param .b32 _);
-; CHECK-NEXT: call (retval0), %rd{{[0-9]+}}, (param0, param1, param2, param3), $L__prototype_0;
+; CHECK: ld.global.{{u|b}}64 {{%rd[0-9]+}}, [prototype_0];
+; CHECK: call (retval0), %rd{{[0-9]+}}, (param0, param1, param2, param3), $L__prototype_0;
; CHECK-NOT: prototype_0 : .callprototype
%fp = load ptr, ptr addrspace(1) @prototype_0, align 8
%ret = call i32 %fp(i32 %a, i32 %b, i32 %c, i32 %d)
>From 7fdc60a94530ac8ce73b4e00829842f659f0ae3a Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Tue, 18 Aug 2026 17:29:07 +0000
Subject: [PATCH 05/13] Add initial way of using MCSymbol to unique the call
prototype name
---
llvm/lib/Target/NVPTX/CMakeLists.txt | 1 +
llvm/lib/Target/NVPTX/NVPTX.h | 3 +
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 15 ++--
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 6 +-
.../Target/NVPTX/NVPTXMachineFunctionInfo.h | 25 +++++-
.../NVPTX/NVPTXPreEmitSymbolLowering.cpp | 90 +++++++++++++++++++
llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp | 5 ++
.../NVPTX/callprototype-dollar-local-label.ll | 19 ++++
8 files changed, 153 insertions(+), 11 deletions(-)
create mode 100644 llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
create mode 100644 llvm/test/CodeGen/NVPTX/callprototype-dollar-local-label.ll
diff --git a/llvm/lib/Target/NVPTX/CMakeLists.txt b/llvm/lib/Target/NVPTX/CMakeLists.txt
index 8eb3471fa4fb2..1f72b6bb04f7a 100644
--- a/llvm/lib/Target/NVPTX/CMakeLists.txt
+++ b/llvm/lib/Target/NVPTX/CMakeLists.txt
@@ -36,6 +36,7 @@ set(NVPTXCodeGen_sources
NVPTXMCExpr.cpp
NVPTXMarkKernelPtrsGlobal.cpp
NVPTXPeephole.cpp
+ NVPTXPreEmitSymbolLowering.cpp
NVPTXPrologEpilogPass.cpp
NVPTXPromoteParamAlign.cpp
NVPTXProxyRegErasure.cpp
diff --git a/llvm/lib/Target/NVPTX/NVPTX.h b/llvm/lib/Target/NVPTX/NVPTX.h
index 6897baaede9c4..02ae074066a93 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.h
+++ b/llvm/lib/Target/NVPTX/NVPTX.h
@@ -67,6 +67,8 @@ MachineFunctionPass *createNVPTXPeepholeLegacyPass();
MachineFunctionPass *createNVPTXProxyRegErasureLegacyPass();
MachineFunctionPass *createNVPTXForwardParamsLegacyPass();
MachineFunctionPass *createNVPTXAddressFolderLegacyPass();
+MachineFunctionPass *
+createNVPTXPreEmitSymbolLoweringLegacyPass(const TargetMachine &TM);
void initializeNVVMReflectLegacyPassPass(PassRegistry &);
void initializeGenericToNVVMLegacyPassPass(PassRegistry &);
@@ -74,6 +76,7 @@ void initializeNVPTXAllocaHoistingLegacyPassPass(PassRegistry &);
void initializeNVPTXAsmPrinterPass(PassRegistry &);
void initializeNVPTXAssignValidGlobalNamesLegacyPassPass(PassRegistry &);
void initializeNVPTXAtomicLowerLegacyPassPass(PassRegistry &);
+void initializeNVPTXPreEmitSymbolLoweringLegacyPass(PassRegistry &);
void initializeNVPTXCtorDtorLoweringLegacyPass(PassRegistry &);
void initializeNVPTXLowerAggrCopiesLegacyPassPass(PassRegistry &);
void initializeNVPTXLowerAllocaLegacyPassPass(PassRegistry &);
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 6a51abcc23521..0392e4789a25f 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -253,7 +253,7 @@ class NVPTXAsmPrinter : public AsmPrinter {
void encodeDebugInfoRegisterNumbers(const MachineFunction &MF);
void printReturnValStr(const Function *, raw_ostream &O);
void printReturnValStr(const MachineFunction &MF, raw_ostream &O);
- void emitCallPrototype(const CallBase &CB, unsigned UniqueCallSite,
+ void emitCallPrototype(const CallBase &CB, MCSymbol *PrototypeSymbol,
raw_ostream &O) const;
void emitJumpTable(const MachineJumpTableEntry &MJT, unsigned MJTI) const;
@@ -605,6 +605,8 @@ MCOperand NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO) {
case MachineOperand::MO_MachineBasicBlock:
return MCOperand::createExpr(
MCSymbolRefExpr::create(MO.getMBB()->getSymbol(), OutContext));
+ case MachineOperand::MO_MCSymbol:
+ return GetSymbolRef(MO.getMCSymbol());
case MachineOperand::MO_ExternalSymbol:
return GetSymbolRef(GetExternalSymbolSymbol(MO.getSymbolName()));
case MachineOperand::MO_MCSymbol:
@@ -729,7 +731,7 @@ void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,
}
void NVPTXAsmPrinter::emitCallPrototype(const CallBase &CB,
- unsigned UniqueCallSite,
+ MCSymbol *PrototypeSymbol,
raw_ostream &O) const {
const DataLayout &DL = getDataLayout();
const NVPTXSubtarget &STI = MF->getSubtarget<NVPTXSubtarget>();
@@ -737,7 +739,8 @@ void NVPTXAsmPrinter::emitCallPrototype(const CallBase &CB,
const auto PtrVT = TLI->getPointerTy(DL);
Type *RetTy = CB.getFunctionType()->getReturnType();
- O << "$L__prototype_" << UniqueCallSite << " : .callprototype ";
+ PrototypeSymbol->print(O, MAI);
+ O << " : .callprototype ";
if (RetTy->isVoidTy() || RetTy->isEmptyTy()) {
O << "()";
@@ -939,8 +942,10 @@ void NVPTXAsmPrinter::emitFunctionBodyStart() {
emitDemotedVars(&MF->getFunction(), O);
const auto *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
- for (const auto &[Id, CB] : MFI->getCallPrototypes())
- emitCallPrototype(*CB, Id, O);
+ for (const auto &Entry : MFI->getCallPrototypes()) {
+ const auto &Prototype = Entry.second;
+ emitCallPrototype(*Prototype.CB, Prototype.Symbol, O);
+ }
OutStreamer->emitRawText(O.str());
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 87cc6f41d60d0..e771b0f86631b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -1983,9 +1983,9 @@ foreach is_convergent = [0, 1] in {
let isCall = 1, isConvergent = is_convergent in {
def CALL # convergent_suffix :
NVPTXInst<(outs),
- (ins ADDR_base:$addr, CallOperand:$rets, CallOperand:$params,
- i32imm:$proto),
- "call${rets:RetList} $addr, (${params:ParamList}), $$L__prototype_$proto;">;
+ (ins ADDR_base:$addr, CallOperand:$rets, CallOperand:$params,
+ ADDR_base:$proto),
+ "call${rets:RetList} $addr, (${params:ParamList}), $proto;">;
def CALL_UNI # convergent_suffix :
NVPTXInst<(outs),
diff --git a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
index c55ccf592f6fb..f8102344c03e4 100644
--- a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
+++ b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
@@ -23,6 +23,12 @@ class CallBase;
class MCSymbol;
class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
+public:
+ struct CallPrototype {
+ const CallBase *CB;
+ MCSymbol *Symbol = nullptr;
+ };
+
private:
/// The parameter symbols whose image handles were replaced with image
/// references.
@@ -30,7 +36,7 @@ class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
/// Stores a mapping from a unique call-site id to the call instruction that
/// needs an indirect-call prototype emitted.
- std::map<unsigned, const CallBase *> CallPrototypes;
+ std::map<unsigned, CallPrototype> CallPrototypes;
public:
NVPTXMachineFunctionInfo(const Function &F, const TargetSubtargetInfo *STI) {}
@@ -53,10 +59,23 @@ class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
}
void addCallPrototype(unsigned Id, const CallBase *CB) {
- CallPrototypes.try_emplace(Id, CB);
+ CallPrototypes.try_emplace(Id, CallPrototype{CB});
+ }
+
+ void setCallPrototypeSymbol(unsigned Id, MCSymbol *Symbol) {
+ auto It = CallPrototypes.find(Id);
+ assert(It != CallPrototypes.end() && "unknown call prototype");
+ It->second.Symbol = Symbol;
+ }
+
+ MCSymbol *getCallPrototypeSymbol(unsigned Id) const {
+ auto It = CallPrototypes.find(Id);
+ assert(It != CallPrototypes.end() && "unknown call prototype");
+ assert(It->second.Symbol && "call prototype symbol not set");
+ return It->second.Symbol;
}
- const std::map<unsigned, const CallBase *> &getCallPrototypes() const {
+ const std::map<unsigned, CallPrototype> &getCallPrototypes() const {
return CallPrototypes;
}
};
diff --git a/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
new file mode 100644
index 0000000000000..722279bbe06e7
--- /dev/null
+++ b/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
@@ -0,0 +1,90 @@
+//===-- NVPTXPreEmitSymbolLowering.cpp - Lower symbols before emission -*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "NVPTX.h"
+#include "NVPTXMachineFunctionInfo.h"
+#include "NVPTXRegisterInfo.h"
+#include "llvm/CodeGen/MachineFunction.h"
+#include "llvm/CodeGen/MachineFunctionPass.h"
+#include "llvm/CodeGen/MachineInstr.h"
+#include "llvm/CodeGen/MachineOperand.h"
+#include "llvm/IR/GlobalValue.h"
+#include "llvm/IR/Module.h"
+#include "llvm/MC/MCSymbol.h"
+#include "llvm/Target/TargetMachine.h"
+
+using namespace llvm;
+
+namespace {
+
+class NVPTXPreEmitSymbolLoweringLegacy : public MachineFunctionPass {
+ const TargetMachine *TM = nullptr;
+
+public:
+ static char ID;
+
+ NVPTXPreEmitSymbolLoweringLegacy() : MachineFunctionPass(ID) {}
+ NVPTXPreEmitSymbolLoweringLegacy(const TargetMachine &TM)
+ : MachineFunctionPass(ID), TM(&TM) {}
+
+ bool doInitialization(Module &M) override {
+ assert(TM && "TargetMachine must be set");
+ for (const GlobalValue &GV : M.global_values())
+ TM->getSymbol(&GV);
+ return false;
+ }
+
+ bool runOnMachineFunction(MachineFunction &MF) override {
+ auto *MFI = MF.getInfo<NVPTXMachineFunctionInfo>();
+ if (MFI->getCallPrototypes().empty())
+ return false;
+
+ for (const auto &[Id, Prototype] : MFI->getCallPrototypes()) {
+ if (!Prototype.Symbol) {
+ MCSymbol *Symbol = MF.getContext().createTempSymbol(
+ "prototype_" + Twine(Id), /*AlwaysAddSuffix=*/false);
+ MFI->setCallPrototypeSymbol(Id, Symbol);
+ }
+ }
+
+ bool Changed = false;
+ for (MachineBasicBlock &MBB : MF) {
+ for (MachineInstr &MI : MBB) {
+ if (MI.getOpcode() != NVPTX::CALL && MI.getOpcode() != NVPTX::CALL_conv)
+ continue;
+
+ MachineOperand &Proto = MI.getOperand(3);
+ if (Proto.isImm()) {
+ Proto.ChangeToMCSymbol(MFI->getCallPrototypeSymbol(Proto.getImm()));
+ Changed = true;
+ }
+ assert(Proto.isMCSymbol() &&
+ "call prototype operand must be rewritten to a symbol");
+ }
+ }
+
+ return Changed;
+ }
+
+ void getAnalysisUsage(AnalysisUsage &AU) const override {
+ AU.setPreservesCFG();
+ MachineFunctionPass::getAnalysisUsage(AU);
+ }
+};
+
+} // namespace
+
+char NVPTXPreEmitSymbolLoweringLegacy::ID = 0;
+
+INITIALIZE_PASS(NVPTXPreEmitSymbolLoweringLegacy, "nvptx-pre-emit-symbol-lowering",
+ "NVPTX Pre-Emit Symbol Lowering", false, false)
+
+MachineFunctionPass *
+llvm::createNVPTXPreEmitSymbolLoweringLegacyPass(const TargetMachine &TM) {
+ return new NVPTXPreEmitSymbolLoweringLegacy(TM);
+}
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
index f4987a8dbfbba..bd12f448a0618 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -135,6 +135,7 @@ class NVPTXPassConfig : public TargetPassConfig {
bool addInstSelector() override;
void addPreRegAlloc() override;
void addPostRegAlloc() override;
+ void addPreEmitPass2() override;
FunctionPass *createTargetRegisterAllocator(bool) override;
void addFastRegAlloc() override;
@@ -355,6 +356,10 @@ void NVPTXPassConfig::addPostRegAlloc() {
}
}
+void NVPTXPassConfig::addPreEmitPass2() {
+ addPass(createNVPTXPreEmitSymbolLoweringLegacyPass(getNVPTXTargetMachine()));
+}
+
FunctionPass *NVPTXPassConfig::createTargetRegisterAllocator(bool) {
return nullptr; // No reg alloc
}
diff --git a/llvm/test/CodeGen/NVPTX/callprototype-dollar-local-label.ll b/llvm/test/CodeGen/NVPTX/callprototype-dollar-local-label.ll
new file mode 100644
index 0000000000000..f0effa497d3ef
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/callprototype-dollar-local-label.ll
@@ -0,0 +1,19 @@
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_75 | FileCheck %s
+; RUN: %if ptxas-sm_75 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_75 | %ptxas-verify --compile-only -arch=sm_75 %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+@"$L__prototype_0" = addrspace(1) global ptr null, align 8
+@"$L__prototype_00" = addrspace(1) global i32 7, align 4
+
+define i32 @call_via_dollar_prototype_0(i32 %a, i32 %b, i32 %c, i32 %d) {
+; CHECK-DAG: .visible .global .align 8 .u64 $L__prototype_0;
+; CHECK-DAG: .visible .global .align 4 .u32 $L__prototype_00 = 7;
+; CHECK-LABEL: call_via_dollar_prototype_0(
+; CHECK: $L__prototype_01 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _, .param .b32 _, .param .b32 _);
+; CHECK: ld.global.{{u|b}}64 {{%rd[0-9]+}}, [$L__prototype_0];
+; CHECK: call (retval0), %rd{{[0-9]+}}, (param0, param1, param2, param3), $L__prototype_01;
+ %fp = load ptr, ptr addrspace(1) @"$L__prototype_0", align 8
+ %ret = call i32 %fp(i32 %a, i32 %b, i32 %c, i32 %d)
+ ret i32 %ret
+}
>From 7df29ea75567941e8241331091d31e0c2bee5e82 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Wed, 26 Aug 2026 21:51:16 +0000
Subject: [PATCH 06/13] Fix build issue
---
llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
index 722279bbe06e7..6463f9e18348f 100644
--- a/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
@@ -8,13 +8,13 @@
#include "NVPTX.h"
#include "NVPTXMachineFunctionInfo.h"
-#include "NVPTXRegisterInfo.h"
#include "llvm/CodeGen/MachineFunction.h"
#include "llvm/CodeGen/MachineFunctionPass.h"
#include "llvm/CodeGen/MachineInstr.h"
#include "llvm/CodeGen/MachineOperand.h"
#include "llvm/IR/GlobalValue.h"
#include "llvm/IR/Module.h"
+#include "llvm/MC/MCContext.h"
#include "llvm/MC/MCSymbol.h"
#include "llvm/Target/TargetMachine.h"
>From 29ee14e070d3dc4517300ee80bc3e90c90b9a098 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Wed, 26 Aug 2026 22:04:23 +0000
Subject: [PATCH 07/13] Fix formatting
---
llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp | 5 +++--
1 file changed, 3 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
index 6463f9e18348f..50340e0c92017 100644
--- a/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
@@ -1,4 +1,4 @@
-//===-- NVPTXPreEmitSymbolLowering.cpp - Lower symbols before emission -*- C++ -*-===//
+//===-- NVPTXPreEmitSymbolLowering.cpp - Lower symbols before emission ----===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -81,7 +81,8 @@ class NVPTXPreEmitSymbolLoweringLegacy : public MachineFunctionPass {
char NVPTXPreEmitSymbolLoweringLegacy::ID = 0;
-INITIALIZE_PASS(NVPTXPreEmitSymbolLoweringLegacy, "nvptx-pre-emit-symbol-lowering",
+INITIALIZE_PASS(NVPTXPreEmitSymbolLoweringLegacy,
+ "nvptx-pre-emit-symbol-lowering",
"NVPTX Pre-Emit Symbol Lowering", false, false)
MachineFunctionPass *
>From 6e34d0f102d7ae66490a82428ff1b3ae50d6cb1e Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Tue, 8 Sep 2026 16:43:21 +0000
Subject: [PATCH 08/13] Move logic to NVPTXAsmPrinter
---
llvm/lib/Target/NVPTX/CMakeLists.txt | 1 -
llvm/lib/Target/NVPTX/NVPTX.h | 3 -
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 39 +++++++-
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 2 +-
.../Target/NVPTX/NVPTXMachineFunctionInfo.h | 25 +----
.../NVPTX/NVPTXPreEmitSymbolLowering.cpp | 91 -------------------
llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp | 5 -
llvm/test/CodeGen/NVPTX/empty-type.ll | 4 +-
8 files changed, 40 insertions(+), 130 deletions(-)
delete mode 100644 llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
diff --git a/llvm/lib/Target/NVPTX/CMakeLists.txt b/llvm/lib/Target/NVPTX/CMakeLists.txt
index 1f72b6bb04f7a..8eb3471fa4fb2 100644
--- a/llvm/lib/Target/NVPTX/CMakeLists.txt
+++ b/llvm/lib/Target/NVPTX/CMakeLists.txt
@@ -36,7 +36,6 @@ set(NVPTXCodeGen_sources
NVPTXMCExpr.cpp
NVPTXMarkKernelPtrsGlobal.cpp
NVPTXPeephole.cpp
- NVPTXPreEmitSymbolLowering.cpp
NVPTXPrologEpilogPass.cpp
NVPTXPromoteParamAlign.cpp
NVPTXProxyRegErasure.cpp
diff --git a/llvm/lib/Target/NVPTX/NVPTX.h b/llvm/lib/Target/NVPTX/NVPTX.h
index 02ae074066a93..6897baaede9c4 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.h
+++ b/llvm/lib/Target/NVPTX/NVPTX.h
@@ -67,8 +67,6 @@ MachineFunctionPass *createNVPTXPeepholeLegacyPass();
MachineFunctionPass *createNVPTXProxyRegErasureLegacyPass();
MachineFunctionPass *createNVPTXForwardParamsLegacyPass();
MachineFunctionPass *createNVPTXAddressFolderLegacyPass();
-MachineFunctionPass *
-createNVPTXPreEmitSymbolLoweringLegacyPass(const TargetMachine &TM);
void initializeNVVMReflectLegacyPassPass(PassRegistry &);
void initializeGenericToNVVMLegacyPassPass(PassRegistry &);
@@ -76,7 +74,6 @@ void initializeNVPTXAllocaHoistingLegacyPassPass(PassRegistry &);
void initializeNVPTXAsmPrinterPass(PassRegistry &);
void initializeNVPTXAssignValidGlobalNamesLegacyPassPass(PassRegistry &);
void initializeNVPTXAtomicLowerLegacyPassPass(PassRegistry &);
-void initializeNVPTXPreEmitSymbolLoweringLegacyPass(PassRegistry &);
void initializeNVPTXCtorDtorLoweringLegacyPass(PassRegistry &);
void initializeNVPTXLowerAggrCopiesLegacyPassPass(PassRegistry &);
void initializeNVPTXLowerAllocaLegacyPassPass(PassRegistry &);
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 0392e4789a25f..99d5db629eb75 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -220,6 +220,7 @@ class NVPTXAsmPrinter : public AsmPrinter {
private:
const Function *F;
+ DenseMap<unsigned, MCSymbol *> CallPrototypeSymbols;
NVPTXTargetStreamer *getTargetStreamer() const;
@@ -232,6 +233,8 @@ class NVPTXAsmPrinter : public AsmPrinter {
void emitInstruction(const MachineInstr *) override;
void lowerToMCInst(const MachineInstr *MI, MCInst &OutMI);
+ MCOperand lowerCallPrototypeArg(const MachineOperand &MO);
+ MCOperand lowerOperand(const MachineInstr &MI, unsigned OpNum);
MCOperand lowerOperand(const MachineOperand &MO);
MCOperand GetSymbolRef(const MCSymbol *Symbol);
MCRegister encodeVirtualRegister(Register Reg);
@@ -590,8 +593,30 @@ void NVPTXAsmPrinter::emitInstruction(const MachineInstr *MI) {
void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) {
OutMI.setOpcode(MI->getOpcode());
- for (const auto MO : MI->operands())
- OutMI.addOperand(lowerOperand(MO));
+ for (unsigned I = 0, E = MI->getNumOperands(); I != E; ++I)
+ OutMI.addOperand(lowerOperand(*MI, I));
+}
+
+static bool isCallPrototypeOperand(const MachineInstr &MI,
+ unsigned OpNum) {
+ return OpNum == 3 &&
+ (MI.getOpcode() == NVPTX::CALL || MI.getOpcode() == NVPTX::CALL_conv);
+}
+
+MCOperand NVPTXAsmPrinter::lowerCallPrototypeArg(const MachineOperand &MO) {
+ assert(MO.isImm() && "call prototype operand must be an ID");
+ auto It = CallPrototypeSymbols.find(MO.getImm());
+ assert(It != CallPrototypeSymbols.end() && "unknown call prototype");
+ return GetSymbolRef(It->second);
+}
+
+MCOperand NVPTXAsmPrinter::lowerOperand(const MachineInstr &MI,
+ unsigned OpNum) {
+ const MachineOperand &MO = MI.getOperand(OpNum);
+ if (isCallPrototypeOperand(MI, OpNum))
+ return lowerCallPrototypeArg(MO);
+
+ return lowerOperand(MO);
}
MCOperand NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO) {
@@ -941,10 +966,13 @@ void NVPTXAsmPrinter::emitFunctionBodyStart() {
raw_svector_ostream O(Str);
emitDemotedVars(&MF->getFunction(), O);
+ CallPrototypeSymbols.clear();
const auto *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
- for (const auto &Entry : MFI->getCallPrototypes()) {
- const auto &Prototype = Entry.second;
- emitCallPrototype(*Prototype.CB, Prototype.Symbol, O);
+ for (const auto &[Id, CB] : MFI->getCallPrototypes()) {
+ MCSymbol *Symbol = OutContext.createTempSymbol(
+ "prototype_" + Twine(Id), /*AlwaysAddSuffix=*/false);
+ CallPrototypeSymbols.try_emplace(Id, Symbol);
+ emitCallPrototype(*CB, Symbol, O);
}
OutStreamer->emitRawText(O.str());
@@ -955,6 +983,7 @@ void NVPTXAsmPrinter::emitFunctionBodyStart() {
}
void NVPTXAsmPrinter::emitFunctionBodyEnd() {
+ CallPrototypeSymbols.clear();
VRegMapping.clear();
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index e771b0f86631b..d56a81abf1aa7 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -1984,7 +1984,7 @@ foreach is_convergent = [0, 1] in {
def CALL # convergent_suffix :
NVPTXInst<(outs),
(ins ADDR_base:$addr, CallOperand:$rets, CallOperand:$params,
- ADDR_base:$proto),
+ i32imm:$proto),
"call${rets:RetList} $addr, (${params:ParamList}), $proto;">;
def CALL_UNI # convergent_suffix :
diff --git a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
index f8102344c03e4..c55ccf592f6fb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
+++ b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
@@ -23,12 +23,6 @@ class CallBase;
class MCSymbol;
class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
-public:
- struct CallPrototype {
- const CallBase *CB;
- MCSymbol *Symbol = nullptr;
- };
-
private:
/// The parameter symbols whose image handles were replaced with image
/// references.
@@ -36,7 +30,7 @@ class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
/// Stores a mapping from a unique call-site id to the call instruction that
/// needs an indirect-call prototype emitted.
- std::map<unsigned, CallPrototype> CallPrototypes;
+ std::map<unsigned, const CallBase *> CallPrototypes;
public:
NVPTXMachineFunctionInfo(const Function &F, const TargetSubtargetInfo *STI) {}
@@ -59,23 +53,10 @@ class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
}
void addCallPrototype(unsigned Id, const CallBase *CB) {
- CallPrototypes.try_emplace(Id, CallPrototype{CB});
- }
-
- void setCallPrototypeSymbol(unsigned Id, MCSymbol *Symbol) {
- auto It = CallPrototypes.find(Id);
- assert(It != CallPrototypes.end() && "unknown call prototype");
- It->second.Symbol = Symbol;
- }
-
- MCSymbol *getCallPrototypeSymbol(unsigned Id) const {
- auto It = CallPrototypes.find(Id);
- assert(It != CallPrototypes.end() && "unknown call prototype");
- assert(It->second.Symbol && "call prototype symbol not set");
- return It->second.Symbol;
+ CallPrototypes.try_emplace(Id, CB);
}
- const std::map<unsigned, CallPrototype> &getCallPrototypes() const {
+ const std::map<unsigned, const CallBase *> &getCallPrototypes() const {
return CallPrototypes;
}
};
diff --git a/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
deleted file mode 100644
index 50340e0c92017..0000000000000
--- a/llvm/lib/Target/NVPTX/NVPTXPreEmitSymbolLowering.cpp
+++ /dev/null
@@ -1,91 +0,0 @@
-//===-- NVPTXPreEmitSymbolLowering.cpp - Lower symbols before emission ----===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-#include "NVPTX.h"
-#include "NVPTXMachineFunctionInfo.h"
-#include "llvm/CodeGen/MachineFunction.h"
-#include "llvm/CodeGen/MachineFunctionPass.h"
-#include "llvm/CodeGen/MachineInstr.h"
-#include "llvm/CodeGen/MachineOperand.h"
-#include "llvm/IR/GlobalValue.h"
-#include "llvm/IR/Module.h"
-#include "llvm/MC/MCContext.h"
-#include "llvm/MC/MCSymbol.h"
-#include "llvm/Target/TargetMachine.h"
-
-using namespace llvm;
-
-namespace {
-
-class NVPTXPreEmitSymbolLoweringLegacy : public MachineFunctionPass {
- const TargetMachine *TM = nullptr;
-
-public:
- static char ID;
-
- NVPTXPreEmitSymbolLoweringLegacy() : MachineFunctionPass(ID) {}
- NVPTXPreEmitSymbolLoweringLegacy(const TargetMachine &TM)
- : MachineFunctionPass(ID), TM(&TM) {}
-
- bool doInitialization(Module &M) override {
- assert(TM && "TargetMachine must be set");
- for (const GlobalValue &GV : M.global_values())
- TM->getSymbol(&GV);
- return false;
- }
-
- bool runOnMachineFunction(MachineFunction &MF) override {
- auto *MFI = MF.getInfo<NVPTXMachineFunctionInfo>();
- if (MFI->getCallPrototypes().empty())
- return false;
-
- for (const auto &[Id, Prototype] : MFI->getCallPrototypes()) {
- if (!Prototype.Symbol) {
- MCSymbol *Symbol = MF.getContext().createTempSymbol(
- "prototype_" + Twine(Id), /*AlwaysAddSuffix=*/false);
- MFI->setCallPrototypeSymbol(Id, Symbol);
- }
- }
-
- bool Changed = false;
- for (MachineBasicBlock &MBB : MF) {
- for (MachineInstr &MI : MBB) {
- if (MI.getOpcode() != NVPTX::CALL && MI.getOpcode() != NVPTX::CALL_conv)
- continue;
-
- MachineOperand &Proto = MI.getOperand(3);
- if (Proto.isImm()) {
- Proto.ChangeToMCSymbol(MFI->getCallPrototypeSymbol(Proto.getImm()));
- Changed = true;
- }
- assert(Proto.isMCSymbol() &&
- "call prototype operand must be rewritten to a symbol");
- }
- }
-
- return Changed;
- }
-
- void getAnalysisUsage(AnalysisUsage &AU) const override {
- AU.setPreservesCFG();
- MachineFunctionPass::getAnalysisUsage(AU);
- }
-};
-
-} // namespace
-
-char NVPTXPreEmitSymbolLoweringLegacy::ID = 0;
-
-INITIALIZE_PASS(NVPTXPreEmitSymbolLoweringLegacy,
- "nvptx-pre-emit-symbol-lowering",
- "NVPTX Pre-Emit Symbol Lowering", false, false)
-
-MachineFunctionPass *
-llvm::createNVPTXPreEmitSymbolLoweringLegacyPass(const TargetMachine &TM) {
- return new NVPTXPreEmitSymbolLoweringLegacy(TM);
-}
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
index bd12f448a0618..f4987a8dbfbba 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -135,7 +135,6 @@ class NVPTXPassConfig : public TargetPassConfig {
bool addInstSelector() override;
void addPreRegAlloc() override;
void addPostRegAlloc() override;
- void addPreEmitPass2() override;
FunctionPass *createTargetRegisterAllocator(bool) override;
void addFastRegAlloc() override;
@@ -356,10 +355,6 @@ void NVPTXPassConfig::addPostRegAlloc() {
}
}
-void NVPTXPassConfig::addPreEmitPass2() {
- addPass(createNVPTXPreEmitSymbolLoweringLegacyPass(getNVPTXTargetMachine()));
-}
-
FunctionPass *NVPTXPassConfig::createTargetRegisterAllocator(bool) {
return nullptr; // No reg alloc
}
diff --git a/llvm/test/CodeGen/NVPTX/empty-type.ll b/llvm/test/CodeGen/NVPTX/empty-type.ll
index f29c3d7c2e21b..04d546e246f47 100644
--- a/llvm/test/CodeGen/NVPTX/empty-type.ll
+++ b/llvm/test/CodeGen/NVPTX/empty-type.ll
@@ -132,13 +132,13 @@ define void @indirect_caller(ptr %fp) {
; CHECK-LABEL: indirect_caller(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: prototype_4 : .callprototype ()_ (.param .b32 _);
+; CHECK-NEXT: $L__prototype_4 : .callprototype ()_ (.param .b32 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [indirect_caller_param_0];
; CHECK-NEXT: { // callseq 4, 0
; CHECK-NEXT: .param .b32 param0;
; CHECK-NEXT: st.param.b32 [param0], 42;
-; CHECK-NEXT: call %rd1, (param0), prototype_4;
+; CHECK-NEXT: call %rd1, (param0), $L__prototype_4;
; CHECK-NEXT: } // callseq 4
; CHECK-NEXT: ret;
call void %fp(%empty zeroinitializer, i32 42)
>From c4eab5f5f2c42fedc5b73a96740bf06592bbec19 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Tue, 8 Sep 2026 16:49:08 +0000
Subject: [PATCH 09/13] Formatting
---
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 7 +++----
1 file changed, 3 insertions(+), 4 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 99d5db629eb75..937ea17bb027e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -597,8 +597,7 @@ void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) {
OutMI.addOperand(lowerOperand(*MI, I));
}
-static bool isCallPrototypeOperand(const MachineInstr &MI,
- unsigned OpNum) {
+static bool isCallPrototypeOperand(const MachineInstr &MI, unsigned OpNum) {
return OpNum == 3 &&
(MI.getOpcode() == NVPTX::CALL || MI.getOpcode() == NVPTX::CALL_conv);
}
@@ -969,8 +968,8 @@ void NVPTXAsmPrinter::emitFunctionBodyStart() {
CallPrototypeSymbols.clear();
const auto *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
for (const auto &[Id, CB] : MFI->getCallPrototypes()) {
- MCSymbol *Symbol = OutContext.createTempSymbol(
- "prototype_" + Twine(Id), /*AlwaysAddSuffix=*/false);
+ MCSymbol *Symbol = OutContext.createTempSymbol("prototype_" + Twine(Id),
+ /*AlwaysAddSuffix=*/false);
CallPrototypeSymbols.try_emplace(Id, Symbol);
emitCallPrototype(*CB, Symbol, O);
}
>From f741d3756618faaddd33ce81d44513b6917f3d31 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Tue, 8 Sep 2026 16:58:38 +0000
Subject: [PATCH 10/13] Fix rebase issue
---
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 2 --
1 file changed, 2 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 937ea17bb027e..1a7afdf35263d 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -633,8 +633,6 @@ MCOperand NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO) {
return GetSymbolRef(MO.getMCSymbol());
case MachineOperand::MO_ExternalSymbol:
return GetSymbolRef(GetExternalSymbolSymbol(MO.getSymbolName()));
- case MachineOperand::MO_MCSymbol:
- return GetSymbolRef(MO.getMCSymbol());
case MachineOperand::MO_JumpTableIndex:
// The jump table index names the .branchtargets list emitted for a brx.idx
// (see emitJumpTable); reference it by that label.
>From f117de2e2a4c5c8c36813a4425681f4d39c50060 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Wed, 9 Sep 2026 14:46:57 +0000
Subject: [PATCH 11/13] Address PR comments
---
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 26 ++++++++++---------
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 2 +-
.../Target/NVPTX/NVPTXMachineFunctionInfo.h | 14 +++++++---
3 files changed, 25 insertions(+), 17 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 1a7afdf35263d..57acb06b47172 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -220,7 +220,6 @@ class NVPTXAsmPrinter : public AsmPrinter {
private:
const Function *F;
- DenseMap<unsigned, MCSymbol *> CallPrototypeSymbols;
NVPTXTargetStreamer *getTargetStreamer() const;
@@ -604,9 +603,11 @@ static bool isCallPrototypeOperand(const MachineInstr &MI, unsigned OpNum) {
MCOperand NVPTXAsmPrinter::lowerCallPrototypeArg(const MachineOperand &MO) {
assert(MO.isImm() && "call prototype operand must be an ID");
- auto It = CallPrototypeSymbols.find(MO.getImm());
- assert(It != CallPrototypeSymbols.end() && "unknown call prototype");
- return GetSymbolRef(It->second);
+ const auto &CallPrototypes =
+ MF->getInfo<NVPTXMachineFunctionInfo>()->getCallPrototypes();
+ auto It = CallPrototypes.find(MO.getImm());
+ assert(It != CallPrototypes.end() && "unknown call prototype");
+ return GetSymbolRef(It->second.second);
}
MCOperand NVPTXAsmPrinter::lowerOperand(const MachineInstr &MI,
@@ -629,10 +630,10 @@ MCOperand NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO) {
case MachineOperand::MO_MachineBasicBlock:
return MCOperand::createExpr(
MCSymbolRefExpr::create(MO.getMBB()->getSymbol(), OutContext));
- case MachineOperand::MO_MCSymbol:
- return GetSymbolRef(MO.getMCSymbol());
case MachineOperand::MO_ExternalSymbol:
return GetSymbolRef(GetExternalSymbolSymbol(MO.getSymbolName()));
+ case MachineOperand::MO_MCSymbol:
+ return GetSymbolRef(MO.getMCSymbol());
case MachineOperand::MO_JumpTableIndex:
// The jump table index names the .branchtargets list emitted for a brx.idx
// (see emitJumpTable); reference it by that label.
@@ -963,12 +964,9 @@ void NVPTXAsmPrinter::emitFunctionBodyStart() {
raw_svector_ostream O(Str);
emitDemotedVars(&MF->getFunction(), O);
- CallPrototypeSymbols.clear();
const auto *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
- for (const auto &[Id, CB] : MFI->getCallPrototypes()) {
- MCSymbol *Symbol = OutContext.createTempSymbol("prototype_" + Twine(Id),
- /*AlwaysAddSuffix=*/false);
- CallPrototypeSymbols.try_emplace(Id, Symbol);
+ for (const auto &Entry : MFI->getCallPrototypes()) {
+ const auto &[CB, Symbol] = Entry.second;
emitCallPrototype(*CB, Symbol, O);
}
@@ -980,7 +978,6 @@ void NVPTXAsmPrinter::emitFunctionBodyStart() {
}
void NVPTXAsmPrinter::emitFunctionBodyEnd() {
- CallPrototypeSymbols.clear();
VRegMapping.clear();
}
@@ -1277,6 +1274,11 @@ bool NVPTXAsmPrinter::doInitialization(Module &M) {
GlobalsEmitted = false;
+ // Ensure globals are in the symbol table before ISel so any temp symbols are
+ // guaranteed not to collide with user symbols
+ for (const GlobalValue &GV : M.global_values())
+ TM.getSymbol(&GV);
+
return Result;
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index a946ceaa35d81..4277bb791e26e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1604,7 +1604,7 @@ SDValue NVPTXTargetLowering::LowerCall(TargetLowering::CallLoweringInfo &CLI,
if (IsIndirectCall)
DAG.getMachineFunction()
.getInfo<NVPTXMachineFunctionInfo>()
- ->addCallPrototype(UniqueCallSite, CB);
+ ->addCallPrototype(UniqueCallSite, CB, DAG.getMachineFunction());
const bool IsUnknownIntrinsic =
CalleeF && CalleeF->isIntrinsic() &&
diff --git a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
index c55ccf592f6fb..bcbd8fc8466d9 100644
--- a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
+++ b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
@@ -28,9 +28,10 @@ class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
/// references.
SmallPtrSet<const MCSymbol *, 8> ImageHandleSymbols;
+ using CallProtoTy = std::pair<const CallBase *, MCSymbol *>;
/// Stores a mapping from a unique call-site id to the call instruction that
/// needs an indirect-call prototype emitted.
- std::map<unsigned, const CallBase *> CallPrototypes;
+ std::map<unsigned, CallProtoTy> CallPrototypes;
public:
NVPTXMachineFunctionInfo(const Function &F, const TargetSubtargetInfo *STI) {}
@@ -52,11 +53,16 @@ class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
return ImageHandleSymbols.contains(Symbol);
}
- void addCallPrototype(unsigned Id, const CallBase *CB) {
- CallPrototypes.try_emplace(Id, CB);
+ void addCallPrototype(unsigned Id, const CallBase *CB, MachineFunction &MF) {
+ if (CallPrototypes.find(Id) == CallPrototypes.end()) {
+ MCSymbol *Symbol =
+ MF.getContext().createTempSymbol("prototype_" + Twine(Id),
+ /*AlwaysAddSuffix=*/false);
+ CallPrototypes.try_emplace(Id, CB, Symbol);
+ }
}
- const std::map<unsigned, const CallBase *> &getCallPrototypes() const {
+ const std::map<unsigned, CallProtoTy> &getCallPrototypes() const {
return CallPrototypes;
}
};
>From b2171ca4b8e7bf86178a83c180488f020b294055 Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Wed, 9 Sep 2026 18:11:07 +0000
Subject: [PATCH 12/13] Fix build issue
---
llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h | 1 +
1 file changed, 1 insertion(+)
diff --git a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
index bcbd8fc8466d9..672444111c95d 100644
--- a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
+++ b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
@@ -16,6 +16,7 @@
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/CodeGen/MachineFunction.h"
+#include "llvm/MC/MCContext.h"
#include <map>
namespace llvm {
>From 4a8eea057d346bd6080a4b5e89caa813b009e47d Mon Sep 17 00:00:00 2001
From: Daniel Donenfeld <ddonenfeld at nvidia.com>
Date: Wed, 9 Sep 2026 21:48:34 +0000
Subject: [PATCH 13/13] additional review changes
---
.../CodeGenCUDA/callprototype-local-label.cu | 5 ++-
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 44 +++++++++----------
.../Target/NVPTX/NVPTXMachineFunctionInfo.h | 5 +--
llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll | 3 +-
llvm/test/CodeGen/NVPTX/callchain.ll | 3 +-
.../NVPTX/callprototype-dollar-local-label.ll | 5 ++-
.../NVPTX/callprototype-local-label.ll | 5 ++-
.../CodeGen/NVPTX/convert-call-to-indirect.ll | 12 +++--
llvm/test/CodeGen/NVPTX/empty-type.ll | 5 ++-
llvm/test/CodeGen/NVPTX/indirect_byval.ll | 6 ++-
.../CodeGen/NVPTX/lower-args-gridconstant.ll | 24 ++++++----
llvm/test/CodeGen/NVPTX/noreturn.ll | 6 ++-
llvm/test/CodeGen/NVPTX/param-align.ll | 15 ++++---
llvm/test/CodeGen/NVPTX/vaargs.ll | 10 +++--
14 files changed, 87 insertions(+), 61 deletions(-)
diff --git a/clang/test/CodeGenCUDA/callprototype-local-label.cu b/clang/test/CodeGenCUDA/callprototype-local-label.cu
index 9e5f3d868a6d5..c9914f6b5bdaa 100644
--- a/clang/test/CodeGenCUDA/callprototype-local-label.cu
+++ b/clang/test/CodeGenCUDA/callprototype-local-label.cu
@@ -24,6 +24,7 @@ __device__ int call_via_prototype_0(int a, int b, int c, int d) {
// CHECK: .visible .global .align 8 .u64 prototype_0;
// CHECK-LABEL: .visible .func (.param .b32 func_retval0) call_via_prototype_0(
-// CHECK: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _, .param .b32 _, .param .b32 _);
+// CHECK: $L__prototype_0:
+// CHECK-NEXT: .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _, .param .b32 _, .param .b32 _);
// CHECK: call (retval0), %rd{{[0-9]+}}, (param0, param1, param2, param3), $L__prototype_0;
-// CHECK-NOT: prototype_0 : .callprototype
+// CHECK-NOT: prototype_0:
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 57acb06b47172..110f858be04c0 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -255,8 +255,7 @@ class NVPTXAsmPrinter : public AsmPrinter {
void encodeDebugInfoRegisterNumbers(const MachineFunction &MF);
void printReturnValStr(const Function *, raw_ostream &O);
void printReturnValStr(const MachineFunction &MF, raw_ostream &O);
- void emitCallPrototype(const CallBase &CB, MCSymbol *PrototypeSymbol,
- raw_ostream &O) const;
+ void emitCallPrototype(const CallBase &CB, MCSymbol *PrototypeSymbol) const;
void emitJumpTable(const MachineJumpTableEntry &MJT, unsigned MJTI) const;
/// Should a .noreturn directive be emitted for \p V, which is either a
@@ -592,13 +591,15 @@ void NVPTXAsmPrinter::emitInstruction(const MachineInstr *MI) {
void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) {
OutMI.setOpcode(MI->getOpcode());
- for (unsigned I = 0, E = MI->getNumOperands(); I != E; ++I)
- OutMI.addOperand(lowerOperand(*MI, I));
+ for (const auto &MO : MI->operands()) {
+ OutMI.addOperand(lowerOperand(MO));
+ }
}
-static bool isCallPrototypeOperand(const MachineInstr &MI, unsigned OpNum) {
- return OpNum == 3 &&
- (MI.getOpcode() == NVPTX::CALL || MI.getOpcode() == NVPTX::CALL_conv);
+static bool isCallPrototypeOperand(const MachineOperand &MO) {
+ return MO.getOperandNo() == 3 &&
+ (MO.getParent()->getOpcode() == NVPTX::CALL ||
+ MO.getParent()->getOpcode() == NVPTX::CALL_conv);
}
MCOperand NVPTXAsmPrinter::lowerCallPrototypeArg(const MachineOperand &MO) {
@@ -610,15 +611,6 @@ MCOperand NVPTXAsmPrinter::lowerCallPrototypeArg(const MachineOperand &MO) {
return GetSymbolRef(It->second.second);
}
-MCOperand NVPTXAsmPrinter::lowerOperand(const MachineInstr &MI,
- unsigned OpNum) {
- const MachineOperand &MO = MI.getOperand(OpNum);
- if (isCallPrototypeOperand(MI, OpNum))
- return lowerCallPrototypeArg(MO);
-
- return lowerOperand(MO);
-}
-
MCOperand NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO) {
switch (MO.getType()) {
default:
@@ -626,6 +618,8 @@ MCOperand NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO) {
case MachineOperand::MO_Register:
return MCOperand::createReg(encodeVirtualRegister(MO.getReg()));
case MachineOperand::MO_Immediate:
+ if (isCallPrototypeOperand(MO))
+ return lowerCallPrototypeArg(MO);
return MCOperand::createImm(MO.getImm());
case MachineOperand::MO_MachineBasicBlock:
return MCOperand::createExpr(
@@ -754,16 +748,19 @@ void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,
}
void NVPTXAsmPrinter::emitCallPrototype(const CallBase &CB,
- MCSymbol *PrototypeSymbol,
- raw_ostream &O) const {
+ MCSymbol *PrototypeSymbol) const {
const DataLayout &DL = getDataLayout();
const NVPTXSubtarget &STI = MF->getSubtarget<NVPTXSubtarget>();
const auto *TLI = cast<NVPTXTargetLowering>(STI.getTargetLowering());
const auto PtrVT = TLI->getPointerTy(DL);
Type *RetTy = CB.getFunctionType()->getReturnType();
- PrototypeSymbol->print(O, MAI);
- O << " : .callprototype ";
+ OutStreamer->emitLabel(PrototypeSymbol);
+
+ SmallString<128> Str;
+ raw_svector_ostream O(Str);
+
+ O << ".callprototype ";
if (RetTy->isVoidTy() || RetTy->isEmptyTy()) {
O << "()";
@@ -849,6 +846,8 @@ void NVPTXAsmPrinter::emitCallPrototype(const CallBase &CB,
if (shouldEmitPTXNoReturn(CB))
O << " .noreturn";
O << ";\n";
+
+ OutStreamer->emitRawText(O.str());
}
void NVPTXAsmPrinter::emitJumpTable(const MachineJumpTableEntry &MJT,
@@ -963,15 +962,14 @@ void NVPTXAsmPrinter::emitFunctionBodyStart() {
SmallString<128> Str;
raw_svector_ostream O(Str);
emitDemotedVars(&MF->getFunction(), O);
+ OutStreamer->emitRawText(O.str());
const auto *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
for (const auto &Entry : MFI->getCallPrototypes()) {
const auto &[CB, Symbol] = Entry.second;
- emitCallPrototype(*CB, Symbol, O);
+ emitCallPrototype(*CB, Symbol);
}
- OutStreamer->emitRawText(O.str());
-
if (const MachineJumpTableInfo *MJTI = MF->getJumpTableInfo())
for (const auto &[Idx, JT] : enumerate(MJTI->getJumpTables()))
emitJumpTable(JT, Idx);
diff --git a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
index 672444111c95d..76dbc02e9be11 100644
--- a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
+++ b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
@@ -17,7 +17,6 @@
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/CodeGen/MachineFunction.h"
#include "llvm/MC/MCContext.h"
-#include <map>
namespace llvm {
class CallBase;
@@ -57,8 +56,8 @@ class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
void addCallPrototype(unsigned Id, const CallBase *CB, MachineFunction &MF) {
if (CallPrototypes.find(Id) == CallPrototypes.end()) {
MCSymbol *Symbol =
- MF.getContext().createTempSymbol("prototype_" + Twine(Id),
- /*AlwaysAddSuffix=*/false);
+ MF.getContext().createTempSymbol("prototype_",
+ /*AlwaysAddSuffix=*/true);
CallPrototypes.try_emplace(Id, CB, Symbol);
}
}
diff --git a/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll b/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
index 3cd64105b65c3..5bd342375a381 100644
--- a/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
@@ -13,7 +13,8 @@ target triple = "nvptx64-nvidia-cuda"
%"class.sycl::_V1::detail::half_impl::half" = type { half }
%complex_half = type { half, half }
-; CHECK: $L__prototype_0 : .callprototype (.param .align 2 .b8 _[4]) _ (.param .b32 _, .param .b32 _, .param .align 2 .b8 _[4]);
+; CHECK: $L__prototype_0:
+; CHECK-NEXT: .callprototype (.param .align 2 .b8 _[4]) _ (.param .b32 _, .param .b32 _, .param .align 2 .b8 _[4]);
; CHECK: .param .align 2 .b8 param2[4];
; CHECK: .param .align 2 .b8 retval0[4];
; CHECK-DAG: st.param.b16 [param2], %rs{{[0-9]+}};
diff --git a/llvm/test/CodeGen/NVPTX/callchain.ll b/llvm/test/CodeGen/NVPTX/callchain.ll
index be61a9fbe3291..9e93b5c60883d 100644
--- a/llvm/test/CodeGen/NVPTX/callchain.ll
+++ b/llvm/test/CodeGen/NVPTX/callchain.ll
@@ -4,7 +4,8 @@
target triple = "nvptx"
define void @foo(ptr %ptr) {
-; CHECK: $L__prototype_0 : .callprototype ()_ ()
+; CHECK: $L__prototype_0:
+; CHECK-NEXT: ()_ ()
tail call void %ptr()
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/callprototype-dollar-local-label.ll b/llvm/test/CodeGen/NVPTX/callprototype-dollar-local-label.ll
index f0effa497d3ef..16aa0e1326af6 100644
--- a/llvm/test/CodeGen/NVPTX/callprototype-dollar-local-label.ll
+++ b/llvm/test/CodeGen/NVPTX/callprototype-dollar-local-label.ll
@@ -10,9 +10,10 @@ define i32 @call_via_dollar_prototype_0(i32 %a, i32 %b, i32 %c, i32 %d) {
; CHECK-DAG: .visible .global .align 8 .u64 $L__prototype_0;
; CHECK-DAG: .visible .global .align 4 .u32 $L__prototype_00 = 7;
; CHECK-LABEL: call_via_dollar_prototype_0(
-; CHECK: $L__prototype_01 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _, .param .b32 _, .param .b32 _);
+; CHECK: $L__prototype_1:
+; CHECK-NEXT: .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _, .param .b32 _, .param .b32 _);
; CHECK: ld.global.{{u|b}}64 {{%rd[0-9]+}}, [$L__prototype_0];
-; CHECK: call (retval0), %rd{{[0-9]+}}, (param0, param1, param2, param3), $L__prototype_01;
+; CHECK: call (retval0), %rd{{[0-9]+}}, (param0, param1, param2, param3), $L__prototype_1;
%fp = load ptr, ptr addrspace(1) @"$L__prototype_0", align 8
%ret = call i32 %fp(i32 %a, i32 %b, i32 %c, i32 %d)
ret i32 %ret
diff --git a/llvm/test/CodeGen/NVPTX/callprototype-local-label.ll b/llvm/test/CodeGen/NVPTX/callprototype-local-label.ll
index 894a5dc97d11b..8bc6fb509035b 100644
--- a/llvm/test/CodeGen/NVPTX/callprototype-local-label.ll
+++ b/llvm/test/CodeGen/NVPTX/callprototype-local-label.ll
@@ -7,10 +7,11 @@ target triple = "nvptx64-nvidia-cuda"
define i32 @call_via_prototype_0(i32 %a, i32 %b, i32 %c, i32 %d) {
; CHECK-LABEL: call_via_prototype_0(
-; CHECK: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _, .param .b32 _, .param .b32 _);
+; CHECK: $L__prototype_0:
+; CHECK-NEXT: .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _, .param .b32 _, .param .b32 _);
; CHECK: ld.global.{{u|b}}64 {{%rd[0-9]+}}, [prototype_0];
; CHECK: call (retval0), %rd{{[0-9]+}}, (param0, param1, param2, param3), $L__prototype_0;
-; CHECK-NOT: prototype_0 : .callprototype
+; CHECK-NOT: prototype_0:
%fp = load ptr, ptr addrspace(1) @prototype_0, align 8
%ret = call i32 %fp(i32 %a, i32 %b, i32 %c, i32 %d)
ret i32 %ret
diff --git a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
index e04c7fdef81df..da55dd1929c05 100644
--- a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
@@ -10,7 +10,8 @@ define %struct.64 @test_return_type_mismatch(ptr %p) {
; CHECK-LABEL: test_return_type_mismatch(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<32>;
-; CHECK-NEXT: $L__prototype_0 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _);
+; CHECK: $L__prototype_0:
+; CHECK-NEXT: .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_return_type_mismatch_param_0];
; CHECK-NEXT: { // callseq 0, 0
@@ -66,7 +67,8 @@ define i64 @test_param_type_mismatch(ptr %p) {
; CHECK-LABEL: test_param_type_mismatch(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-NEXT: $L__prototype_1 : .callprototype (.param .b64 _) _ (.param .b64 _);
+; CHECK: $L__prototype_1:
+; CHECK-NEXT: .callprototype (.param .b64 _) _ (.param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: { // callseq 1, 0
; CHECK-NEXT: .param .b64 param0;
@@ -86,7 +88,8 @@ define i64 @test_param_count_mismatch(ptr %p) {
; CHECK-LABEL: test_param_count_mismatch(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<4>;
-; CHECK-NEXT: $L__prototype_2 : .callprototype (.param .b64 _) _ (.param .b64 _, .param .b64 _);
+; CHECK: $L__prototype_2:
+; CHECK-NEXT: .callprototype (.param .b64 _) _ (.param .b64 _, .param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_param_count_mismatch_param_0];
; CHECK-NEXT: { // callseq 2, 0
@@ -112,7 +115,8 @@ define %struct.64 @test_return_type_mismatch_variadic(ptr %p) {
; CHECK-NEXT: .reg .b64 %SP;
; CHECK-NEXT: .reg .b64 %SPL;
; CHECK-NEXT: .reg .b64 %rd<33>;
-; CHECK-NEXT: $L__prototype_3 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _, .param .b64 _);
+; CHECK: $L__prototype_3:
+; CHECK-NEXT: .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _, .param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: mov.b64 %SPL, __local_depot3;
; CHECK-NEXT: ld.param.b64 %rd1, [test_return_type_mismatch_variadic_param_0];
diff --git a/llvm/test/CodeGen/NVPTX/empty-type.ll b/llvm/test/CodeGen/NVPTX/empty-type.ll
index 04d546e246f47..55c82c3f1b43c 100644
--- a/llvm/test/CodeGen/NVPTX/empty-type.ll
+++ b/llvm/test/CodeGen/NVPTX/empty-type.ll
@@ -132,13 +132,14 @@ define void @indirect_caller(ptr %fp) {
; CHECK-LABEL: indirect_caller(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: $L__prototype_4 : .callprototype ()_ (.param .b32 _);
+; CHECK: $L__prototype_0:
+; CHECK-NEXT: .callprototype ()_ (.param .b32 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [indirect_caller_param_0];
; CHECK-NEXT: { // callseq 4, 0
; CHECK-NEXT: .param .b32 param0;
; CHECK-NEXT: st.param.b32 [param0], 42;
-; CHECK-NEXT: call %rd1, (param0), $L__prototype_4;
+; CHECK-NEXT: call %rd1, (param0), $L__prototype_0;
; CHECK-NEXT: } // callseq 4
; CHECK-NEXT: ret;
call void %fp(%empty zeroinitializer, i32 42)
diff --git a/llvm/test/CodeGen/NVPTX/indirect_byval.ll b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
index a5a394386dc76..4c14ab13fc281 100644
--- a/llvm/test/CodeGen/NVPTX/indirect_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
@@ -18,7 +18,8 @@ define internal i32 @foo() {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-NEXT: .reg .b64 %rd<4>;
-; CHECK-NEXT: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .align 1 .b8 _[1], .param .b64 _);
+; CHECK: $L__prototype_0:
+; CHECK-NEXT: .callprototype (.param .b32 _) _ (.param .align 1 .b8 _[1], .param .b64 _);
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: mov.b64 %SPL, __local_depot0;
; CHECK-NEXT: add.u64 %rd1, %SPL, 0;
@@ -53,7 +54,8 @@ define internal i32 @bar() {
; CHECK-NEXT: .reg .b64 %SPL;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-NEXT: .reg .b64 %rd<5>;
-; CHECK-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .align 8 .b8 _[8], .param .b64 _);
+; CHECK: $L__prototype_1:
+; CHECK-NEXT: .callprototype (.param .b32 _) _ (.param .align 8 .b8 _[8], .param .b64 _);
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: mov.b64 %SPL, __local_depot1;
; CHECK-NEXT: add.u64 %rd1, %SPL, 0;
diff --git a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
index 90e1e29410027..7d86249ddc235 100644
--- a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
@@ -154,7 +154,8 @@ define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.gri
; PTX-DEFAULT-LABEL: grid_const_escape(
; PTX-DEFAULT: {
; PTX-DEFAULT-NEXT: .reg .b64 %rd<4>;
-; PTX-DEFAULT-NEXT: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-DEFAULT: $L__prototype_0:
+; PTX-DEFAULT-NEXT: .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-DEFAULT-NEXT: // %bb.0:
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_escape_param_0;
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd2, %rd1;
@@ -171,7 +172,8 @@ define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.gri
; PTX-SHORT-PTR: {
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<2>;
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<4>;
-; PTX-SHORT-PTR-NEXT: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-SHORT-PTR: $L__prototype_0:
+; PTX-SHORT-PTR-NEXT: .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-SHORT-PTR-NEXT: // %bb.0:
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_escape_param_0;
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd1, %r1;
@@ -201,7 +203,8 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-DEFAULT-NEXT: .reg .b64 %SPL;
; PTX-DEFAULT-NEXT: .reg .b32 %r<2>;
; PTX-DEFAULT-NEXT: .reg .b64 %rd<8>;
-; PTX-DEFAULT-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
+; PTX-DEFAULT: $L__prototype_1:
+; PTX-DEFAULT-NEXT: .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
; PTX-DEFAULT-NEXT: // %bb.0:
; PTX-DEFAULT-NEXT: mov.b64 %SPL, __local_depot4;
; PTX-DEFAULT-NEXT: mov.b64 %rd1, multiple_grid_const_escape_param_0;
@@ -232,7 +235,8 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-SHORT-PTR-NEXT: .reg .b32 %SPL;
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<5>;
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<8>;
-; PTX-SHORT-PTR-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
+; PTX-SHORT-PTR: $L__prototype_1:
+; PTX-SHORT-PTR-NEXT: .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
; PTX-SHORT-PTR-NEXT: // %bb.0:
; PTX-SHORT-PTR-NEXT: mov.b32 %SPL, __local_depot4;
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, multiple_grid_const_escape_param_0;
@@ -362,7 +366,8 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g
; PTX-DEFAULT: {
; PTX-DEFAULT-NEXT: .reg .b32 %r<3>;
; PTX-DEFAULT-NEXT: .reg .b64 %rd<6>;
-; PTX-DEFAULT-NEXT: $L__prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-DEFAULT: $L__prototype_2:
+; PTX-DEFAULT-NEXT: .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-DEFAULT-NEXT: // %bb.0:
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_partial_escape_param_0;
; PTX-DEFAULT-NEXT: ld.param.b64 %rd2, [grid_const_partial_escape_param_1];
@@ -384,7 +389,8 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g
; PTX-SHORT-PTR: {
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<4>;
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<6>;
-; PTX-SHORT-PTR-NEXT: $L__prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-SHORT-PTR: $L__prototype_2:
+; PTX-SHORT-PTR-NEXT: .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-SHORT-PTR-NEXT: // %bb.0:
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_partial_escape_param_0;
; PTX-SHORT-PTR-NEXT: ld.param.b64 %rd1, [grid_const_partial_escape_param_1];
@@ -422,7 +428,8 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4
; PTX-DEFAULT: {
; PTX-DEFAULT-NEXT: .reg .b32 %r<4>;
; PTX-DEFAULT-NEXT: .reg .b64 %rd<6>;
-; PTX-DEFAULT-NEXT: $L__prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-DEFAULT: $L__prototype_3:
+; PTX-DEFAULT-NEXT: .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-DEFAULT-NEXT: // %bb.0:
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_partial_escapemem_param_0;
; PTX-DEFAULT-NEXT: ld.param.b64 %rd2, [grid_const_partial_escapemem_param_1];
@@ -446,7 +453,8 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4
; PTX-SHORT-PTR: {
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<5>;
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<6>;
-; PTX-SHORT-PTR-NEXT: $L__prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-SHORT-PTR: $L__prototype_3:
+; PTX-SHORT-PTR-NEXT: .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-SHORT-PTR-NEXT: // %bb.0:
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_partial_escapemem_param_0;
; PTX-SHORT-PTR-NEXT: ld.param.b64 %rd1, [grid_const_partial_escapemem_param_1];
diff --git a/llvm/test/CodeGen/NVPTX/noreturn.ll b/llvm/test/CodeGen/NVPTX/noreturn.ll
index 416418c7ffde7..6f3d9c3e70269 100644
--- a/llvm/test/CodeGen/NVPTX/noreturn.ll
+++ b/llvm/test/CodeGen/NVPTX/noreturn.ll
@@ -32,8 +32,10 @@ define ptx_kernel void @ignore_kernel_noreturn() #0 {
}
; CHECK-LABEL: .entry callprototype_noreturn(
-; CHECK: {{[$]}}L__prototype_{{[0-9]+}} : .callprototype ()_ (.param .b32 _) .noreturn;
-; CHECK: {{[$]}}L__prototype_{{[0-9]+}} : .callprototype (.param .b32 _) _ (.param .b32 _);
+; CHECK: {{[$]}}L__prototype_{{[0-9]+}}:
+; CHECK-NEXT: ()_ (.param .b32 _) .noreturn;
+; CHECK: {{[$]}}L__prototype_{{[0-9]+}}:
+; CHECK-NEXT: .callprototype (.param .b32 _) _ (.param .b32 _);
define ptx_kernel void @callprototype_noreturn(i32) {
%fn = load ptr, ptr addrspace(1) @function_pointer
diff --git a/llvm/test/CodeGen/NVPTX/param-align.ll b/llvm/test/CodeGen/NVPTX/param-align.ll
index 7f0a1646962ae..ce8f47d8048e2 100644
--- a/llvm/test/CodeGen/NVPTX/param-align.ll
+++ b/llvm/test/CodeGen/NVPTX/param-align.ll
@@ -55,17 +55,22 @@ declare ptr @getfp(i32 %n)
define ptx_device void @t6() {
; CHECK: .func t6
%fp = call ptr @getfp(i32 0)
-; CHECK: $L__prototype_2 : .callprototype ()_ (.param .align 8 .b8 _[8]);
+; CHECK: $L__prototype_0:
+; CHECK-NEXT: ()_ (.param .align 8 .b8 _[8]);
call void %fp(ptr byval(double) null);
%fp2 = call ptr @getfp(i32 1)
-; NOALIGN4: $L__prototype_4 : .callprototype ()_ (.param .align 2 .b8 _[4]);
-; ALIGN4: $L__prototype_4 : .callprototype ()_ (.param .align 4 .b8 _[4]);
+; NOALIGN4: $L__prototype_1:
+; NOALIGN4-NEXT: ()_ (.param .align 2 .b8 _[4]);
+; ALIGN4: $L__prototype_1:
+; ALIGN4-NEXT: ()_ (.param .align 4 .b8 _[4]);
call void %fp(ptr byval(%struct.half2) null);
%fp3 = call ptr @getfp(i32 2)
-; NOALIGN4: $L__prototype_6 : .callprototype ()_ (.param .align 1 .b8 _[1]);
-; ALIGN4: $L__prototype_6 : .callprototype ()_ (.param .align 4 .b8 _[1]);
+; NOALIGN4: $L__prototype_2:
+; NOALIGN4-NEXT: ()_ (.param .align 1 .b8 _[1]);
+; ALIGN4: $L__prototype_2:
+; ALIGN4-NEXT: ()_ (.param .align 4 .b8 _[1]);
call void %fp(ptr byval(i8) null);
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/vaargs.ll b/llvm/test/CodeGen/NVPTX/vaargs.ll
index 34ff8c18e2ca0..3c40bd47c7984 100644
--- a/llvm/test/CodeGen/NVPTX/vaargs.ll
+++ b/llvm/test/CodeGen/NVPTX/vaargs.ll
@@ -216,7 +216,8 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK32-NEXT: .reg .b32 %SPL;
; CHECK32-NEXT: .reg .b32 %r<9>;
; CHECK32-NEXT: .reg .b64 %rd<3>;
-; CHECK32-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _);
+; CHECK32: $L__prototype_0:
+; CHECK32-NEXT: .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _);
; CHECK32-NEXT: // %bb.0: // %entry
; CHECK32-NEXT: mov.b32 %SPL, __local_depot1;
; CHECK32-NEXT: ld.param.b32 %r2, [test_foo_param_3];
@@ -238,7 +239,7 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK32-NEXT: .param .b32 retval0;
; CHECK32-NEXT: st.param.b32 [param1], %r3;
; CHECK32-NEXT: st.param.b32 [param0], 4;
-; CHECK32-NEXT: call (retval0), %r7, (param0, param1), $L__prototype_1;
+; CHECK32-NEXT: call (retval0), %r7, (param0, param1), $L__prototype_0;
; CHECK32-NEXT: ld.param.b32 %r8, [retval0];
; CHECK32-NEXT: } // callseq 1
; CHECK32-NEXT: st.param.b32 [func_retval0], %r8;
@@ -251,7 +252,8 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK64-NEXT: .reg .b64 %SPL;
; CHECK64-NEXT: .reg .b32 %r<3>;
; CHECK64-NEXT: .reg .b64 %rd<9>;
-; CHECK64-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b64 _);
+; CHECK64: $L__prototype_0:
+; CHECK64-NEXT: .callprototype (.param .b32 _) _ (.param .b32 _, .param .b64 _);
; CHECK64-NEXT: // %bb.0: // %entry
; CHECK64-NEXT: mov.b64 %SPL, __local_depot1;
; CHECK64-NEXT: ld.param.b64 %rd3, [test_foo_param_3];
@@ -273,7 +275,7 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK64-NEXT: .param .b32 retval0;
; CHECK64-NEXT: st.param.b64 [param1], %rd4;
; CHECK64-NEXT: st.param.b32 [param0], 4;
-; CHECK64-NEXT: call (retval0), %rd8, (param0, param1), $L__prototype_1;
+; CHECK64-NEXT: call (retval0), %rd8, (param0, param1), $L__prototype_0;
; CHECK64-NEXT: ld.param.b32 %r2, [retval0];
; CHECK64-NEXT: } // callseq 1
; CHECK64-NEXT: st.param.b32 [func_retval0], %r2;
More information about the cfe-commits
mailing list