https://github.com/daniel-donenfeld updated https://github.com/llvm/llvm-project/pull/205639
>From b772f93433fc53b6959c084014c2bbcfe8d6778f Mon Sep 17 00:00:00 2001 From: Daniel Donenfeld <[email protected]> Date: Tue, 2 Jun 2026 20:15:39 +0000 Subject: [PATCH 1/4] 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 +- .../CodeGen/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 | 16 ++++++++-------- llvm/test/CodeGen/NVPTX/noreturn.ll | 4 ++-- llvm/test/CodeGen/NVPTX/param-align.ll | 10 +++++----- 10 files changed, 48 insertions(+), 31 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 94acde4b22b5e..5a474e458fc14 100644 --- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp @@ -323,7 +323,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()) { O << "()"; diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td index 65dd146926f8b..e162736cf04f3 100644 --- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td +++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td @@ -1845,7 +1845,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 1510be1b0daa9..9f763eca19915 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 c982aaf70d4a2..db70140c16ff4 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: cvta.local.u64 %SP, %SPL; @@ -32,7 +32,7 @@ define internal i32 @foo() { ; CHECK-NEXT: add.u64 %rd3, %SPL, 1; ; CHECK-NEXT: ld.local.b8 %rs1, [%rd3]; ; CHECK-NEXT: st.param.b8 [param0], %rs1; -; CHECK-NEXT: call (retval0), %rd1, (param0, param1), prototype_0; +; CHECK-NEXT: call (retval0), %rd1, (param0, param1), $L__prototype_0; ; CHECK-NEXT: ld.param.b32 %r1, [retval0]; ; CHECK-NEXT: } // callseq 0 ; CHECK-NEXT: st.param.b32 [func_retval0], %r1; @@ -54,7 +54,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: cvta.local.u64 %SP, %SPL; @@ -68,7 +68,7 @@ define internal i32 @bar() { ; CHECK-NEXT: add.u64 %rd3, %SPL, 8; ; CHECK-NEXT: ld.local.b64 %rd4, [%rd3]; ; CHECK-NEXT: st.param.b64 [param0], %rd4; -; CHECK-NEXT: call (retval0), %rd1, (param0, param1), prototype_1; +; CHECK-NEXT: call (retval0), %rd1, (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 da8c440512af3..ee4f852ca586d 100644 --- a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll +++ b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll @@ -122,7 +122,7 @@ define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.gri ; PTX-LABEL: grid_const_escape( ; PTX: { ; PTX-NEXT: .reg .b64 %rd<4>; -; PTX-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _); +; PTX-NEXT: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _); ; PTX-NEXT: // %bb.0: ; PTX-NEXT: mov.b64 %rd1, grid_const_escape_param_0; ; PTX-NEXT: cvta.param.u64 %rd2, %rd1; @@ -131,7 +131,7 @@ define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.gri ; PTX-NEXT: .param .b32 retval0; ; PTX-NEXT: st.param.b64 [param0], %rd2; ; PTX-NEXT: mov.b64 %rd3, escape; -; PTX-NEXT: call (retval0), %rd3, (param0), prototype_0; +; PTX-NEXT: call (retval0), %rd3, (param0), $L__prototype_0; ; PTX-NEXT: } // callseq 0 ; PTX-NEXT: ret; ; OPT-LABEL: define ptx_kernel void @grid_const_escape( @@ -152,7 +152,7 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4 ; PTX-NEXT: .reg .b64 %SPL; ; PTX-NEXT: .reg .b32 %r<2>; ; PTX-NEXT: .reg .b64 %rd<8>; -; PTX-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _); +; PTX-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _); ; PTX-NEXT: // %bb.0: ; PTX-NEXT: mov.b64 %SPL, __local_depot4; ; PTX-NEXT: cvta.local.u64 %SP, %SPL; @@ -173,7 +173,7 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4 ; PTX-NEXT: st.param.b64 [param1], %rd5; ; PTX-NEXT: st.param.b64 [param0], %rd4; ; PTX-NEXT: mov.b64 %rd7, escape3; -; PTX-NEXT: call (retval0), %rd7, (param0, param1, param2), prototype_1; +; PTX-NEXT: call (retval0), %rd7, (param0, param1, param2), $L__prototype_1; ; PTX-NEXT: } // callseq 1 ; PTX-NEXT: ret; ; OPT-LABEL: define ptx_kernel void @multiple_grid_const_escape( @@ -251,7 +251,7 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g ; PTX: { ; PTX-NEXT: .reg .b32 %r<3>; ; PTX-NEXT: .reg .b64 %rd<6>; -; PTX-NEXT: prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _); +; PTX-NEXT: $L__prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _); ; PTX-NEXT: // %bb.0: ; PTX-NEXT: mov.b64 %rd1, grid_const_partial_escape_param_0; ; PTX-NEXT: ld.param.b64 %rd2, [grid_const_partial_escape_param_1]; @@ -265,7 +265,7 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g ; PTX-NEXT: .param .b32 retval0; ; PTX-NEXT: st.param.b64 [param0], %rd4; ; PTX-NEXT: mov.b64 %rd5, escape; -; PTX-NEXT: call (retval0), %rd5, (param0), prototype_2; +; PTX-NEXT: call (retval0), %rd5, (param0), $L__prototype_2; ; PTX-NEXT: } // callseq 2 ; PTX-NEXT: ret; ; OPT-LABEL: define ptx_kernel void @grid_const_partial_escape( @@ -289,7 +289,7 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4 ; PTX: { ; PTX-NEXT: .reg .b32 %r<4>; ; PTX-NEXT: .reg .b64 %rd<6>; -; PTX-NEXT: prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _); +; PTX-NEXT: $L__prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _); ; PTX-NEXT: // %bb.0: ; PTX-NEXT: mov.b64 %rd1, grid_const_partial_escapemem_param_0; ; PTX-NEXT: ld.param.b64 %rd2, [grid_const_partial_escapemem_param_1]; @@ -304,7 +304,7 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4 ; PTX-NEXT: .param .b32 retval0; ; PTX-NEXT: st.param.b64 [param0], %rd4; ; PTX-NEXT: mov.b64 %rd5, escape; -; PTX-NEXT: call (retval0), %rd5, (param0), prototype_3; +; PTX-NEXT: call (retval0), %rd5, (param0), $L__prototype_3; ; PTX-NEXT: } // callseq 3 ; PTX-NEXT: st.param.b32 [func_retval0], %r3; ; PTX-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 af56fee30c3e08cfacdebd2ea3eba60c569dac99 Mon Sep 17 00:00:00 2001 From: Daniel Donenfeld <[email protected]> Date: Wed, 24 Jun 2026 21:34:42 +0000 Subject: [PATCH 2/4] 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 ea0807a88bd68..f11c41cf68034 100644 --- a/llvm/test/CodeGen/NVPTX/vaargs.ll +++ b/llvm/test/CodeGen/NVPTX/vaargs.ll @@ -188,7 +188,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: cvta.local.u32 %SP, %SPL; @@ -211,7 +211,7 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) { ; CHECK32-NEXT: .param .b32 retval0; ; CHECK32-NEXT: st.param.b32 [param1], %r7; ; CHECK32-NEXT: st.param.b32 [param0], 4; -; CHECK32-NEXT: call (retval0), %r5, (param0, param1), prototype_1; +; CHECK32-NEXT: call (retval0), %r5, (param0, param1), $L__prototype_1; ; CHECK32-NEXT: ld.param.b32 %r8, [retval0]; ; CHECK32-NEXT: } // callseq 1 ; CHECK32-NEXT: st.param.b32 [func_retval0], %r8; @@ -224,7 +224,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: cvta.local.u64 %SP, %SPL; @@ -247,7 +247,7 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) { ; CHECK64-NEXT: .param .b32 retval0; ; CHECK64-NEXT: st.param.b64 [param1], %rd8; ; CHECK64-NEXT: st.param.b32 [param0], 4; -; CHECK64-NEXT: call (retval0), %rd6, (param0, param1), prototype_1; +; CHECK64-NEXT: call (retval0), %rd6, (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 eb992a362888de4b89e6a30525eb74178433ae8e Mon Sep 17 00:00:00 2001 From: Daniel Donenfeld <[email protected]> Date: Mon, 29 Jun 2026 19:54:26 +0000 Subject: [PATCH 3/4] 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 e4a8998aba70fcbbec0f22d4648c4966f8542837 Mon Sep 17 00:00:00 2001 From: Daniel Donenfeld <[email protected]> Date: Mon, 29 Jun 2026 20:31:48 +0000 Subject: [PATCH 4/4] 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) _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
