From 82d0173f72735404098cfcecc8f511e2f8a95cb1 Mon Sep 17 00:00:00 2001
From: Dmitry Sidorov <Dmitry.Sidorov@amd.com>
Date: Thu, 26 Mar 2026 14:57:41 +0100
Subject: [PATCH] [HIP][CUDA] Apply protected visibility to kernels and globals
 (#187784)

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().
---
 clang/lib/CodeGen/CodeGenModule.cpp           | 21 ++++++++
 .../CodeGenHIP/amdgcnspirv-visibility.cpp     | 50 +++++++++++++++++++
 2 files changed, 71 insertions(+)
 create mode 100644 clang/test/CodeGenHIP/amdgcnspirv-visibility.cpp

diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index ed517d244f9a..b4a24bcf03d7 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -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;
diff --git a/clang/test/CodeGenHIP/amdgcnspirv-visibility.cpp b/clang/test/CodeGenHIP/amdgcnspirv-visibility.cpp
new file mode 100644
index 000000000000..d1b42e236897
--- /dev/null
+++ b/clang/test/CodeGenHIP/amdgcnspirv-visibility.cpp
@@ -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;
+}