https://github.com/skc7 created https://github.com/llvm/llvm-project/pull/188189
Upstreaming clangIR PR: https://github.com/llvm/clangir/pull/2092 This PR adds support for emitting llvm.used and llvm.compiler.used global arrays in CIR. Added addUsedGlobal() and addCompilerUsedGlobal() methods to CIRGenModule Adds __hip_cuid_* to llvm.compiler.used for HIP compilation. Followed OGCG implementation in clang/lib/CodeGen/CodeGenModule.cpp >From 7f6c96b2cd67319be1375cbb98e05e7b9f710eca Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Tue, 24 Mar 2026 12:56:23 +0530 Subject: [PATCH] [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 68db08a5580ca..12d1297f53c0e 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 f3ab733bf4c6a..969f7cfc04dca 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) { @@ -3066,6 +3142,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() {} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
