https://github.com/steffenlarsen created https://github.com/llvm/llvm-project/pull/229060
The host-side kernel handle is a global variable, but getKernelHandle copied dso_local from the device stub while the stub was still a function declaration. emitDeviceStub then updated the handle's linkage, but not its dso_local. As a result, dso_local followed the rules for function declarations rather than those for the handle: - Under PIE, and with -fno-plt or -fno-direct-access-external-data in static builds, defined handles were not dso_local, though a definition in an executable can't be preempted. Every reference to them went through the GOT. - With static -fno-plt, a declared handle was not dso_local, though data can use a copy relocation. -fno-plt only affects calls. - On MinGW, a declared handle was dso_local, though a variable may be auto-imported from a DLL. Compute dso_local with CGM.setDSOLocal when the handle is created and again once emitDeviceStub gives it its final linkage. The result is what any other global variable with the same linkage and visibility gets. CUDA is unaffected, as its kernel handle is the stub itself. Assisted-by: Claude Opus 5.5 >From c173f9b42d8eb946b2bc836c308b196d02fe85cf Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 5 Oct 2026 07:00:48 -0500 Subject: [PATCH] [HIP] Compute kernel handle dso_local as for a variable The host-side kernel handle is a global variable, but getKernelHandle copied dso_local from the device stub while the stub was still a function declaration. emitDeviceStub then updated the handle's linkage, but not its dso_local. As a result, dso_local followed the rules for function declarations rather than those for the handle: - Under PIE, and with -fno-plt or -fno-direct-access-external-data in static builds, defined handles were not dso_local, though a definition in an executable can't be preempted. Every reference to them went through the GOT. - With static -fno-plt, a declared handle was not dso_local, though data can use a copy relocation. -fno-plt only affects calls. - On MinGW, a declared handle was dso_local, though a variable may be auto-imported from a DLL. Compute dso_local with CGM.setDSOLocal when the handle is created and again once emitDeviceStub gives it its final linkage. The result is what any other global variable with the same linkage and visibility gets. CUDA is unaffected, as its kernel handle is the stub itself. Assisted-by: Claude Opus 5.5 Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CodeGen/CGCUDANV.cpp | 3 +- .../CodeGenCUDA/kernel-handle-dso-local.cu | 57 +++++++++++++++++++ 2 files changed, 59 insertions(+), 1 deletion(-) create mode 100644 clang/test/CodeGenCUDA/kernel-handle-dso-local.cu diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp index 260c2aed44ee3..ac2ed9a109c4f 100644 --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -335,6 +335,7 @@ void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF, dyn_cast<llvm::GlobalVariable>(KernelHandles[CGF.CurFn->getName()])) { GV->setLinkage(CGF.CurFn->getLinkage()); GV->setInitializer(CGF.CurFn); + CGM.setDSOLocal(GV); } if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(), CudaFeature::CUDA_USES_NEW_LAUNCH) || @@ -1532,8 +1533,8 @@ llvm::GlobalValue *CGNVCUDARuntime::getKernelHandle(llvm::Function *F, CGM.getMangledName( GD.getWithKernelReferenceKind(KernelReferenceKind::Kernel))); Var->setAlignment(CGM.getPointerAlign().getAsAlign()); - Var->setDSOLocal(F->isDSOLocal()); Var->setVisibility(F->getVisibility()); + CGM.setDSOLocal(Var); auto *FD = cast<FunctionDecl>(GD.getDecl()); auto *FT = FD->getPrimaryTemplate(); if (!FT || FT->isThisDeclarationADefinition()) diff --git a/clang/test/CodeGenCUDA/kernel-handle-dso-local.cu b/clang/test/CodeGenCUDA/kernel-handle-dso-local.cu new file mode 100644 index 0000000000000..af649873704b6 --- /dev/null +++ b/clang/test/CodeGenCUDA/kernel-handle-dso-local.cu @@ -0,0 +1,57 @@ +// Check that HIP kernel handles are dso_local according to the rules for global +// variables, with the handle's final linkage. + +// Shared library. +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip \ +// RUN: -mrelocation-model pic -pic-level 2 \ +// RUN: | FileCheck -check-prefixes=CHECK,DEF-NONLOCAL,DECL-NONLOCAL %s + +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip \ +// RUN: -mrelocation-model pic -pic-level 2 -pic-is-pie \ +// RUN: | FileCheck -check-prefixes=CHECK,DEF-LOCAL,DECL-NONLOCAL %s + +// A declared handle may be reached through a copy relocation, which -fno-plt +// does not affect, as the handle is not a function. +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip \ +// RUN: -mrelocation-model static \ +// RUN: | FileCheck -check-prefixes=CHECK,DEF-LOCAL,DECL-LOCAL %s +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip \ +// RUN: -mrelocation-model static -fno-plt \ +// RUN: | FileCheck -check-prefixes=CHECK,DEF-LOCAL,DECL-LOCAL %s +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -o - -x hip \ +// RUN: -mrelocation-model static -fno-direct-access-external-data \ +// RUN: | FileCheck -check-prefixes=CHECK,DEF-LOCAL,DECL-NONLOCAL %s + +// A declared handle may be auto-imported from a DLL. +// RUN: %clang_cc1 -triple x86_64-w64-windows-gnu -emit-llvm %s -o - -x hip \ +// RUN: | FileCheck -check-prefixes=CHECK,DEF-LOCAL,DECL-NONLOCAL %s + +#include "Inputs/cuda.h" + +__global__ void ext_kernel() {} +template <class T> __global__ void tmpl_kernel() {} +template <class T> __global__ void inst_kernel() {} +template __global__ void inst_kernel<float>(); +__global__ void decl_kernel(); +static __global__ void static_kernel() {} + +void launch() { + ext_kernel<<<1, 1>>>(); + tmpl_kernel<int><<<1, 1>>>(); + decl_kernel<<<1, 1>>>(); + static_kernel<<<1, 1>>>(); +} + +// DEF-LOCAL-DAG: @_Z10ext_kernelv = dso_local constant ptr @_Z25__device_stub__ext_kernelv, align 8 +// DEF-LOCAL-DAG: @_Z11tmpl_kernelIiEvv = linkonce_odr dso_local constant ptr @_Z26__device_stub__tmpl_kernelIiEvv, comdat, align 8 +// DEF-LOCAL-DAG: @_Z11inst_kernelIfEvv = weak_odr dso_local constant ptr @_Z26__device_stub__inst_kernelIfEvv, comdat, align 8 + +// DEF-NONLOCAL-DAG: @_Z10ext_kernelv = constant ptr @_Z25__device_stub__ext_kernelv, align 8 +// DEF-NONLOCAL-DAG: @_Z11tmpl_kernelIiEvv = linkonce_odr constant ptr @_Z26__device_stub__tmpl_kernelIiEvv, comdat, align 8 +// DEF-NONLOCAL-DAG: @_Z11inst_kernelIfEvv = weak_odr constant ptr @_Z26__device_stub__inst_kernelIfEvv, comdat, align 8 + +// DECL-LOCAL-DAG: @_Z11decl_kernelv = external dso_local constant ptr, align 8 +// DECL-NONLOCAL-DAG: @_Z11decl_kernelv = external constant ptr, align 8 + +// Internal handles are always dso_local, so it isn't printed. +// CHECK-DAG: @_ZL13static_kernelv = internal constant ptr @_ZL28__device_stub__static_kernelv, align 8 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
