[clang] [CIR] Implement setGlobalVisibility (PR #195257)
Zaky Hermawan via cfe-commits
cfe-commits at lists.llvm.org
Fri May 15 03:22:59 PDT 2026
https://github.com/ZakyHermawan updated https://github.com/llvm/llvm-project/pull/195257
>From 9e693e6e5a3ccde79dc4e7c3ce4c186b6b1c47e2 Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 1 May 2026 18:26:16 +0700
Subject: [PATCH 1/5] [CIR] Implement setGlobalVisibility
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
.../clang/CIR/Interfaces/CIROpInterfaces.td | 12 +++
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 74 ++++++++++++++++++-
clang/lib/CIR/CodeGen/CIRGenModule.h | 15 +++-
3 files changed, 97 insertions(+), 4 deletions(-)
diff --git a/clang/include/clang/CIR/Interfaces/CIROpInterfaces.td b/clang/include/clang/CIR/Interfaces/CIROpInterfaces.td
index 898e28964eef0..181397c1809aa 100644
--- a/clang/include/clang/CIR/Interfaces/CIROpInterfaces.td
+++ b/clang/include/clang/CIR/Interfaces/CIROpInterfaces.td
@@ -145,6 +145,18 @@ let cppNamespace = "::cir" in {
}]
>,
InterfaceMethod<"",
+ "void", "setGlobalVisibility", (ins "cir::VisibilityKind":$val), [{}],
+ /*defaultImplementation=*/[{
+ $_op.setGlobalVisibility(val);
+ }]
+ >,
+ InterfaceMethod<"",
+ "void", "setLinkage", (ins "cir::GlobalLinkageKind":$val), [{}],
+ /*defaultImplementation=*/[{
+ $_op.setLinkage(val);
+ }]
+ >,
+ InterfaceMethod<"",
"bool", "isDSOLocal", (ins), [{}],
/*defaultImplementation=*/[{
return $_op.getDsoLocal();
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index af8fd52bef017..2239a07e3ef4b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -2702,9 +2702,77 @@ static bool shouldAssumeDSOLocal(const CIRGenModule &cgm,
return false;
}
-void CIRGenModule::setGlobalVisibility(mlir::Operation *gv,
+static void setGlobalVisibilityHelper(const CIRGenModule &cgm,
+ cir::CIRGlobalValueInterface gv,
+ cir::VisibilityKind visibility) {
+ gv.setGlobalVisibility(cir::VisibilityKind::Default);
+ // Also update MLIR symbol visibility to match linkage
+ if (auto globalOp = dyn_cast<cir::GlobalOp>(gv.getOperation()))
+ mlir::SymbolTable::setSymbolVisibility(globalOp,
+ cgm.getMLIRVisibility(globalOp));
+ else if (auto funcOp = dyn_cast<cir::FuncOp>(gv.getOperation()))
+ mlir::SymbolTable::setSymbolVisibility(
+ funcOp, cgm.getMLIRVisibilityFromCIRLinkage(funcOp.getLinkage()));
+}
+
+void CIRGenModule::setGlobalVisibility(cir::CIRGlobalValueInterface gv,
const NamedDecl *d) const {
- assert(!cir::MissingFeatures::opGlobalVisibility());
+ // Internal definitions always have default visibility.
+ if (gv.hasLocalLinkage()) {
+ setGlobalVisibilityHelper(*this, gv, cir::VisibilityKind::Default);
+ return;
+ }
+ if (!d)
+ return;
+
+ // Set visibility for definitions, and for declarations if requested globally
+ // or set explicitly.
+ LinkageInfo lv = d->getLinkageAndVisibility();
+
+ // OpenMP declare target variables must be visible to the host so they can
+ // be registered. We require protected visibility unless the variable has
+ // the DT_nohost modifier and does not need to be registered.
+ if (getASTContext().getLangOpts().OpenMP &&
+ getASTContext().getLangOpts().OpenMPIsTargetDevice && isa<VarDecl>(d) &&
+ d->hasAttr<OMPDeclareTargetDeclAttr>() &&
+ d->getAttr<OMPDeclareTargetDeclAttr>()->getDevType() !=
+ OMPDeclareTargetDeclAttr::DT_NoHost &&
+ lv.getVisibility() == HiddenVisibility) {
+ gv.setGlobalVisibility(cir::VisibilityKind::Protected);
+ 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 (getASTContext().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.setGlobalVisibility(cir::VisibilityKind::Protected);
+ return;
+ }
+ }
+
+ if (getASTContext().getLangOpts().HLSL && !d->isInExportDeclContext()) {
+ gv.setGlobalVisibility(cir::VisibilityKind::Hidden);
+ return;
+ }
+
+ assert(!cir::MissingFeatures::opGlobalDLLImportExport());
+
+ if (lv.isVisibilityExplicit() || getLangOpts().SetVisibilityForExternDecls ||
+ !gv.isDeclarationForLinker())
+ gv.setGlobalVisibility(getCIRVisibilityKind(lv.getVisibility()));
}
void CIRGenModule::setDSOLocal(cir::CIRGlobalValueInterface gv) const {
@@ -2724,7 +2792,7 @@ void CIRGenModule::setGVProperties(mlir::Operation *op,
void CIRGenModule::setGVPropertiesAux(mlir::Operation *op,
const NamedDecl *d) const {
- setGlobalVisibility(op, d);
+ setGlobalVisibility(cast<cir::CIRGlobalValueInterface>(op), d);
setDSOLocal(op);
assert(!cir::MissingFeatures::opGlobalPartition());
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 2869411015bc5..f388e84ad55ed 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -431,6 +431,19 @@ class CIRGenModule : public CIRGenTypeCache {
llvm_unreachable("unknown visibility!");
}
+ static cir::VisibilityKind getCIRVisibilityKind(Visibility v) {
+ switch (v) {
+ case DefaultVisibility:
+ return cir::VisibilityKind::Default;
+ case HiddenVisibility:
+ return cir::VisibilityKind::Hidden;
+ case ProtectedVisibility:
+ return cir::VisibilityKind::Protected;
+ }
+
+ llvm_unreachable("unknown visibility!");
+ }
+
llvm::DenseMap<mlir::Attribute, cir::GlobalOp> constantStringMap;
llvm::DenseMap<const UnnamedGlobalConstantDecl *, cir::GlobalOp>
unnamedGlobalConstantDeclMap;
@@ -599,7 +612,7 @@ class CIRGenModule : public CIRGenTypeCache {
mlir::Type convertType(clang::QualType type);
/// Set the visibility for the given global.
- void setGlobalVisibility(mlir::Operation *op, const NamedDecl *d) const;
+ void setGlobalVisibility(cir::CIRGlobalValueInterface gv, const NamedDecl *d) const;
void setDSOLocal(mlir::Operation *op) const;
void setDSOLocal(cir::CIRGlobalValueInterface gv) const;
>From 61051c6a6c4b5cb4ca5384acba93ce10822fa314 Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 1 May 2026 18:34:30 +0700
Subject: [PATCH 2/5] clang-format
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
clang/lib/CIR/CodeGen/CIRGenModule.h | 5 +++--
1 file changed, 3 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index f388e84ad55ed..66c222b2a0742 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -440,7 +440,7 @@ class CIRGenModule : public CIRGenTypeCache {
case ProtectedVisibility:
return cir::VisibilityKind::Protected;
}
-
+
llvm_unreachable("unknown visibility!");
}
@@ -612,7 +612,8 @@ class CIRGenModule : public CIRGenTypeCache {
mlir::Type convertType(clang::QualType type);
/// Set the visibility for the given global.
- void setGlobalVisibility(cir::CIRGlobalValueInterface gv, const NamedDecl *d) const;
+ void setGlobalVisibility(cir::CIRGlobalValueInterface gv,
+ const NamedDecl *d) const;
void setDSOLocal(mlir::Operation *op) const;
void setDSOLocal(cir::CIRGlobalValueInterface gv) const;
>From b750c0f052e05786f37d088a485699c5bce22c5e Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 15 May 2026 09:30:51 +0700
Subject: [PATCH 3/5] [CIR] Remove unnecessary helper, setGlobalVisibility
call, and add tests
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 23 ++---------
.../CIR/CodeGenCUDA/attribute-visibility.cu | 40 +++++++++++++++++++
2 files changed, 44 insertions(+), 19 deletions(-)
create mode 100644 clang/test/CIR/CodeGenCUDA/attribute-visibility.cu
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 2239a07e3ef4b..79135457b4a81 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -1170,7 +1170,6 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty,
if (const SectionAttr *sa = d->getAttr<SectionAttr>())
gv.setSectionAttr(builder.getStringAttr(sa->getName()));
}
- gv.setGlobalVisibility(getGlobalVisibilityAttrFromDecl(d).getValue());
// Handle XCore specific ABI requirements.
if (getTriple().getArch() == llvm::Triple::xcore)
@@ -2702,24 +2701,11 @@ static bool shouldAssumeDSOLocal(const CIRGenModule &cgm,
return false;
}
-static void setGlobalVisibilityHelper(const CIRGenModule &cgm,
- cir::CIRGlobalValueInterface gv,
- cir::VisibilityKind visibility) {
- gv.setGlobalVisibility(cir::VisibilityKind::Default);
- // Also update MLIR symbol visibility to match linkage
- if (auto globalOp = dyn_cast<cir::GlobalOp>(gv.getOperation()))
- mlir::SymbolTable::setSymbolVisibility(globalOp,
- cgm.getMLIRVisibility(globalOp));
- else if (auto funcOp = dyn_cast<cir::FuncOp>(gv.getOperation()))
- mlir::SymbolTable::setSymbolVisibility(
- funcOp, cgm.getMLIRVisibilityFromCIRLinkage(funcOp.getLinkage()));
-}
-
void CIRGenModule::setGlobalVisibility(cir::CIRGlobalValueInterface gv,
const NamedDecl *d) const {
// Internal definitions always have default visibility.
if (gv.hasLocalLinkage()) {
- setGlobalVisibilityHelper(*this, gv, cir::VisibilityKind::Default);
+ gv.setGlobalVisibility(cir::VisibilityKind::Default);
return;
}
if (!d)
@@ -2752,11 +2738,12 @@ void CIRGenModule::setGlobalVisibility(cir::CIRGlobalValueInterface gv,
if (isa<FunctionDecl>(d)) {
needsProtected =
d->hasAttr<CUDAGlobalAttr>() || d->hasAttr<DeviceKernelAttr>();
- } else if (const auto *vd = dyn_cast<VarDecl>(d))
+ } 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.setGlobalVisibility(cir::VisibilityKind::Protected);
return;
@@ -2764,8 +2751,7 @@ void CIRGenModule::setGlobalVisibility(cir::CIRGlobalValueInterface gv,
}
if (getASTContext().getLangOpts().HLSL && !d->isInExportDeclContext()) {
- gv.setGlobalVisibility(cir::VisibilityKind::Hidden);
- return;
+ llvm_unreachable("setGlobalVisibility: HLSL is NYI");
}
assert(!cir::MissingFeatures::opGlobalDLLImportExport());
@@ -2895,7 +2881,6 @@ void CIRGenModule::setFunctionAttributes(GlobalDecl globalDecl,
// recompute it here. This is a minimal fix for now.
if (!isLocalLinkage(getFunctionLinkage(globalDecl))) {
const Decl *decl = globalDecl.getDecl();
- func.setGlobalVisibility(getGlobalVisibilityAttrFromDecl(decl).getValue());
}
// If we plan on emitting this inline builtin, we can't treat it as a builtin.
diff --git a/clang/test/CIR/CodeGenCUDA/attribute-visibility.cu b/clang/test/CIR/CodeGenCUDA/attribute-visibility.cu
new file mode 100644
index 0000000000000..867c1a13bd11b
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/attribute-visibility.cu
@@ -0,0 +1,40 @@
+#include "Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
+// RUN: -fcuda-is-device -emit-cir -target-sdk-version=12.3 \
+// RUN: -fvisibility=hidden -fapply-global-visibility-to-externs \
+// RUN: -I%S/Inputs/ %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
+// RUN: -fcuda-is-device -emit-llvm -target-sdk-version=12.3 \
+// RUN: -fvisibility=hidden -fapply-global-visibility-to-externs \
+// RUN: -I%S/Inputs/ %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda \
+// RUN: -fcuda-is-device -emit-llvm -target-sdk-version=12.3 \
+// RUN: -fvisibility=hidden -fapply-global-visibility-to-externs \
+// RUN: -I%S/Inputs/ %s -o %t.ll
+// RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %s
+
+// CIR-DEVICE: cir.global protected {{.*}} @deviceVar = #cir.int<0>
+// LLVM-DEVICE: @deviceVar = protected addrspace(1) externally_initialized global i32 0
+// OGCG-DEVICE: @deviceVar = protected addrspace(1) externally_initialized global i32 0
+__attribute__((device)) __device__ int deviceVar;
+
+// CIR-DEVICE: cir.global protected constant {{.*}} @constantVar = #cir.int<0>
+// LLVM-DEVICE: @constantVar = protected addrspace(4) externally_initialized constant i32 0
+// OGCG-DEVICE: @constantVar = protected addrspace(4) externally_initialized constant i32 0
+__attribute__((constant)) __constant__ int constantVar;
+
+// CIR-DEVICE: cir.global protected {{.*}} @nonconstVal = #cir.int<42>
+// LLVM-DEVICE: @nonconstVal = protected addrspace(1) externally_initialized global i32 42
+// OGCG-DEVICE: @nonconstVal = protected addrspace(1) externally_initialized global i32 42
+__device__ int nonconstVal = 42;
+
+// CIR-DEVICE: cir.func {{.*}} protected {{.*}} @_Z10kernelFuncv()
+// LLVM-DEVICE: define protected void @_Z10kernelFuncv()
+// OGCG-DEVICE: define protected ptx_kernel void @_Z10kernelFuncv()
+__attribute__((global)) __global__ void kernelFunc() {
+}
>From aeaa448dc4bc6c0dd5c9ea49abfcbd04bca9a070 Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 15 May 2026 09:39:47 +0700
Subject: [PATCH 4/5] [CIR] setGlobalVisibility for OpenMP is NYI
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 79135457b4a81..31ae3afe98e1a 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -2724,7 +2724,7 @@ void CIRGenModule::setGlobalVisibility(cir::CIRGlobalValueInterface gv,
d->getAttr<OMPDeclareTargetDeclAttr>()->getDevType() !=
OMPDeclareTargetDeclAttr::DT_NoHost &&
lv.getVisibility() == HiddenVisibility) {
- gv.setGlobalVisibility(cir::VisibilityKind::Protected);
+ llvm_unreachable("setGlobalVisibility: OpenMP is NYI");
return;
}
>From 4e35b28fc4c95f042fb3aa4bae6a1cb0d7677e81 Mon Sep 17 00:00:00 2001
From: ZakyHermawan <zaky.hermawan9615 at gmail.com>
Date: Fri, 15 May 2026 17:22:38 +0700
Subject: [PATCH 5/5] [CIR] Remove duplicate on setLinkage and update test
Signed-off-by: ZakyHermawan <zaky.hermawan9615 at gmail.com>
---
clang/include/clang/CIR/Interfaces/CIROpInterfaces.td | 6 ------
clang/test/CIR/CodeGenCUDA/attribute-visibility.cu | 2 +-
2 files changed, 1 insertion(+), 7 deletions(-)
diff --git a/clang/include/clang/CIR/Interfaces/CIROpInterfaces.td b/clang/include/clang/CIR/Interfaces/CIROpInterfaces.td
index 79ef196150c50..fb256c4a26c2e 100644
--- a/clang/include/clang/CIR/Interfaces/CIROpInterfaces.td
+++ b/clang/include/clang/CIR/Interfaces/CIROpInterfaces.td
@@ -151,12 +151,6 @@ let cppNamespace = "::cir" in {
}]
>,
InterfaceMethod<"",
- "void", "setLinkage", (ins "cir::GlobalLinkageKind":$val), [{}],
- /*defaultImplementation=*/[{
- $_op.setLinkage(val);
- }]
- >,
- InterfaceMethod<"",
"bool", "isDSOLocal", (ins), [{}],
/*defaultImplementation=*/[{
return $_op.getDsoLocal();
diff --git a/clang/test/CIR/CodeGenCUDA/attribute-visibility.cu b/clang/test/CIR/CodeGenCUDA/attribute-visibility.cu
index 867c1a13bd11b..74af9cf01b344 100644
--- a/clang/test/CIR/CodeGenCUDA/attribute-visibility.cu
+++ b/clang/test/CIR/CodeGenCUDA/attribute-visibility.cu
@@ -34,7 +34,7 @@ __attribute__((constant)) __constant__ int constantVar;
__device__ int nonconstVal = 42;
// CIR-DEVICE: cir.func {{.*}} protected {{.*}} @_Z10kernelFuncv()
-// LLVM-DEVICE: define protected void @_Z10kernelFuncv()
+// LLVM-DEVICE: define protected ptx_kernel void @_Z10kernelFuncv()
// OGCG-DEVICE: define protected ptx_kernel void @_Z10kernelFuncv()
__attribute__((global)) __global__ void kernelFunc() {
}
More information about the cfe-commits
mailing list