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 1/4] [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 >From da4b85a52628a01a8a4a2bf0240e40bc307b7cc6 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Wed, 30 Sep 2026 08:29:14 -0500 Subject: [PATCH 2/4] Move parts to getKernelHandle Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp | 18 ++++++++------- .../CIR/CodeGenHIP/kernel-handle-linkage.hip | 22 +++++++++++++++---- 2 files changed, 28 insertions(+), 12 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp index 22bbcae08006d..02e0be7d12e41 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp @@ -358,17 +358,12 @@ void CIRGenNVCUDARuntime::emitDeviceStub(CIRGenFunction &cgf, cir::FuncOp fn, globalOp->setAttr("alignment", builder.getI64IntegerAttr( cgm.getPointerAlign().getQuantity())); - // The handle must track the kernel stub's linkage/visibility, not the - // global-op default (external). + // 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())); - 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); + cgm.setDSOLocal(static_cast<mlir::Operation *>(globalOp)); } // CUDA 9.0 changed the way to launch kernels. @@ -429,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 index a5c10d356b07a..3e6077912b744 100644 --- a/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip +++ b/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip @@ -10,10 +10,10 @@ // 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. +// 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) {} @@ -37,6 +37,14 @@ template <> __global__ void decl_only<int>(int *p) {} // 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); + +// 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); @@ -46,6 +54,8 @@ void launch(int *p) { hidden_kernel<<<1, 1>>>(p); decl_only<int><<<1, 1>>>(p); fwd_kernel<<<1, 1>>>(p); + hidden_ext_kernel<<<1, 1>>>(p); + static_hidden_kernel<<<1, 1>>>(p); } static __global__ void fwd_kernel(int *p) {} @@ -59,6 +69,8 @@ static __global__ void fwd_kernel(int *p) {} // 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 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 @@ -69,3 +81,5 @@ static __global__ void fwd_kernel(int *p) {} // 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 >From edcd7e529a19b9bff4fd5b0463e3db5842075ae7 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 5 Oct 2026 05:57:28 -0500 Subject: [PATCH 3/4] Add PIE test case with deviation explanation Signed-off-by: Steffen Holst Larsen <[email protected]> --- .../CIR/CodeGenHIP/kernel-handle-linkage.hip | 30 +++++++++++++++++++ 1 file changed, 30 insertions(+) diff --git a/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip b/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip index 3e6077912b744..fc8e9b52923ad 100644 --- a/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip +++ b/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip @@ -10,6 +10,15 @@ // 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-prefixes=PIE,PIE-CIR %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-prefixes=PIE,PIE-OGCG %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 @@ -40,6 +49,7 @@ 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 @@ -55,6 +65,7 @@ void launch(int *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); } @@ -70,6 +81,7 @@ static __global__ void fwd_kernel(int *p) {} // 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 @@ -83,3 +95,21 @@ static __global__ void fwd_kernel(int *p) {} // 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 + +// Under PIE, a handle defined in the executable can't be replaced by another +// definition, so CIR marks it dso_local as it would any global variable. +// Classic codegen copies the stub's dso_local while the stub is still a +// declaration, so its handles are not. Both are correct. A handle that is only +// declared is dso_local in neither, 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-CIR-DAG: @_Z10ext_kernelPi = dso_local constant ptr @_Z25__device_stub__ext_kernelPi, align 8 +// PIE-CIR-DAG: @_Z4tmplIiEvPT_ = linkonce_odr dso_local constant ptr @_Z19__device_stub__tmplIiEvPT_, comdat, align 8 +// PIE-CIR-DAG: @_Z4instIfEvPT_ = weak_odr dso_local constant ptr @_Z19__device_stub__instIfEvPT_, comdat, align 8 +// PIE-CIR-DAG: @_Z9decl_onlyIiEvPT_ = dso_local constant ptr @_Z24__device_stub__decl_onlyIiEvPT_, align 8 +// PIE-OGCG-DAG: @_Z10ext_kernelPi = constant ptr @_Z25__device_stub__ext_kernelPi, align 8 +// PIE-OGCG-DAG: @_Z4tmplIiEvPT_ = linkonce_odr constant ptr @_Z19__device_stub__tmplIiEvPT_, comdat, align 8 +// PIE-OGCG-DAG: @_Z4instIfEvPT_ = weak_odr constant ptr @_Z19__device_stub__instIfEvPT_, comdat, align 8 +// PIE-OGCG-DAG: @_Z9decl_onlyIiEvPT_ = constant ptr @_Z24__device_stub__decl_onlyIiEvPT_, align 8 >From 585ae8e56264fb48f955b309199cd8f9c2ae8083 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 5 Oct 2026 06:30:44 -0500 Subject: [PATCH 4/4] Add cases that show dso_local printed in LLVM IR Signed-off-by: Steffen Holst Larsen <[email protected]> --- .../CIR/CodeGenHIP/kernel-handle-linkage.hip | 21 +++++++++++++++++++ 1 file changed, 21 insertions(+) diff --git a/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip b/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip index fc8e9b52923ad..13671b9ef7061 100644 --- a/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip +++ b/clang/test/CIR/CodeGenHIP/kernel-handle-linkage.hip @@ -19,6 +19,15 @@ // RUN: -I%S/../CodeGenCUDA/Inputs/ -emit-llvm %s -o - \ // RUN: | FileCheck --check-prefixes=PIE,PIE-OGCG %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 @@ -97,6 +106,18 @@ static __global__ void fwd_kernel(int *p) {} // 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 CIR marks it dso_local as it would any global variable. // Classic codegen copies the stub's dso_local while the stub is still a _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
