Author: Steffen Larsen Date: 2026-10-20T07:38:57Z New Revision: 6cc9f0e55f06c81e2a4a8cd4e344803dd9ffc987
URL: https://github.com/llvm/llvm-project/commit/6cc9f0e55f06c81e2a4a8cd4e344803dd9ffc987 DIFF: https://github.com/llvm/llvm-project/commit/6cc9f0e55f06c81e2a4a8cd4e344803dd9ffc987.diff LOG: [CIR][HIP] Match kernel handle linkage, visibility and comdat (#227267) 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. Assisted-by: Claude Opus 5.5 --------- Signed-off-by: Steffen Holst Larsen <[email protected]> Added: clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip Modified: clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp index 8f7d65b827507..dfd311de78d54 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp @@ -355,9 +355,15 @@ void CIRGenNVCUDARuntime::emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn, auto gv = cir::GlobalViewAttr::get(fnPtrTy, sym); globalOp->setAttr("initial_value", gv); - globalOp->removeAttr("sym_visibility"); globalOp->setAttr("alignment", builder.getI64IntegerAttr( cgm.getPointerAlign().getQuantity())); + + // The stub's linkage is only final once it is defined, so the handle + // takes it over here. + globalOp.setLinkage(fn.getLinkage()); + mlir::SymbolTable::setSymbolVisibility( + globalOp, cgm.getMLIRVisibilityFromCIRLinkage(fn.getLinkage())); + cgm.setDSOLocal(static_cast<mlir::Operation *>(globalOp)); } // CUDA 9.0 changed the way to launch kernels. @@ -418,6 +424,13 @@ mlir::Operation *CIRGenNVCUDARuntime::getKernelHandle(cir::FuncOp fn, globalOp->setAttr("alignment", builder.getI64IntegerAttr( cgm.getPointerAlign().getQuantity())); + // Inherit visibility and comdat from the kernel's declaration. + auto *fd = cast<FunctionDecl>(gd.getDecl()); + cgm.setGVProperties(globalOp, fd); + FunctionTemplateDecl *ft = fd->getPrimaryTemplate(); + if (!ft || ft->isThisDeclarationADefinition()) + cgm.maybeSetTrivialComdat(*fd, globalOp); + // Store references kernelHandles[fn.getSymName()] = globalOp; kernelStubs[globalOp] = fn; 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..c518a0f7feb7d --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip @@ -0,0 +1,130 @@ +#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 + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api \ +// RUN: -mrelocation-model pic -pic-level 2 -pic-is-pie \ +// RUN: -fclangir -I%S/../CodeGenCUDA/Inputs/ -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=PIE %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api \ +// RUN: -mrelocation-model pic -pic-level 2 -pic-is-pie \ +// RUN: -I%S/../CodeGenCUDA/Inputs/ -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=PIE %s + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api \ +// RUN: -mrelocation-model static \ +// RUN: -fclangir -I%S/../CodeGenCUDA/Inputs/ -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=STATIC %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api \ +// RUN: -mrelocation-model static \ +// RUN: -I%S/../CodeGenCUDA/Inputs/ -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=STATIC %s + +// The host-side kernel handle takes the linkage of the kernel's device stub and +// the visibility of the kernel's declaration. Kernels with internal linkage +// must get an internal handle in every TU that uses them as 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); + +// Defined in another TU, so the stub is never emitted here; the visibility +// must come from the declaration. +__attribute__((visibility("hidden"))) __global__ void hidden_ext_kernel(int *p); +__global__ void ext_decl_kernel(int *p); + +// Internal linkage wins over the explicit visibility. +static __attribute__((visibility("hidden"))) __global__ void +static_hidden_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); + hidden_ext_kernel<<<1, 1>>>(p); + ext_decl_kernel<<<1, 1>>>(p); + static_hidden_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_> +// CIR-DAG: cir.global "private" hidden constant external @_Z17hidden_ext_kernelPi : +// CIR-DAG: cir.global "private" constant external @_Z15ext_decl_kernelPi : +// CIR-DAG: cir.global "private" constant internal dso_local @_ZL20static_hidden_kernelPi = #cir.global_view<@_ZL35__device_stub__static_hidden_kernelPi> + +// 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 +// LLVM-DAG: @_Z17hidden_ext_kernelPi = external hidden constant ptr, align 8 +// LLVM-DAG: @_ZL20static_hidden_kernelPi = internal constant ptr @_ZL35__device_stub__static_hidden_kernelPi, align 8 +// LLVM-DAG: @_Z15ext_decl_kernelPi = external constant ptr, align 8 + +// In a non-PIC executable, handles are dso_local, even when only declared, as +// a handle defined in a shared library is reached through a copy relocation. +// Internal and hidden handles are implicitly dso_local, so it isn't printed. +// STATIC-DAG: @_Z10ext_kernelPi = dso_local constant ptr @_Z25__device_stub__ext_kernelPi, align 8 +// STATIC-DAG: @_Z13hidden_kernelPi = hidden constant ptr @_Z28__device_stub__hidden_kernelPi, align 8 +// STATIC-DAG: @_ZL13static_kernelPi = internal constant ptr @_ZL28__device_stub__static_kernelPi, align 8 +// STATIC-DAG: @_Z4tmplIiEvPT_ = linkonce_odr dso_local constant ptr @_Z19__device_stub__tmplIiEvPT_, comdat, align 8 +// STATIC-DAG: @_Z4instIfEvPT_ = weak_odr dso_local constant ptr @_Z19__device_stub__instIfEvPT_, comdat, align 8 +// STATIC-DAG: @_Z9decl_onlyIiEvPT_ = dso_local constant ptr @_Z24__device_stub__decl_onlyIiEvPT_, align 8 +// STATIC-DAG: @_Z17hidden_ext_kernelPi = external hidden constant ptr, align 8 +// STATIC-DAG: @_Z15ext_decl_kernelPi = external dso_local constant ptr, align 8 + +// Under PIE, a handle defined in the executable can't be replaced by another +// definition, so it is dso_local as any global variable would be. A handle that +// is only declared is not, as it may be defined in a shared library. +// PIE-DAG: @_ZL13static_kernelPi = internal constant ptr @_ZL28__device_stub__static_kernelPi, align 8 +// PIE-DAG: @_Z17hidden_ext_kernelPi = external hidden constant ptr, align 8 +// PIE-DAG: @_Z15ext_decl_kernelPi = external constant ptr, align 8 +// PIE-DAG: @_Z10ext_kernelPi = dso_local constant ptr @_Z25__device_stub__ext_kernelPi, align 8 +// PIE-DAG: @_Z4tmplIiEvPT_ = linkonce_odr dso_local constant ptr @_Z19__device_stub__tmplIiEvPT_, comdat, align 8 +// PIE-DAG: @_Z4instIfEvPT_ = weak_odr dso_local constant ptr @_Z19__device_stub__instIfEvPT_, comdat, align 8 +// PIE-DAG: @_Z9decl_onlyIiEvPT_ = dso_local 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
