[clang] [CIR] Add addLLVMUsed and addLLVMCompilerUsed methods to CIRGenModule (PR #188189)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Apr 2 00:27:30 PDT 2026
https://github.com/skc7 updated https://github.com/llvm/llvm-project/pull/188189
>From 0f80a8a4345fb11d6f9846775f332806ba762de0 Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Tue, 24 Mar 2026 12:56:23 +0530
Subject: [PATCH 1/4] [CIR] Add addLLVMUsed and addLLVMCompilerUsed methods to
CIRGenModule
---
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 102 ++++++++++++++++++++++++-
clang/lib/CIR/CodeGen/CIRGenModule.h | 19 +++++
clang/test/CIR/CodeGenHIP/hip-cuid.hip | 27 +++++++
3 files changed, 146 insertions(+), 2 deletions(-)
create mode 100644 clang/test/CIR/CodeGenHIP/hip-cuid.hip
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index e0681eb760249..43600d33c87a9 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -669,7 +669,20 @@ void CIRGenModule::setCommonAttributes(GlobalDecl gd, mlir::Operation *gv) {
if (isa_and_nonnull<NamedDecl>(d))
setGVProperties(gv, dyn_cast<NamedDecl>(d));
assert(!cir::MissingFeatures::defaultVisibility());
- assert(!cir::MissingFeatures::opGlobalUsedOrCompilerUsed());
+
+ if (auto globalOp = mlir::dyn_cast<cir::GlobalOp>(gv)) {
+ if (d && d->hasAttr<UsedAttr>())
+ addUsedOrCompilerUsedGlobal(globalOp);
+
+ if (const auto *vd = dyn_cast_if_present<VarDecl>(d);
+ vd && ((codeGenOpts.KeepPersistentStorageVariables &&
+ (vd->getStorageDuration() == SD_Static ||
+ vd->getStorageDuration() == SD_Thread)) ||
+ (codeGenOpts.KeepStaticConsts &&
+ vd->getStorageDuration() == SD_Static &&
+ vd->getType().isConstQualified())))
+ addUsedOrCompilerUsedGlobal(globalOp);
+ }
}
void CIRGenModule::setNonAliasAttributes(GlobalDecl gd, mlir::Operation *op) {
@@ -684,7 +697,10 @@ void CIRGenModule::setNonAliasAttributes(GlobalDecl gd, mlir::Operation *op) {
}
assert(!cir::MissingFeatures::opGlobalPragmaClangSection());
- assert(!cir::MissingFeatures::opGlobalUsedOrCompilerUsed());
+ if (auto globalOp = mlir::dyn_cast<cir::GlobalOp>(op)) {
+ if (d->hasAttr<RetainAttr>())
+ addUsedGlobal(globalOp);
+ }
assert(!cir::MissingFeatures::opFuncCPUAndFeaturesAttributes());
getTargetCIRGenInfo().setTargetAttributes(gd.getDecl(), op, *this);
@@ -1095,6 +1111,62 @@ cir::GlobalViewAttr CIRGenModule::getAddrOfGlobalVarAttr(const VarDecl *d) {
return builder.getGlobalViewAttr(ptrTy, globalOp);
}
+void CIRGenModule::addUsedGlobal(cir::GlobalOp gv) {
+ assert(!gv.isDeclaration() &&
+ "Only globals with definition can force usage.");
+ LLVMUsed.emplace_back(gv);
+}
+
+void CIRGenModule::addCompilerUsedGlobal(cir::GlobalOp gv) {
+ assert(!gv.isDeclaration() &&
+ "Only globals with definition can force usage.");
+ LLVMCompilerUsed.emplace_back(gv);
+}
+
+void CIRGenModule::addUsedOrCompilerUsedGlobal(cir::GlobalOp gv) {
+ assert(!gv.isDeclaration() &&
+ "Only globals with definition can force usage.");
+ if (getTriple().isOSBinFormatELF())
+ LLVMCompilerUsed.emplace_back(gv);
+ else
+ LLVMUsed.emplace_back(gv);
+}
+
+static void emitUsed(CIRGenModule &cgm, StringRef name,
+ std::vector<cir::GlobalOp> &list) {
+ // Don't create llvm.used if there is no need.
+ if (list.empty())
+ return;
+
+ // Convert List to what ConstantArray needs.
+ auto &builder = cgm.getBuilder();
+ auto loc = builder.getUnknownLoc();
+ llvm::SmallVector<mlir::Attribute, 8> usedArray;
+ usedArray.resize(list.size());
+ for (unsigned i = 0, e = list.size(); i != e; ++i) {
+ usedArray[i] = cir::GlobalViewAttr::get(
+ cgm.voidPtrTy, mlir::FlatSymbolRefAttr::get(list[i].getSymNameAttr()));
+ }
+
+ if (usedArray.empty())
+ return;
+ auto arrayTy = cir::ArrayType::get(cgm.voidPtrTy, usedArray.size());
+
+ auto initAttr = cir::ConstArrayAttr::get(
+ arrayTy, mlir::ArrayAttr::get(&cgm.getMLIRContext(), usedArray));
+
+ auto gv = CIRGenModule::createGlobalOp(cgm, loc, name, arrayTy,
+ /*isConstant=*/false);
+ gv.setLinkage(cir::GlobalLinkageKind::AppendingLinkage);
+ gv.setInitialValueAttr(initAttr);
+ // TODO(CIR): Set section to "llvm.metadata" once GlobalOp supports sections.
+}
+
+void CIRGenModule::emitLLVMUsed() {
+ emitUsed(*this, "llvm.used", LLVMUsed);
+ emitUsed(*this, "llvm.compiler.used", LLVMCompilerUsed);
+}
+
void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd,
bool isTentative) {
if (getLangOpts().OpenCL || getLangOpts().OpenMPIsTargetDevice) {
@@ -3100,6 +3172,32 @@ void CIRGenModule::release() {
(getTriple().isSPIRV() && getTriple().getVendor() == llvm::Triple::AMD))
emitAMDGPUMetadata();
+ if (getLangOpts().HIP) {
+ // Emit a unique ID so that host and device binaries from the same
+ // compilation unit can be associated.
+ std::string cuidName =
+ ("__hip_cuid_" + getASTContext().getCUIDHash()).str();
+ auto int8Ty = cir::IntType::get(&getMLIRContext(), 8, /*isSigned=*/false);
+ auto loc = builder.getUnknownLoc();
+ mlir::ptr::MemorySpaceAttrInterface addrSpace =
+ cir::LangAddressSpaceAttr::get(&getMLIRContext(),
+ getGlobalVarAddressSpace(nullptr));
+
+ auto gv = createGlobalOp(*this, loc, cuidName, int8Ty,
+ /*isConstant=*/false, addrSpace);
+ gv.setLinkage(cir::GlobalLinkageKind::ExternalLinkage);
+ // Initialize with zero
+ auto zeroAttr = cir::IntAttr::get(int8Ty, 0);
+ gv.setInitialValueAttr(zeroAttr);
+ // External linkage requires public visibility
+ mlir::SymbolTable::setSymbolVisibility(
+ gv, mlir::SymbolTable::Visibility::Public);
+
+ addCompilerUsedGlobal(gv);
+ }
+
+ emitLLVMUsed();
+
// There's a lot of code that is not implemented yet.
assert(!cir::MissingFeatures::cgmRelease());
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 266510de84fd0..08d6965899d93 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -176,6 +176,19 @@ class CIRGenModule : public CIRGenTypeCache {
void mapResolvedBlockAddress(cir::BlockAddressOp op, cir::LabelOp);
void updateResolvedBlockAddress(cir::BlockAddressOp op,
cir::LabelOp newLabel);
+
+ /// Add a global value to the LLVMUsed list.
+ void addUsedGlobal(cir::GlobalOp gv);
+
+ /// Add a global value to the LLVMCompilerUsed list.
+ void addCompilerUsedGlobal(cir::GlobalOp gv);
+
+ /// Add a global to a list to be added to the llvm.compiler.used metadata.
+ void addUsedOrCompilerUsedGlobal(cir::GlobalOp gv);
+
+ /// Emit llvm.used and llvm.compiler.used globals.
+ void emitLLVMUsed();
+
/// Tell the consumer that this variable has been instantiated.
void handleCXXStaticMemberVarInstantiation(VarDecl *vd);
@@ -440,6 +453,12 @@ class CIRGenModule : public CIRGenTypeCache {
cir::FuncType fnType = nullptr, bool dontDefer = false,
ForDefinition_t isForDefinition = NotForDefinition);
+ /// List of global values which are required to be present in the object file;
+ /// This is used for forcing visibility of symbols which may otherwise be
+ /// optimized out.
+ std::vector<cir::GlobalOp> LLVMUsed;
+ std::vector<cir::GlobalOp> LLVMCompilerUsed;
+
mlir::Type getVTableComponentType();
CIRGenVTables &getVTables() { return vtables; }
diff --git a/clang/test/CIR/CodeGenHIP/hip-cuid.hip b/clang/test/CIR/CodeGenHIP/hip-cuid.hip
new file mode 100644
index 0000000000000..8622ae75bc34d
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/hip-cuid.hip
@@ -0,0 +1,27 @@
+#include "../CodeGenCUDA/Inputs/cuda.h"
+
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ll
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \
+// RUN: -fcuda-is-device -emit-llvm %s -o %t.ogcg.ll
+// RUN: FileCheck --check-prefix=OGCG %s --input-file=%t.ogcg.ll
+
+// Test that HIP compiler unit ID global is emitted
+
+// CIR: cir.global external lang_address_space(offload_global) @__hip_cuid_{{.*}} = #cir.int<0> : !u8i
+
+// TODO(CIR): Should emit addrspace(1) once LangAddressSpace lowering is supported.
+// LLVM: @__hip_cuid_{{.*}} = global i8 0
+// LLVM: @llvm.compiler.used = {{.*}}@__hip_cuid_
+
+// OGCG: @__hip_cuid_{{.*}} = addrspace(1) global i8 0
+// OGCG: @llvm.compiler.used = {{.*}}@__hip_cuid_
+
+__global__ void kernel() {}
>From 30fefdb4fac1f53d3c09592dad968d5d62cb47a2 Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Fri, 27 Mar 2026 16:27:04 +0530
Subject: [PATCH 2/4] use CIRGlobalValueInterface
---
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 59 +++++++++++++-------------
clang/lib/CIR/CodeGen/CIRGenModule.h | 14 +++---
clang/test/CIR/CodeGen/attr-retain.c | 18 ++++++++
clang/test/CIR/CodeGen/attr-used.c | 14 ++++++
4 files changed, 69 insertions(+), 36 deletions(-)
create mode 100644 clang/test/CIR/CodeGen/attr-retain.c
create mode 100644 clang/test/CIR/CodeGen/attr-used.c
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 43600d33c87a9..b4e306c4b2b1a 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -670,9 +670,9 @@ void CIRGenModule::setCommonAttributes(GlobalDecl gd, mlir::Operation *gv) {
setGVProperties(gv, dyn_cast<NamedDecl>(d));
assert(!cir::MissingFeatures::defaultVisibility());
- if (auto globalOp = mlir::dyn_cast<cir::GlobalOp>(gv)) {
+ if (auto gvi = mlir::dyn_cast<cir::CIRGlobalValueInterface>(gv)) {
if (d && d->hasAttr<UsedAttr>())
- addUsedOrCompilerUsedGlobal(globalOp);
+ addUsedOrCompilerUsedGlobal(gvi);
if (const auto *vd = dyn_cast_if_present<VarDecl>(d);
vd && ((codeGenOpts.KeepPersistentStorageVariables &&
@@ -681,7 +681,7 @@ void CIRGenModule::setCommonAttributes(GlobalDecl gd, mlir::Operation *gv) {
(codeGenOpts.KeepStaticConsts &&
vd->getStorageDuration() == SD_Static &&
vd->getType().isConstQualified())))
- addUsedOrCompilerUsedGlobal(globalOp);
+ addUsedOrCompilerUsedGlobal(gvi);
}
}
@@ -693,6 +693,8 @@ void CIRGenModule::setNonAliasAttributes(GlobalDecl gd, mlir::Operation *op) {
if (auto gvi = mlir::dyn_cast<cir::CIRGlobalValueInterface>(op)) {
if (const auto *sa = d->getAttr<SectionAttr>())
gvi.setSection(builder.getStringAttr(sa->getName()));
+ if (d->hasAttr<RetainAttr>())
+ addUsedGlobal(gvi);
}
}
@@ -1111,60 +1113,59 @@ cir::GlobalViewAttr CIRGenModule::getAddrOfGlobalVarAttr(const VarDecl *d) {
return builder.getGlobalViewAttr(ptrTy, globalOp);
}
-void CIRGenModule::addUsedGlobal(cir::GlobalOp gv) {
- assert(!gv.isDeclaration() &&
+void CIRGenModule::addUsedGlobal(cir::CIRGlobalValueInterface gv) {
+ assert((mlir::isa<cir::FuncOp>(gv.getOperation()) ||
+ !gv.isDeclarationForLinker()) &&
"Only globals with definition can force usage.");
- LLVMUsed.emplace_back(gv);
+ llvmUsed.emplace_back(gv);
}
-void CIRGenModule::addCompilerUsedGlobal(cir::GlobalOp gv) {
- assert(!gv.isDeclaration() &&
+void CIRGenModule::addCompilerUsedGlobal(cir::CIRGlobalValueInterface gv) {
+ assert(!gv.isDeclarationForLinker() &&
"Only globals with definition can force usage.");
- LLVMCompilerUsed.emplace_back(gv);
+ llvmCompilerUsed.emplace_back(gv);
}
-void CIRGenModule::addUsedOrCompilerUsedGlobal(cir::GlobalOp gv) {
- assert(!gv.isDeclaration() &&
+void CIRGenModule::addUsedOrCompilerUsedGlobal(
+ cir::CIRGlobalValueInterface gv) {
+ assert((mlir::isa<cir::FuncOp>(gv.getOperation()) ||
+ !gv.isDeclarationForLinker()) &&
"Only globals with definition can force usage.");
if (getTriple().isOSBinFormatELF())
- LLVMCompilerUsed.emplace_back(gv);
+ llvmCompilerUsed.emplace_back(gv);
else
- LLVMUsed.emplace_back(gv);
+ llvmUsed.emplace_back(gv);
}
static void emitUsed(CIRGenModule &cgm, StringRef name,
- std::vector<cir::GlobalOp> &list) {
- // Don't create llvm.used if there is no need.
+ std::vector<cir::CIRGlobalValueInterface> &list) {
if (list.empty())
return;
- // Convert List to what ConstantArray needs.
- auto &builder = cgm.getBuilder();
- auto loc = builder.getUnknownLoc();
- llvm::SmallVector<mlir::Attribute, 8> usedArray;
+ CIRGenBuilderTy &builder = cgm.getBuilder();
+ mlir::Location loc = builder.getUnknownLoc();
+ llvm::SmallVector<mlir::Attribute> usedArray;
usedArray.resize(list.size());
- for (unsigned i = 0, e = list.size(); i != e; ++i) {
+ for (auto [i, op] : llvm::enumerate(list)) {
usedArray[i] = cir::GlobalViewAttr::get(
- cgm.voidPtrTy, mlir::FlatSymbolRefAttr::get(list[i].getSymNameAttr()));
+ cgm.voidPtrTy, mlir::FlatSymbolRefAttr::get(op.getNameAttr()));
}
- if (usedArray.empty())
- return;
- auto arrayTy = cir::ArrayType::get(cgm.voidPtrTy, usedArray.size());
+ cir::ArrayType arrayTy = cir::ArrayType::get(cgm.voidPtrTy, usedArray.size());
- auto initAttr = cir::ConstArrayAttr::get(
+ cir::ConstArrayAttr initAttr = cir::ConstArrayAttr::get(
arrayTy, mlir::ArrayAttr::get(&cgm.getMLIRContext(), usedArray));
- auto gv = CIRGenModule::createGlobalOp(cgm, loc, name, arrayTy,
- /*isConstant=*/false);
+ cir::GlobalOp gv = CIRGenModule::createGlobalOp(cgm, loc, name, arrayTy,
+ /*isConstant=*/false);
gv.setLinkage(cir::GlobalLinkageKind::AppendingLinkage);
gv.setInitialValueAttr(initAttr);
// TODO(CIR): Set section to "llvm.metadata" once GlobalOp supports sections.
}
void CIRGenModule::emitLLVMUsed() {
- emitUsed(*this, "llvm.used", LLVMUsed);
- emitUsed(*this, "llvm.compiler.used", LLVMCompilerUsed);
+ emitUsed(*this, "llvm.used", llvmUsed);
+ emitUsed(*this, "llvm.compiler.used", llvmCompilerUsed);
}
void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd,
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 08d6965899d93..ac202fd096608 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -177,14 +177,14 @@ class CIRGenModule : public CIRGenTypeCache {
void updateResolvedBlockAddress(cir::BlockAddressOp op,
cir::LabelOp newLabel);
- /// Add a global value to the LLVMUsed list.
- void addUsedGlobal(cir::GlobalOp gv);
+ /// Add a global value to the llvmUsed list.
+ void addUsedGlobal(cir::CIRGlobalValueInterface gv);
- /// Add a global value to the LLVMCompilerUsed list.
- void addCompilerUsedGlobal(cir::GlobalOp gv);
+ /// Add a global value to the llvmCompilerUsed list.
+ void addCompilerUsedGlobal(cir::CIRGlobalValueInterface gv);
/// Add a global to a list to be added to the llvm.compiler.used metadata.
- void addUsedOrCompilerUsedGlobal(cir::GlobalOp gv);
+ void addUsedOrCompilerUsedGlobal(cir::CIRGlobalValueInterface gv);
/// Emit llvm.used and llvm.compiler.used globals.
void emitLLVMUsed();
@@ -456,8 +456,8 @@ class CIRGenModule : public CIRGenTypeCache {
/// List of global values which are required to be present in the object file;
/// This is used for forcing visibility of symbols which may otherwise be
/// optimized out.
- std::vector<cir::GlobalOp> LLVMUsed;
- std::vector<cir::GlobalOp> LLVMCompilerUsed;
+ std::vector<cir::CIRGlobalValueInterface> llvmUsed;
+ std::vector<cir::CIRGlobalValueInterface> llvmCompilerUsed;
mlir::Type getVTableComponentType();
CIRGenVTables &getVTables() { return vtables; }
diff --git a/clang/test/CIR/CodeGen/attr-retain.c b/clang/test/CIR/CodeGen/attr-retain.c
new file mode 100644
index 0000000000000..28597aa5adce2
--- /dev/null
+++ b/clang/test/CIR/CodeGen/attr-retain.c
@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir %s -o - | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm %s -o - | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=LLVM
+
+const int c0 __attribute__((retain)) = 42;
+
+__attribute__((retain)) int g0;
+int g1 __attribute__((retain));
+__attribute__((used, retain)) static int g3;
+
+void __attribute__((retain)) f0(void) {}
+static void __attribute__((used, retain)) f2(void) {}
+
+// CIR: cir.global "private" appending @llvm.used = #cir.const_array<[#cir.global_view<@c0> : !cir.ptr<!void>, #cir.global_view<@f0> : !cir.ptr<!void>, #cir.global_view<@f2> : !cir.ptr<!void>, #cir.global_view<@g0> : !cir.ptr<!void>, #cir.global_view<@g1> : !cir.ptr<!void>, #cir.global_view<@g3> : !cir.ptr<!void>]> : !cir.array<!cir.ptr<!void> x 6>
+// CIR: cir.global "private" appending @llvm.compiler.used = #cir.const_array<[#cir.global_view<@f2> : !cir.ptr<!void>, #cir.global_view<@g3> : !cir.ptr<!void>]> : !cir.array<!cir.ptr<!void> x 2>
+
+// LLVM: @llvm.used = appending global [6 x ptr] [ptr @c0, ptr @f0, ptr @f2, ptr @g0, ptr @g1, ptr @g3]
+// LLVM: @llvm.compiler.used = appending global [2 x ptr] [ptr @f2, ptr @g3]
diff --git a/clang/test/CIR/CodeGen/attr-used.c b/clang/test/CIR/CodeGen/attr-used.c
new file mode 100644
index 0000000000000..f609f822e39df
--- /dev/null
+++ b/clang/test/CIR/CodeGen/attr-used.c
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir %s -o - | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm %s -o - | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=LLVM
+
+int g0 __attribute__((used));
+
+static void __attribute__((used)) f0(void) {
+}
+
+__attribute__((used)) int a0;
+void pr27535(void) { (void)a0; }
+
+// CIR: cir.global "private" appending @llvm.compiler.used = #cir.const_array<[#cir.global_view<@f0> : !cir.ptr<!void>, #cir.global_view<@g0> : !cir.ptr<!void>, #cir.global_view<@a0> : !cir.ptr<!void>]> : !cir.array<!cir.ptr<!void> x 3>
+// LLVM: @llvm.compiler.used = appending global [3 x ptr] [ptr @f0, ptr @g0, ptr @a0]
>From 0981c8827ccf52149dc5f9c5598cd303407be19c Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Tue, 31 Mar 2026 12:48:55 +0530
Subject: [PATCH 3/4] add tests persistent-storage-variables and
keep-static-consts
---
.../keep-persistent-storage-variables.cpp | 20 +++++++++++++++++++
clang/test/CIR/CodeGen/keep-static-consts.cpp | 11 ++++++++++
2 files changed, 31 insertions(+)
create mode 100644 clang/test/CIR/CodeGen/keep-persistent-storage-variables.cpp
create mode 100644 clang/test/CIR/CodeGen/keep-static-consts.cpp
diff --git a/clang/test/CIR/CodeGen/keep-persistent-storage-variables.cpp b/clang/test/CIR/CodeGen/keep-persistent-storage-variables.cpp
new file mode 100644
index 0000000000000..a2520675052b0
--- /dev/null
+++ b/clang/test/CIR/CodeGen/keep-persistent-storage-variables.cpp
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -fkeep-persistent-storage-variables -triple x86_64-unknown-linux-gnu -fclangir -emit-cir %s -o - | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -fkeep-persistent-storage-variables -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm %s -o - | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -fkeep-persistent-storage-variables -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=LLVM
+
+static int g1;
+static int g2 = 1;
+int g5;
+int g6 = 6;
+
+namespace {
+ int s4 = 42;
+}
+
+struct ST {
+ static int s6;
+};
+int ST::s6 = 7;
+
+// CIR: cir.global "private" appending @llvm.compiler.used = #cir.const_array<[#cir.global_view<@_ZL2g1> : !cir.ptr<!void>, #cir.global_view<@_ZL2g2> : !cir.ptr<!void>, #cir.global_view<@g5> : !cir.ptr<!void>, #cir.global_view<@g6> : !cir.ptr<!void>, #cir.global_view<@_ZN12_GLOBAL__N_12s4E> : !cir.ptr<!void>, #cir.global_view<@_ZN2ST2s6E> : !cir.ptr<!void>]> : !cir.array<!cir.ptr<!void> x 6>
+// LLVM: @llvm.compiler.used = appending global [6 x ptr] [ptr @_ZL2g1, ptr @_ZL2g2, ptr @g5, ptr @g6, ptr @_ZN12_GLOBAL__N_12s4E, ptr @_ZN2ST2s6E]
diff --git a/clang/test/CIR/CodeGen/keep-static-consts.cpp b/clang/test/CIR/CodeGen/keep-static-consts.cpp
new file mode 100644
index 0000000000000..ef88e54c187a0
--- /dev/null
+++ b/clang/test/CIR/CodeGen/keep-static-consts.cpp
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 -fkeep-static-consts -triple x86_64-unknown-linux-gnu -fclangir -emit-cir %s -o - | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -fkeep-static-consts -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm %s -o - | FileCheck %s --check-prefix=LLVM
+// RUN: %clang_cc1 -fkeep-static-consts -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix=LLVM
+
+static const char srcvers[] = "xyz";
+extern const int b = 1;
+const char srcvers2[] = "abc";
+constexpr int N = 2;
+
+// CIR: cir.global "private" appending @llvm.compiler.used = #cir.const_array<[#cir.global_view<@_ZL7srcvers> : !cir.ptr<!void>, #cir.global_view<@b> : !cir.ptr<!void>, #cir.global_view<@_ZL8srcvers2> : !cir.ptr<!void>, #cir.global_view<@_ZL1N> : !cir.ptr<!void>]> : !cir.array<!cir.ptr<!void> x 4>
+// LLVM: @llvm.compiler.used = appending global [4 x ptr] [ptr @_ZL7srcvers, ptr @b, ptr @_ZL8srcvers2, ptr @_ZL1N]
>From 29810d7566a574592df5863679dbb85bfb989c7a Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Thu, 2 Apr 2026 11:24:14 +0530
Subject: [PATCH 4/4] add gv section attribute
---
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 6 +-----
clang/test/CIR/CodeGen/attr-retain.c | 4 ++--
clang/test/CIR/CodeGen/attr-used.c | 2 +-
.../test/CIR/CodeGen/keep-persistent-storage-variables.cpp | 2 +-
clang/test/CIR/CodeGen/keep-static-consts.cpp | 2 +-
5 files changed, 6 insertions(+), 10 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index b4e306c4b2b1a..dea2bbde34717 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -699,10 +699,6 @@ void CIRGenModule::setNonAliasAttributes(GlobalDecl gd, mlir::Operation *op) {
}
assert(!cir::MissingFeatures::opGlobalPragmaClangSection());
- if (auto globalOp = mlir::dyn_cast<cir::GlobalOp>(op)) {
- if (d->hasAttr<RetainAttr>())
- addUsedGlobal(globalOp);
- }
assert(!cir::MissingFeatures::opFuncCPUAndFeaturesAttributes());
getTargetCIRGenInfo().setTargetAttributes(gd.getDecl(), op, *this);
@@ -1160,7 +1156,7 @@ static void emitUsed(CIRGenModule &cgm, StringRef name,
/*isConstant=*/false);
gv.setLinkage(cir::GlobalLinkageKind::AppendingLinkage);
gv.setInitialValueAttr(initAttr);
- // TODO(CIR): Set section to "llvm.metadata" once GlobalOp supports sections.
+ gv.setSectionAttr(builder.getStringAttr("llvm.metadata"));
}
void CIRGenModule::emitLLVMUsed() {
diff --git a/clang/test/CIR/CodeGen/attr-retain.c b/clang/test/CIR/CodeGen/attr-retain.c
index 28597aa5adce2..dd409e4713444 100644
--- a/clang/test/CIR/CodeGen/attr-retain.c
+++ b/clang/test/CIR/CodeGen/attr-retain.c
@@ -14,5 +14,5 @@ static void __attribute__((used, retain)) f2(void) {}
// CIR: cir.global "private" appending @llvm.used = #cir.const_array<[#cir.global_view<@c0> : !cir.ptr<!void>, #cir.global_view<@f0> : !cir.ptr<!void>, #cir.global_view<@f2> : !cir.ptr<!void>, #cir.global_view<@g0> : !cir.ptr<!void>, #cir.global_view<@g1> : !cir.ptr<!void>, #cir.global_view<@g3> : !cir.ptr<!void>]> : !cir.array<!cir.ptr<!void> x 6>
// CIR: cir.global "private" appending @llvm.compiler.used = #cir.const_array<[#cir.global_view<@f2> : !cir.ptr<!void>, #cir.global_view<@g3> : !cir.ptr<!void>]> : !cir.array<!cir.ptr<!void> x 2>
-// LLVM: @llvm.used = appending global [6 x ptr] [ptr @c0, ptr @f0, ptr @f2, ptr @g0, ptr @g1, ptr @g3]
-// LLVM: @llvm.compiler.used = appending global [2 x ptr] [ptr @f2, ptr @g3]
+// LLVM: @llvm.used = appending global [6 x ptr] [ptr @c0, ptr @f0, ptr @f2, ptr @g0, ptr @g1, ptr @g3], section "llvm.metadata"
+// LLVM: @llvm.compiler.used = appending global [2 x ptr] [ptr @f2, ptr @g3], section "llvm.metadata"
diff --git a/clang/test/CIR/CodeGen/attr-used.c b/clang/test/CIR/CodeGen/attr-used.c
index f609f822e39df..51c961992e062 100644
--- a/clang/test/CIR/CodeGen/attr-used.c
+++ b/clang/test/CIR/CodeGen/attr-used.c
@@ -11,4 +11,4 @@ __attribute__((used)) int a0;
void pr27535(void) { (void)a0; }
// CIR: cir.global "private" appending @llvm.compiler.used = #cir.const_array<[#cir.global_view<@f0> : !cir.ptr<!void>, #cir.global_view<@g0> : !cir.ptr<!void>, #cir.global_view<@a0> : !cir.ptr<!void>]> : !cir.array<!cir.ptr<!void> x 3>
-// LLVM: @llvm.compiler.used = appending global [3 x ptr] [ptr @f0, ptr @g0, ptr @a0]
+// LLVM: @llvm.compiler.used = appending global [3 x ptr] [ptr @f0, ptr @g0, ptr @a0], section "llvm.metadata"
diff --git a/clang/test/CIR/CodeGen/keep-persistent-storage-variables.cpp b/clang/test/CIR/CodeGen/keep-persistent-storage-variables.cpp
index a2520675052b0..71aea031e1f58 100644
--- a/clang/test/CIR/CodeGen/keep-persistent-storage-variables.cpp
+++ b/clang/test/CIR/CodeGen/keep-persistent-storage-variables.cpp
@@ -17,4 +17,4 @@ struct ST {
int ST::s6 = 7;
// CIR: cir.global "private" appending @llvm.compiler.used = #cir.const_array<[#cir.global_view<@_ZL2g1> : !cir.ptr<!void>, #cir.global_view<@_ZL2g2> : !cir.ptr<!void>, #cir.global_view<@g5> : !cir.ptr<!void>, #cir.global_view<@g6> : !cir.ptr<!void>, #cir.global_view<@_ZN12_GLOBAL__N_12s4E> : !cir.ptr<!void>, #cir.global_view<@_ZN2ST2s6E> : !cir.ptr<!void>]> : !cir.array<!cir.ptr<!void> x 6>
-// LLVM: @llvm.compiler.used = appending global [6 x ptr] [ptr @_ZL2g1, ptr @_ZL2g2, ptr @g5, ptr @g6, ptr @_ZN12_GLOBAL__N_12s4E, ptr @_ZN2ST2s6E]
+// LLVM: @llvm.compiler.used = appending global [6 x ptr] [ptr @_ZL2g1, ptr @_ZL2g2, ptr @g5, ptr @g6, ptr @_ZN12_GLOBAL__N_12s4E, ptr @_ZN2ST2s6E], section "llvm.metadata"
diff --git a/clang/test/CIR/CodeGen/keep-static-consts.cpp b/clang/test/CIR/CodeGen/keep-static-consts.cpp
index ef88e54c187a0..76e953574704f 100644
--- a/clang/test/CIR/CodeGen/keep-static-consts.cpp
+++ b/clang/test/CIR/CodeGen/keep-static-consts.cpp
@@ -8,4 +8,4 @@ const char srcvers2[] = "abc";
constexpr int N = 2;
// CIR: cir.global "private" appending @llvm.compiler.used = #cir.const_array<[#cir.global_view<@_ZL7srcvers> : !cir.ptr<!void>, #cir.global_view<@b> : !cir.ptr<!void>, #cir.global_view<@_ZL8srcvers2> : !cir.ptr<!void>, #cir.global_view<@_ZL1N> : !cir.ptr<!void>]> : !cir.array<!cir.ptr<!void> x 4>
-// LLVM: @llvm.compiler.used = appending global [4 x ptr] [ptr @_ZL7srcvers, ptr @b, ptr @_ZL8srcvers2, ptr @_ZL1N]
+// LLVM: @llvm.compiler.used = appending global [4 x ptr] [ptr @_ZL7srcvers, ptr @b, ptr @_ZL8srcvers2, ptr @_ZL1N], section "llvm.metadata"
More information about the cfe-commits
mailing list