https://github.com/Sirraide updated https://github.com/llvm/llvm-project/pull/231628
>From f8168305e175671d2d21cfe2ab8341c312f47ecb Mon Sep 17 00:00:00 2001 From: Ambrose Leeb <[email protected]> Date: Tue, 13 Oct 2026 19:20:33 +0000 Subject: [PATCH] [Clang] [CUDA] Fix calling convention for forward-declared kernels The calling convention was not being set properly, leading to linker errors in nvlink. We already set the calling convention this way for AMDGPU, so do the same for NVPTX. Fixes #118212. Assisted-by: Codex --- clang/docs/ReleaseNotes.md | 2 + clang/lib/CodeGen/Targets/NVPTX.cpp | 7 +++- clang/test/CodeGenCUDA/kernel-declarations.cu | 41 +++++++++++++++++++ 3 files changed, 48 insertions(+), 2 deletions(-) create mode 100644 clang/test/CodeGenCUDA/kernel-declarations.cu diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index 1dc1764978ea27..52437f15f8a7b3 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -1119,6 +1119,8 @@ features cannot lower the translation-unit ABI level; virtual function that could not be optimised out would cause the program to fail to assemble. This is now fixed. (#GH49183) (#GH67533) +- External CUDA kernel declarations now use the correct calling convention. (#GH118212) + #### AIX Support #### NetBSD Support diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp index 0229389fe17161..ed70b01fcf48bb 100644 --- a/clang/lib/CodeGen/Targets/NVPTX.cpp +++ b/clang/lib/CodeGen/Targets/NVPTX.cpp @@ -91,6 +91,11 @@ class NVPTXTargetCodeGenInfo : public TargetCodeGenInfo { return llvm::CallingConv::PTX_Kernel; } + void setCUDAKernelCallingConvention(const FunctionType *&FT) const override { + FT = getABIInfo().getContext().adjustFunctionType( + FT, FT->getExtInfo().withCallingConv(CC_DeviceKernel)); + } + // Adds a NamedMDNode with GV, Name, and Operand as operands, and adds the // resulting MDNode to the nvvm.annotations MDNode. static void addNVVMMetadata(llvm::GlobalValue *GV, StringRef Name, @@ -273,8 +278,6 @@ void NVPTXTargetCodeGenInfo::setTargetAttributes( // And kernel functions are not subject to inlining F->addFnAttr(llvm::Attribute::NoInline); if (FD->hasAttr<CUDAGlobalAttr>()) { - F->setCallingConv(getDeviceKernelCallingConv()); - for (auto IV : llvm::enumerate(FD->parameters())) if (IV.value()->hasAttr<CUDAGridConstantAttr>()) F->addParamAttr(IV.index(), diff --git a/clang/test/CodeGenCUDA/kernel-declarations.cu b/clang/test/CodeGenCUDA/kernel-declarations.cu new file mode 100644 index 00000000000000..9828afbf35b7bb --- /dev/null +++ b/clang/test/CodeGenCUDA/kernel-declarations.cu @@ -0,0 +1,41 @@ +// REQUIRES: nvptx-registered-target +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fgpu-rdc \ +// RUN: -emit-llvm -o - %s | FileCheck %s --check-prefix=IR +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_89 \ +// RUN: -fcuda-is-device -fgpu-rdc -S -o - %s | FileCheck %s --check-prefix=PTX + +#include "Inputs/cuda.h" + +// Kernels declared in another translation unit must have the kernel calling +// convention even when only referenced by host-side launches. +extern "C" __global__ void kernel(int); +struct Arg { + int x; +}; +template <class T> __global__ void templated_kernel(T); +extern template __global__ void templated_kernel<Arg>(Arg); +struct Holder { + static __global__ void member(int); +}; + +// IR-DAG: declare ptx_kernel void @kernel(i32 noundef) +// IR-DAG: declare ptx_kernel void @_Z16templated_kernelI3ArgEvT_(ptr noundef byval(%struct.Arg) align 4) +// IR-DAG: declare ptx_kernel void @_ZN6Holder6memberEi(i32 noundef) +// PTX-DAG: .extern .entry kernel +// PTX-DAG: .extern .entry _Z16templated_kernelI3ArgEvT_ +// PTX-DAG: .extern .entry _ZN6Holder6memberEi + +// Definitions continue to distinguish kernels from ordinary device functions. +// IR-DAG: define{{.*}} ptx_kernel void @defined_kernel() +// IR-DAG: define{{.*}} void @device_function() +// PTX-DAG: .visible .entry defined_kernel +// PTX-DAG: .visible .func device_function +extern "C" __global__ void defined_kernel() {} +extern "C" __device__ void device_function() {} + +void launch() { + kernel<<<1, 1>>>(1); + templated_kernel<Arg><<<1, 1>>>(Arg{1}); + Holder::member<<<1, 1>>>(1); + defined_kernel<<<1, 1>>>(); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
