https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/227267
>From b82ae82755a3f42cd74b4ad09d8bdb367f6a2536 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Wed, 23 Sep 2026 02:02:47 -0500 Subject: [PATCH] [CIR][HIP] Match kernel handle linkage, visibility and comdat For HIP, CIRGen emits a host-side kernel handle, i.e. a global named after the kernel that points to its __device_stub__ function. The handle was always created with external linkage and default visibility. For kernels with internal linkage and for linkonce_odr template instantiations, every TU that launched the kernel therefore defined a strong external handle, and linking two such TUs failed with multiple-definition errors. Matching classic codegen, we give the handle the stub's linkage, visibility and dso_local, and put it in a trivial comdat under the same conditions. This is done in emitDeviceStub rather than getKernelHandle because the handle is created on the kernel's first reference, which can precede the point where the stub's linkage is set. Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp | 12 ++++ .../CIR/CodeGenHIP/kernel-handle-linkage.hip | 71 +++++++++++++++++++ 2 files changed, 83 insertions(+) create mode 100644 clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp index 79153ca788756..22bbcae08006d 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp @@ -357,6 +357,18 @@ void CIRGenNVCUDARuntime::emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn, globalOp->removeAttr("sym_visibility"); globalOp->setAttr("alignment", builder.getI64IntegerAttr( cgm.getPointerAlign().getQuantity())); + + // The handle must track the kernel stub's linkage/visibility, not the + // global-op default (external). + globalOp.setLinkage(fn.getLinkage()); + mlir::SymbolTable::setSymbolVisibility( + globalOp, cgm.getMLIRVisibilityFromCIRLinkage(fn.getLinkage())); + globalOp.setDSOLocal(fn.isDSOLocal()); + globalOp.setGlobalVisibility(fn.getGlobalVisibility()); + auto *fd = cast<FunctionDecl>(cgf.curGD.getDecl()); + FunctionTemplateDecl *ft = fd->getPrimaryTemplate(); + if (!ft || ft->isThisDeclarationADefinition()) + cgm.maybeSetTrivialComdat(*fd, globalOp); } // CUDA 9.0 changed the way to launch kernels. diff --git a/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip b/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip new file mode 100644 index 0000000000000..a5c10d356b07a --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip @@ -0,0 +1,71 @@ +#include "cuda.h" + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api \ +// RUN: -fclangir -I%S/../CodeGenCUDA/Inputs/ -emit-cir %s -o - \ +// RUN: | FileCheck --check-prefix=CIR %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api \ +// RUN: -fclangir -I%S/../CodeGenCUDA/Inputs/ -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api \ +// RUN: -I%S/../CodeGenCUDA/Inputs/ -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s + +// The host-side kernel handle takes the linkage and visibility of the kernel's +// device stub. Kernels with internal linkage must get an internal handle in +// every TU that uses them; an external one causes multiple-definition link +// errors. + +template <class T> static __global__ void static_tmpl(T *p) {} +template <class T> __global__ void tmpl(T *p) {} +static __global__ void static_kernel(int *p) {} +namespace { +__global__ void anon_kernel(int *p) {} +} +__global__ void ext_kernel(int *p) {} +__attribute__((visibility("hidden"))) __global__ void hidden_kernel(int *p) {} + +// Explicit instantiation definition: weak_odr, in a comdat. +template <class T> __global__ void inst(T *p) {} +template __global__ void inst<float>(float *p); + +// Explicit specialization of a template that is only declared: strong +// external, no comdat. +template <class T> __global__ void decl_only(T *p); +template <> __global__ void decl_only<int>(int *p) {} + +// Referenced before its definition, so the handle is created before the +// stub's linkage is known. +static __global__ void fwd_kernel(int *p); + +void launch(int *p) { + static_tmpl<int><<<1, 1>>>(p); + tmpl<int><<<1, 1>>>(p); + static_kernel<<<1, 1>>>(p); + anon_kernel<<<1, 1>>>(p); + ext_kernel<<<1, 1>>>(p); + hidden_kernel<<<1, 1>>>(p); + decl_only<int><<<1, 1>>>(p); + fwd_kernel<<<1, 1>>>(p); +} + +static __global__ void fwd_kernel(int *p) {} + +// CIR-DAG: cir.global constant external @_Z10ext_kernelPi = #cir.global_view<@_Z25__device_stub__ext_kernelPi> +// CIR-DAG: cir.global hidden constant external @_Z13hidden_kernelPi = #cir.global_view<@_Z28__device_stub__hidden_kernelPi> +// CIR-DAG: cir.global "private" constant internal dso_local @_ZL11static_tmplIiEvPT_ = #cir.global_view<@_ZL26__device_stub__static_tmplIiEvPT_> +// CIR-DAG: cir.global constant linkonce_odr comdat @_Z4tmplIiEvPT_ = #cir.global_view<@_Z19__device_stub__tmplIiEvPT_> +// CIR-DAG: cir.global "private" constant internal dso_local @_ZL13static_kernelPi = #cir.global_view<@_ZL28__device_stub__static_kernelPi> +// CIR-DAG: cir.global "private" constant internal dso_local @_ZN12_GLOBAL__N_111anon_kernelEPi = #cir.global_view<@_ZN12_GLOBAL__N_126__device_stub__anon_kernelEPi> +// CIR-DAG: cir.global "private" constant internal dso_local @_ZL10fwd_kernelPi = #cir.global_view<@_ZL25__device_stub__fwd_kernelPi> +// CIR-DAG: cir.global constant weak_odr comdat @_Z4instIfEvPT_ = #cir.global_view<@_Z19__device_stub__instIfEvPT_> +// CIR-DAG: cir.global constant external @_Z9decl_onlyIiEvPT_ = #cir.global_view<@_Z24__device_stub__decl_onlyIiEvPT_> + +// LLVM-DAG: @_Z10ext_kernelPi = constant ptr @_Z25__device_stub__ext_kernelPi, align 8 +// LLVM-DAG: @_Z13hidden_kernelPi = hidden constant ptr @_Z28__device_stub__hidden_kernelPi, align 8 +// LLVM-DAG: @_ZL11static_tmplIiEvPT_ = internal constant ptr @_ZL26__device_stub__static_tmplIiEvPT_, align 8 +// LLVM-DAG: @_ZL13static_kernelPi = internal constant ptr @_ZL28__device_stub__static_kernelPi, align 8 +// LLVM-DAG: @_ZN12_GLOBAL__N_111anon_kernelEPi = internal constant ptr @_ZN12_GLOBAL__N_126__device_stub__anon_kernelEPi, align 8 +// LLVM-DAG: @_ZL10fwd_kernelPi = internal constant ptr @_ZL25__device_stub__fwd_kernelPi, align 8 +// LLVM-DAG: @_Z4tmplIiEvPT_ = linkonce_odr constant ptr @_Z19__device_stub__tmplIiEvPT_, comdat, align 8 +// LLVM-DAG: @_Z4instIfEvPT_ = weak_odr constant ptr @_Z19__device_stub__instIfEvPT_, comdat, align 8 +// LLVM-DAG: @_Z9decl_onlyIiEvPT_ = constant ptr @_Z24__device_stub__decl_onlyIiEvPT_, align 8 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
