Author: David Rivera Date: 2026-10-23T07:16:45-04:00 New Revision: beded2c4e47599416bf6a96e5972ff688454dfb6
URL: https://github.com/llvm/llvm-project/commit/beded2c4e47599416bf6a96e5972ff688454dfb6 DIFF: https://github.com/llvm/llvm-project/commit/beded2c4e47599416bf6a96e5972ff688454dfb6.diff LOG: [CIR][CUDA] Register only defined stubs (#231390) Added: clang/test/CIR/CodeGenCUDA/kernel-declaration.cu Modified: clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp index dfd311de78d54..680486411a479 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp @@ -538,15 +538,19 @@ void CIRGenNVCUDARuntime::handleGlobalReplace(cir::GlobalOp oldGV, } /// Whether this translation unit has anything for the CUDA runtime to register. -/// These are the same two attributes LoweringPrepare collects to decide whether -/// to build a module ctor, so both sides answer the question from one source. +/// Like LoweringPrepare, count only kernel stub definitions and device +/// variables carrying registration metadata, not kernel declarations or +/// launch-site calls. static bool hasEntitiesToRegister(mlir::ModuleOp module) { // A walk, not a scan of the module body, so this stays in agreement with the // recursive walk LoweringPrepare collects them with. return module ->walk([](mlir::Operation *op) { - if (op->hasAttr(cir::CUDAKernelNameAttr::getMnemonic()) || - op->hasAttr(cir::CUDAVarRegistrationInfoAttr::getMnemonic())) + if (auto fn = mlir::dyn_cast<cir::FuncOp>(op); + fn && !fn.isDeclaration() && + fn->hasAttr(cir::CUDAKernelNameAttr::getMnemonic())) + return mlir::WalkResult::interrupt(); + if (op->hasAttr(cir::CUDAVarRegistrationInfoAttr::getMnemonic())) return mlir::WalkResult::interrupt(); return mlir::WalkResult::advance(); }) diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp index f39a2e499fa57..10f36c5ec8f0e 100644 --- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp +++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp @@ -2805,7 +2805,8 @@ void LoweringPreparePass::runOnOp(mlir::Operation *op) { globalDtorList.emplace_back(fnOp.getName(), globalDtor.value()); if (mlir::Attribute attr = - fnOp->getAttr(cir::CUDAKernelNameAttr::getMnemonic())) { + fnOp->getAttr(cir::CUDAKernelNameAttr::getMnemonic()); + attr && !fnOp.isDeclaration()) { auto kernelNameAttr = dyn_cast<CUDAKernelNameAttr>(attr); llvm::StringRef kernelName = kernelNameAttr.getKernelName(); cudaKernelMap[kernelName] = fnOp; diff --git a/clang/test/CIR/CodeGenCUDA/kernel-declaration.cu b/clang/test/CIR/CodeGenCUDA/kernel-declaration.cu new file mode 100644 index 0000000000000..46bfc85f9af66 --- /dev/null +++ b/clang/test/CIR/CodeGenCUDA/kernel-declaration.cu @@ -0,0 +1,72 @@ +// A launch of a kernel defined in another TU must not register that kernel or +// embed this TU's fat binary. Match classic codegen for both declaration-only +// TUs and TUs that also define a diff erent kernel. + +// RUN: echo -n "GPU binary would be here." > %t.fatbin + +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-cir %s -x cuda \ +// RUN: -target-sdk-version=12.3 -fcuda-include-gpubinary %t.fatbin -o - \ +// RUN: | FileCheck %s --check-prefix=CIR --implicit-check-not=fatbin \ +// RUN: --implicit-check-not=Register --implicit-check-not=module_ctor + +// With no entities to register, the binary must not even be read. +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-cir %s -x cuda \ +// RUN: -target-sdk-version=12.3 -fcuda-include-gpubinary %t.nonexistent -o - \ +// RUN: | FileCheck %s --check-prefix=CIR --implicit-check-not=fatbin \ +// RUN: --implicit-check-not=Register --implicit-check-not=module_ctor + +// RUN: %clang_cc1 -triple x86_64-linux-gnu -fclangir -emit-llvm %s -x cuda \ +// RUN: -target-sdk-version=12.3 -fcuda-include-gpubinary %t.fatbin -o - \ +// RUN: | FileCheck %s --check-prefix=LAUNCH --implicit-check-not=fatbin \ +// RUN: --implicit-check-not=Register --implicit-check-not=module_ctor \ +// RUN: --implicit-check-not=llvm.global_ctors +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -x cuda \ +// RUN: -target-sdk-version=12.3 -fcuda-include-gpubinary %t.fatbin -o - \ +// RUN: | FileCheck %s --check-prefix=LAUNCH --implicit-check-not=fatbin \ +// RUN: --implicit-check-not=Register --implicit-check-not=module_ctor \ +// RUN: --implicit-check-not=llvm.global_ctors + +// RUN: %clang_cc1 -triple x86_64-linux-gnu -fclangir -emit-llvm %s -x hip \ +// RUN: -fhip-new-launch-api -fcuda-include-gpubinary %t.fatbin -o - \ +// RUN: | FileCheck %s --check-prefix=LAUNCH --implicit-check-not=fatbin \ +// RUN: --implicit-check-not=Register --implicit-check-not=module_ctor \ +// RUN: --implicit-check-not=llvm.global_ctors + +// RUN: %clang_cc1 -triple x86_64-linux-gnu -fclangir -emit-llvm %s -x cuda \ +// RUN: -DDEFINE_KERNEL -target-sdk-version=12.3 -fcuda-include-gpubinary %t.fatbin -o - \ +// RUN: | FileCheck %s --check-prefixes=LAUNCH,MIXED \ +// RUN: --implicit-check-not='call{{.*}} @__{{cuda|hip}}RegisterFunction({{.*}}ptr @{{[^,]*}}k0' +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -x cuda \ +// RUN: -DDEFINE_KERNEL -target-sdk-version=12.3 -fcuda-include-gpubinary %t.fatbin -o - \ +// RUN: | FileCheck %s --check-prefixes=LAUNCH,MIXED \ +// RUN: --implicit-check-not='call{{.*}} @__{{cuda|hip}}RegisterFunction({{.*}}ptr @{{[^,]*}}k0' +// RUN: %clang_cc1 -triple x86_64-linux-gnu -fclangir -emit-llvm %s -x hip \ +// RUN: -DDEFINE_KERNEL -fhip-new-launch-api -fcuda-include-gpubinary %t.fatbin -o - \ +// RUN: | FileCheck %s --check-prefixes=LAUNCH,MIXED \ +// RUN: --implicit-check-not='call{{.*}} @__{{cuda|hip}}RegisterFunction({{.*}}ptr @{{[^,]*}}k0' + +#include "Inputs/cuda.h" + +__global__ void k0(int); +#ifdef DEFINE_KERNEL +__global__ void k1(int); +#endif + +void launch() { + k0<<<1, 1>>>(42); +#ifdef DEFINE_KERNEL + k1<<<1, 1>>>(42); +#endif +} + +// Define after the launch to exercise upgrading an existing stub declaration. +#ifdef DEFINE_KERNEL +__global__ void k1(int) {} +#endif + +// CIR: cir.call @_Z17__device_stub__k0i( +// LAUNCH-DAG: call void @_Z17__device_stub__k0i(i32{{.*}} 42) +// LAUNCH-DAG: declare void @_Z17__device_stub__k0i( +// MIXED-DAG: define{{.*}} void @_Z17__device_stub__k1i( +// MIXED-DAG: call{{.*}} @__{{cuda|hip}}RegisterFunction({{.*}}ptr @{{[^,]*}}k1{{[^,]*}}, +// MIXED-DAG: define internal void @__{{cuda|hip}}_module_ctor() _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
