[clang] [CIR] Register only CUDA/HIP kernels with emitted host stubs (PR #234911)
Martin Paul Lücke via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 23 04:06:50 PDT 2026
https://github.com/martin-luecke created https://github.com/llvm/llvm-project/pull/234911
CIR currently attaches `cu.kernel_name` to external kernel declarations and references. The device image and runtime registration passes consequently try to register those kernels against translation units that do not define them, causing runtime lookup failures.
Move `cu.kernel_name` to the host stub definition. Lowering then generates registration calls from the same set of kernels as conventional CodeGen. This also avoids loading the device image when the translation unit only references external kernels.
Add CUDA and HIP tests that reference external kernels without a device image and verify that only kernels defined in the current translation unit are registered.
>From d637cf32c50da4ce24f4bf81b1062280bec87fec Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Martin=20Paul=20L=C3=BCcke?= <martin.luecke at amd.com>
Date: Tue, 20 Oct 2026 10:36:24 +0000
Subject: [PATCH] [CIR] Register only CUDA/HIP kernels with emitted host stubs
CIR attaches `cu.kernel_name` to host-side kernel calls. Calls to kernels
defined outside the current translation unit therefore cause lowering to emit
registration calls for kernels that are not present in the device image.
Attach the attribute to host stub definitions instead. This matches
conventional CodeGen, which records kernels for registration in
`emitDeviceStub`, and avoids reading the device image when there is nothing to
register.
Add CUDA and HIP coverage for external kernel references and check that
registration is emitted only for kernels with emitted host stubs.
Assisted-by: GPT 5.6 Sol
---
clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp | 8 ++-
clang/lib/CIR/CodeGen/CIRGenCall.cpp | 11 ----
clang/test/CIR/CodeGenCUDA/kernel-call.cu | 4 +-
.../CIR/CodeGenCUDA/kernel-registration.cu | 50 +++++++++++++++++++
4 files changed, 59 insertions(+), 14 deletions(-)
create mode 100644 clang/test/CIR/CodeGenCUDA/kernel-registration.cu
diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
index dfd311de78d542..3f1b1e6cdd868f 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
@@ -346,10 +346,16 @@ void CIRGenNVCUDARuntime::emitDeviceStubBodyNew(CIRGenFunction &cgf,
void CIRGenNVCUDARuntime::emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn,
FunctionArgList &args) {
+ CIRGenBuilderTy &builder = cgm.getBuilder();
+ StringRef kernelName = cgm.getMangledName(
+ cgf.curGD.getWithKernelReferenceKind(KernelReferenceKind::Kernel));
+ auto kernelNameAttr = cir::CUDAKernelNameAttr::get(
+ builder.getContext(), builder.getStringAttr(kernelName));
+ fn->setDiscardableAttr(cir::CUDAKernelNameAttr::getMnemonic(),
+ kernelNameAttr);
if (auto globalOp =
llvm::dyn_cast<cir::GlobalOp>(kernelHandles[fn.getSymName()])) {
- CIRGenBuilderTy &builder = cgm.getBuilder();
mlir::Type fnPtrTy = globalOp.getSymType();
auto sym = mlir::FlatSymbolRefAttr::get(fn.getSymNameAttr());
auto gv = cir::GlobalViewAttr::get(fnPtrTy, sym);
diff --git a/clang/lib/CIR/CodeGen/CIRGenCall.cpp b/clang/lib/CIR/CodeGen/CIRGenCall.cpp
index 2d9113c2789933..8e5021c1b2fb69 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCall.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCall.cpp
@@ -484,17 +484,6 @@ void CIRGenModule::constructAttributeList(
if (langOpts.OffloadUniformBlock)
addUnitAttr(cir::CIRDialect::getUniformWorkGroupSizeAttrName());
- if (langOpts.CUDA && !langOpts.CUDAIsDevice &&
- targetDecl->hasAttr<CUDAGlobalAttr>()) {
- GlobalDecl kernel(calleeInfo.getCalleeDecl());
- llvm::StringRef kernelName = getMangledName(
- kernel.getWithKernelReferenceKind(KernelReferenceKind::Kernel));
- auto attr = cir::CUDAKernelNameAttr::get(
- &getMLIRContext(),
- mlir::StringAttr::get(&getMLIRContext(), kernelName));
- attrs.set(attr.getMnemonic(), attr);
- }
-
// TODO(cir): we should also do 'aarch64_pstate_sm_body' here.
if (auto *modularFormat = targetDecl->getAttr<ModularFormatAttr>()) {
diff --git a/clang/test/CIR/CodeGenCUDA/kernel-call.cu b/clang/test/CIR/CodeGenCUDA/kernel-call.cu
index 58da391503faf5..392d649341e806 100644
--- a/clang/test/CIR/CodeGenCUDA/kernel-call.cu
+++ b/clang/test/CIR/CodeGenCUDA/kernel-call.cu
@@ -118,13 +118,13 @@ int main(void) {
// CUDA-NEW: } else {
// CUDA-NEW: cir.const #cir.int<42> : !s32i
// CUDA-NEW: cir.const #cir.fp<1.000000e+00> : !cir.float
- // CUDA-NEW: cir.call @_Z21__device_stub__kernelif({{.*}}) {cu.kernel_name = #cir.cu.kernel_name<"_Z6kernelif">, uniform_work_group_size} : (!s32i {llvm.noundef}, !cir.float {llvm.noundef}) -> ()
+ // CUDA-NEW: cir.call @_Z21__device_stub__kernelif({{.*}}) {uniform_work_group_size} : (!s32i {llvm.noundef}, !cir.float {llvm.noundef}) -> ()
// CUDA-NEW: }
// HIP-NEW: cir.if %{{.*}} {
// HIP-NEW: } else {
// HIP-NEW: cir.const #cir.int<42> : !s32i
// HIP-NEW: cir.const #cir.fp<1.000000e+00> : !cir.float
- // HIP-NEW: cir.call @_Z21__device_stub__kernelif({{.*}}) {cu.kernel_name = #cir.cu.kernel_name<"_Z6kernelif">, uniform_work_group_size} : (!s32i {llvm.noundef}, !cir.float {llvm.noundef}) -> ()
+ // HIP-NEW: cir.call @_Z21__device_stub__kernelif({{.*}}) {uniform_work_group_size} : (!s32i {llvm.noundef}, !cir.float {llvm.noundef}) -> ()
// HIP-NEW: }
kernel<<<1, 1>>>(42, 1.0f);
}
diff --git a/clang/test/CIR/CodeGenCUDA/kernel-registration.cu b/clang/test/CIR/CodeGenCUDA/kernel-registration.cu
new file mode 100644
index 00000000000000..adba0f36015b95
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/kernel-registration.cu
@@ -0,0 +1,50 @@
+// External references must not read the image or carry kernel metadata.
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -x hip -fhip-new-launch-api \
+// RUN: -fcuda-include-gpubinary %t.nonexistent -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefix=DECL --implicit-check-not=cu.kernel_name
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -x cuda -target-sdk-version=12.3 \
+// RUN: -fcuda-include-gpubinary %t.nonexistent -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefix=DECL --implicit-check-not=cu.kernel_name
+
+// RUN: echo -n "GPU binary would be here." > %t.fatbin
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -x hip -fhip-new-launch-api \
+// RUN: -fcuda-include-gpubinary %t.fatbin -DEMIT_DEFINITIONS -fclangir -emit-llvm %s -o %t.ll
+// RUN: FileCheck %s --check-prefix=REGISTER --input-file=%t.ll
+// RUN: FileCheck %s --check-prefix=EXACT --input-file=%t.ll \
+// RUN: --implicit-check-not='call{{.+}}RegisterFunction'
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -x cuda -target-sdk-version=12.3 \
+// RUN: -fcuda-include-gpubinary %t.fatbin -DEMIT_DEFINITIONS -fclangir -emit-llvm %s -o %t.ll
+// RUN: FileCheck %s --check-prefix=REGISTER --input-file=%t.ll
+// RUN: FileCheck %s --check-prefix=EXACT --input-file=%t.ll \
+// RUN: --implicit-check-not='call{{.+}}RegisterFunction'
+
+// EXACT-COUNT-2: call{{.+}}@__{{(cuda|hip)}}RegisterFunction(
+// REGISTER-DAG: call{{.+}}@__{{(cuda|hip)}}RegisterFunction({{.+}}@{{(__device_stub__)?}}defined_kernel,
+// REGISTER-DAG: call{{.+}}@__{{(cuda|hip)}}RegisterFunction({{.+}}@{{.+}}template_kernelIfEvv,
+
+#include "Inputs/cuda.h"
+
+extern "C" __global__ void external_kernel();
+extern "C" __global__ void defined_kernel();
+template <class T> __global__ void template_kernel() {}
+extern template __global__ void template_kernel<int>();
+
+// DECL: cir.global external @external_address = #cir.global_view<@{{(__device_stub__)?}}external_kernel>
+const void *external_address = (const void *)external_kernel;
+
+// DECL-LABEL: cir.func{{.+}}@_Z6launchv(
+void launch() {
+ // DECL: cir.call @__device_stub__external_kernel()
+ external_kernel<<<1, 1>>>();
+ // DECL: cir.call @_Z30__device_stub__template_kernelIiEvv()
+ template_kernel<int><<<1, 1>>>();
+ // DECL: cir.call @__device_stub__defined_kernel()
+ defined_kernel<<<1, 1>>>();
+#ifdef EMIT_DEFINITIONS
+ template_kernel<float><<<1, 1>>>();
+#endif
+}
+
+#ifdef EMIT_DEFINITIONS
+extern "C" __global__ void defined_kernel() {}
+#endif
More information about the cfe-commits
mailing list