[clang] [llvm] [NVPTX] Fix for device function pointer having the same name as generated callprototype (PR #205639)
Daniel Donenfeld via llvm-commits
llvm-commits at lists.llvm.org
Mon Jun 29 12:55:01 PDT 2026
https://github.com/daniel-donenfeld updated https://github.com/llvm/llvm-project/pull/205639
>From 6d246b8f91bea2fc384a29094653af39afac9f9a 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 1/3] Bug fix for device function pointer having same name as a
generated callprototype
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.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/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 7854309c42e5d..425f480d83615 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1205,7 +1205,7 @@ std::string NVPTXTargetLowering::getPrototype(
std::string Prototype;
raw_string_ostream O(Prototype);
- 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 ff3e5144c6fbe..ac44c2fcc5604 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 483d48a1012c6..576af8d5db0cb 100644
--- a/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
@@ -17,7 +17,7 @@ target triple = "nvptx64-nvidia-cuda"
; CHECK: .param .align 2 .b8 retval0[4];
; CHECK-DAG: st.param.b16 [param2], %rs{{[0-9]+}};
; CHECK-DAG: st.param.b16 [param2+2], %rs{{[0-9]+}};
-; 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: call (retval0),
define weak_odr void @foo() {
entry:
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 d82a42ae5d56e..f881ac5f05d44 100644
--- a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
@@ -17,9 +17,9 @@ define %struct.64 @test_return_type_mismatch(ptr %p) {
; CHECK-NEXT: .param .b64 param0;
; CHECK-NEXT: .param .align 1 .b8 retval0[8];
; CHECK-NEXT: st.param.b64 [param0], %rd1;
-; 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: 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];
@@ -72,10 +72,10 @@ define i64 @test_param_type_mismatch(ptr %p) {
; CHECK-NEXT: { // callseq 1, 0
; CHECK-NEXT: .param .b64 param0;
; CHECK-NEXT: .param .b64 retval0;
-; CHECK-NEXT: prototype_1 : .callprototype (.param .b64 _) _ (.param .b64 _);
+; CHECK-NEXT: $L__prototype_1 : .callprototype (.param .b64 _) _ (.param .b64 _);
; 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;
@@ -96,10 +96,10 @@ define i64 @test_param_count_mismatch(ptr %p) {
; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .b64 retval0;
; CHECK-NEXT: st.param.b64 [param0], %rd1;
-; 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: 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;
@@ -127,9 +127,9 @@ define %struct.64 @test_return_type_mismatch_variadic(ptr %p) {
; CHECK-NEXT: add.u64 %rd2, %SP, 0;
; CHECK-NEXT: st.param.b64 [param1], %rd2;
; CHECK-NEXT: st.param.b64 [param0], %rd1;
-; 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: 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 e1fecdb76bd4d..6bee09ae21ff8 100644
--- a/llvm/test/CodeGen/NVPTX/indirect_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
@@ -32,8 +32,8 @@ 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: prototype_0 : .callprototype (.param .b32 _) _ (.param .align 1 .b8 _[1], .param .b64 _);
-; CHECK-NEXT: call (retval0), %rd1, (param0, param1), prototype_0;
+; CHECK-NEXT: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .align 1 .b8 _[1], .param .b64 _);
+; 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;
@@ -69,8 +69,8 @@ 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: prototype_1 : .callprototype (.param .b32 _) _ (.param .align 8 .b8 _[8], .param .b64 _);
-; CHECK-NEXT: call (retval0), %rd1, (param0, param1), prototype_1;
+; CHECK-NEXT: $L__prototype_1 : .callprototype (.param .b32 _) _ (.param .align 8 .b8 _[8], .param .b64 _);
+; 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 4d9cb6ff97413..805fa1bef2296 100644
--- a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
@@ -130,9 +130,9 @@ define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.gri
; PTX-NEXT: .param .b64 param0;
; PTX-NEXT: .param .b32 retval0;
; PTX-NEXT: st.param.b64 [param0], %rd2;
-; PTX-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-NEXT: $L__prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
; 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(
@@ -173,9 +173,9 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-NEXT: st.param.b64 [param2], %rd3;
; PTX-NEXT: st.param.b64 [param1], %rd5;
; PTX-NEXT: st.param.b64 [param0], %rd4;
-; 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: 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(
@@ -267,9 +267,9 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g
; PTX-NEXT: .param .b64 param0;
; PTX-NEXT: .param .b32 retval0;
; PTX-NEXT: st.param.b64 [param0], %rd4;
-; PTX-NEXT: prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-NEXT: $L__prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
; 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(
@@ -307,9 +307,9 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4
; PTX-NEXT: .param .b64 param0;
; PTX-NEXT: .param .b32 retval0;
; PTX-NEXT: st.param.b64 [param0], %rd4;
-; PTX-NEXT: prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
+; PTX-NEXT: $L__prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
; 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 b8d1f5c819bd3a7958f4674fa531f096957d5b1c 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 2/3] 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 a03cf90f61a34..6769086ce93d0 100644
--- a/llvm/test/CodeGen/NVPTX/vaargs.ll
+++ b/llvm/test/CodeGen/NVPTX/vaargs.ll
@@ -207,9 +207,9 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK32-NEXT: .param .b32 retval0;
; CHECK32-NEXT: add.u32 %r6, %SP, 0;
; CHECK32-NEXT: st.param.b32 [param1], %r6;
-; 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: 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 %r7, [retval0];
; CHECK32-NEXT: } // callseq 1
; CHECK32-NEXT: st.param.b32 [func_retval0], %r7;
@@ -243,9 +243,9 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK64-NEXT: .param .b32 retval0;
; CHECK64-NEXT: add.u64 %rd7, %SP, 0;
; CHECK64-NEXT: st.param.b64 [param1], %rd7;
-; 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: 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 f9d28d3eaa396bc7b4a909c8c737f02c43df8af6 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 3/3] 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
More information about the llvm-commits
mailing list