llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Martin Paul Lücke (martin-luecke) <details> <summary>Changes</summary> 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. --- Full diff: https://github.com/llvm/llvm-project/pull/234911.diff 4 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp (+7-1) - (modified) clang/lib/CIR/CodeGen/CIRGenCall.cpp (-11) - (modified) clang/test/CIR/CodeGenCUDA/kernel-call.cu (+2-2) - (added) clang/test/CIR/CodeGenCUDA/kernel-registration.cu (+50) ``````````diff 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 `````````` </details> https://github.com/llvm/llvm-project/pull/234911 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
