https://github.com/ZakyHermawan updated https://github.com/llvm/llvm-project/pull/195257
>From 9e693e6e5a3ccde79dc4e7c3ce4c186b6b1c47e2 Mon Sep 17 00:00:00 2001 From: ZakyHermawan <[email protected]> Date: Fri, 1 May 2026 18:26:16 +0700 Subject: [PATCH 1/7] [CIR] Implement setGlobalVisibility Signed-off-by: ZakyHermawan <[email protected]> --- .../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 <[email protected]> Date: Fri, 1 May 2026 18:34:30 +0700 Subject: [PATCH 2/7] clang-format Signed-off-by: ZakyHermawan <[email protected]> --- 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 <[email protected]> Date: Fri, 15 May 2026 09:30:51 +0700 Subject: [PATCH 3/7] [CIR] Remove unnecessary helper, setGlobalVisibility call, and add tests Signed-off-by: ZakyHermawan <[email protected]> --- 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 <[email protected]> Date: Fri, 15 May 2026 09:39:47 +0700 Subject: [PATCH 4/7] [CIR] setGlobalVisibility for OpenMP is NYI Signed-off-by: ZakyHermawan <[email protected]> --- 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 <[email protected]> Date: Fri, 15 May 2026 17:22:38 +0700 Subject: [PATCH 5/7] [CIR] Remove duplicate on setLinkage and update test Signed-off-by: ZakyHermawan <[email protected]> --- 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() { } >From cddac59b846aecee48878478edef8efe946a3242 Mon Sep 17 00:00:00 2001 From: ZakyHermawan <[email protected]> Date: Fri, 15 May 2026 17:34:05 +0700 Subject: [PATCH 6/7] [CIR] Fix unused error Signed-off-by: ZakyHermawan <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 8 -------- 1 file changed, 8 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index fd8abb2c7cf98..c2c6df65c2026 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -2927,14 +2927,6 @@ void CIRGenModule::setFunctionAttributes(GlobalDecl globalDecl, if (!isIncompleteFunction && func.isDeclaration()) getTargetCIRGenInfo().setTargetAttributes(funcDecl, func, *this); - // TODO(cir): This needs a lot of work to better match CodeGen. That - // ultimately ends up in setGlobalVisibility, which already has the linkage of - // the LLVM GV (corresponding to our FuncOp) computed, so it doesn't have to - // recompute it here. This is a minimal fix for now. - if (!isLocalLinkage(getFunctionLinkage(globalDecl))) { - const Decl *decl = globalDecl.getDecl(); - } - // If we plan on emitting this inline builtin, we can't treat it as a builtin. if (funcDecl->isInlineBuiltinDeclaration()) { const FunctionDecl *fdBody; >From 55e2f45839bf803e1e422033754014565efcb4a5 Mon Sep 17 00:00:00 2001 From: ZakyHermawan <[email protected]> Date: Fri, 15 May 2026 17:51:42 +0700 Subject: [PATCH 7/7] [CIR] Recover hlsl setGlobalVisibility handling Signed-off-by: ZakyHermawan <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index c2c6df65c2026..7b94d459c99fd 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -2796,7 +2796,8 @@ void CIRGenModule::setGlobalVisibility(cir::CIRGlobalValueInterface gv, } if (getASTContext().getLangOpts().HLSL && !d->isInExportDeclContext()) { - llvm_unreachable("setGlobalVisibility: HLSL is NYI"); + gv.setGlobalVisibility(cir::VisibilityKind::Hidden); + return; } assert(!cir::MissingFeatures::opGlobalDLLImportExport()); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
