https://github.com/skc7 updated https://github.com/llvm/llvm-project/pull/188189
>From 85d3261d357a089e064b96db16d0b4e0d1ff9a10 Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Tue, 24 Mar 2026 12:56:23 +0530 Subject: [PATCH 1/2] [CIR] Add addLLVMUsed and addLLVMCompilerUsed methods to CIRGenModule --- clang/include/clang/CIR/MissingFeatures.h | 1 - clang/lib/CIR/CodeGen/CIRGenModule.cpp | 106 +++++++++++++++++++++- clang/lib/CIR/CodeGen/CIRGenModule.h | 19 ++++ clang/test/CIR/CodeGenHIP/hip-cuid.hip | 27 ++++++ 4 files changed, 150 insertions(+), 3 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/hip-cuid.hip diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h index 5051044eb5d50..b580bdb7387b6 100644 --- a/clang/include/clang/CIR/MissingFeatures.h +++ b/clang/include/clang/CIR/MissingFeatures.h @@ -35,7 +35,6 @@ struct MissingFeatures { static bool opGlobalVisibility() { return false; } static bool opGlobalDLLImportExport() { return false; } static bool opGlobalPartition() { return false; } - static bool opGlobalUsedOrCompilerUsed() { return false; } static bool opGlobalAnnotations() { return false; } static bool opGlobalCtorPriority() { return false; } static bool setDSOLocal() { return false; } diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index de68927089873..119d79b861943 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -658,13 +658,33 @@ 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) { setCommonAttributes(gd, op); - assert(!cir::MissingFeatures::opGlobalUsedOrCompilerUsed()); + const Decl *d = gd.getDecl(); + if (d) { + if (auto globalOp = mlir::dyn_cast<cir::GlobalOp>(op)) { + if (d->hasAttr<RetainAttr>()) + addUsedGlobal(globalOp); + } + } + assert(!cir::MissingFeatures::opGlobalSection()); assert(!cir::MissingFeatures::opFuncCPUAndFeaturesAttributes()); assert(!cir::MissingFeatures::opFuncSection()); @@ -1071,6 +1091,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) { @@ -3071,6 +3147,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 7bf50e9b340a58fa8e5cb62c367d91eba2fd959a Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Fri, 27 Mar 2026 16:27:04 +0530 Subject: [PATCH 2/2] use CIRGlobalValueInterface --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 61 +++++++++++++------------- 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(+), 38 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 119d79b861943..1b15273d6f065 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -659,9 +659,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 && @@ -670,7 +670,7 @@ void CIRGenModule::setCommonAttributes(GlobalDecl gd, mlir::Operation *gv) { (codeGenOpts.KeepStaticConsts && vd->getStorageDuration() == SD_Static && vd->getType().isConstQualified()))) - addUsedOrCompilerUsedGlobal(globalOp); + addUsedOrCompilerUsedGlobal(gvi); } } @@ -679,9 +679,9 @@ void CIRGenModule::setNonAliasAttributes(GlobalDecl gd, mlir::Operation *op) { const Decl *d = gd.getDecl(); if (d) { - if (auto globalOp = mlir::dyn_cast<cir::GlobalOp>(op)) { + if (auto gvi = mlir::dyn_cast<cir::CIRGlobalValueInterface>(op)) { if (d->hasAttr<RetainAttr>()) - addUsedGlobal(globalOp); + addUsedGlobal(gvi); } } @@ -1091,60 +1091,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] _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
