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 <[email protected]> Date: Tue, 2 Jun 2026 20:15:39 +0000 Subject: [PATCH 01/11] 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" + +@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 <[email protected]> Date: Wed, 24 Jun 2026 21:34:42 +0000 Subject: [PATCH 02/11] 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 <[email protected]> Date: Mon, 29 Jun 2026 19:54:26 +0000 Subject: [PATCH 03/11] 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 <[email protected]> Date: Mon, 29 Jun 2026 20:31:48 +0000 Subject: [PATCH 04/11] 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 <[email protected]> Date: Tue, 18 Aug 2026 17:29:07 +0000 Subject: [PATCH 05/11] 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 <[email protected]> Date: Wed, 26 Aug 2026 21:51:16 +0000 Subject: [PATCH 06/11] 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 <[email protected]> Date: Wed, 26 Aug 2026 22:04:23 +0000 Subject: [PATCH 07/11] 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 <[email protected]> Date: Tue, 8 Sep 2026 16:43:21 +0000 Subject: [PATCH 08/11] 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 <[email protected]> Date: Tue, 8 Sep 2026 16:49:08 +0000 Subject: [PATCH 09/11] 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 <[email protected]> Date: Tue, 8 Sep 2026 16:58:38 +0000 Subject: [PATCH 10/11] 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 <[email protected]> Date: Wed, 9 Sep 2026 14:46:57 +0000 Subject: [PATCH 11/11] 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; } }; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
