Author: Ambrose Leeb Date: 2026-10-14T17:53:10Z New Revision: 60994af739ad7cec81f846c358520f5880fceb7c
URL: https://github.com/llvm/llvm-project/commit/60994af739ad7cec81f846c358520f5880fceb7c DIFF: https://github.com/llvm/llvm-project/commit/60994af739ad7cec81f846c358520f5880fceb7c.diff LOG: [Clang] [CUDA] Fix calling convention for forward-declared kernels (#231628) 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 Added: clang/test/CodeGenCUDA/kernel-declarations.cu Modified: clang/docs/ReleaseNotes.md clang/lib/CodeGen/Targets/NVPTX.cpp Removed: ################################################################################ diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index aba1dba5854fea..bdb35c448d0d99 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -1168,6 +1168,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
