shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

[HIP][CUDA] Apply protected visibility to kernels and globals (#187784)

Browse Source 
Add the visibility override in setGlobalVisibility(), following the
existing OpenMP precedent. Unlike the AMDGPU post-hoc override, this
check respects explicit [[gnu::visibility("hidden")]] attributes
via isVisibilityExplicit().
This commit is contained in:
Dmitry Sidorov 2026-03-26 14:57:41 +01:00 committed by GitHub
parent cc4727ae3b
commit 82d0173f72
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
2 changed files with 71 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

21
clang/lib/CodeGen/CodeGenModule.cpp
View File

@ -1901,6 +1901,27 @@ void CodeGenModule::setGlobalVisibility(llvm::GlobalValue *GV,
return;
}
// CUDA/HIP device kernels and global variables must be visible to the host
// so they can be registered / initialized. We require protected visibility
// unless the user explicitly requested hidden via an attribute.
if (Context.getLangOpts().CUDAIsDevice &&
LV.getVisibility() == HiddenVisibility && !LV.isVisibilityExplicit() &&
!D->hasAttr<OMPDeclareTargetDeclAttr>()) {
bool NeedsProtected = false;
if (isa<FunctionDecl>(D))
NeedsProtected =
D->hasAttr<CUDAGlobalAttr>() || D->hasAttr<DeviceKernelAttr>();
else if (const auto *VD = dyn_cast<VarDecl>(D))
NeedsProtected = VD->hasAttr<CUDADeviceAttr>() ||
VD->hasAttr<CUDAConstantAttr>() ||
VD->getType()->isCUDADeviceBuiltinSurfaceType() ||
VD->getType()->isCUDADeviceBuiltinTextureType();
if (NeedsProtected) {
GV->setVisibility(llvm::GlobalValue::ProtectedVisibility);
return;
}
}
if (Context.getLangOpts().HLSL && !D->isInExportDeclContext()) {
GV->setVisibility(llvm::GlobalValue::HiddenVisibility);
return;

50
clang/test/CodeGenHIP/amdgcnspirv-visibility.cpp Normal file
View File

@ -0,0 +1,50 @@
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility=default -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-DEFAULT %s
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility=protected -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-PROTECTED %s
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility=hidden -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-HIDDEN %s
// Mirrors clang/test/CodeGenCUDA/amdgpu-visibility.cu for the SPIR-V AMDGCN
// target. Verifies that device kernels and variables with hidden visibility get
// upgraded to protected, matching native AMDGPU behavior.
#define __device__ __attribute__((device))
#define __constant__ __attribute__((constant))
#define __global__ __attribute__((global))
// CHECK-DEFAULT-DAG: @c ={{.*}} addrspace(1) externally_initialized constant
// CHECK-DEFAULT-DAG: @g ={{.*}} addrspace(1) externally_initialized global
// CHECK-DEFAULT-DAG: @e = external addrspace(1) global
// CHECK-PROTECTED-DAG: @c = protected addrspace(1) externally_initialized constant
// CHECK-PROTECTED-DAG: @g = protected addrspace(1) externally_initialized global
// CHECK-PROTECTED-DAG: @e = external protected addrspace(1) global
// CHECK-HIDDEN-DAG: @c = protected addrspace(1) externally_initialized constant
// CHECK-HIDDEN-DAG: @g = protected addrspace(1) externally_initialized global
// CHECK-HIDDEN-DAG: @e = external protected addrspace(1) global
__constant__ int c;
__device__ int g;
extern __device__ int e;
// Explicit [[gnu::visibility("hidden")]] must be respected (not upgraded to
// protected), unlike the implicit -fvisibility=hidden flag.
// CHECK-DEFAULT-DAG: @h = hidden addrspace(1) externally_initialized global
// CHECK-PROTECTED-DAG: @h = hidden addrspace(1) externally_initialized global
// CHECK-HIDDEN-DAG: @h = hidden addrspace(1) externally_initialized global
__attribute__((visibility("hidden"))) __device__ int h;
// dummy one to hold reference to `e`.
__device__ int f() {
return e;
}
// CHECK-DEFAULT: define{{.*}} spir_kernel void @_Z3foov()
// CHECK-PROTECTED: define protected spir_kernel void @_Z3foov()
// CHECK-HIDDEN: define protected spir_kernel void @_Z3foov()
__global__ void foo() {
g = c;
}
// CHECK-DEFAULT: define hidden spir_kernel void @_Z3barv()
// CHECK-PROTECTED: define hidden spir_kernel void @_Z3barv()
// CHECK-HIDDEN: define hidden spir_kernel void @_Z3barv()
__attribute__((visibility("hidden"))) __global__ void bar() {
h = 1;
}