[clang] [CIR] Address Space support for GlobalOps (PR #179082)
David Rivera via cfe-commits
cfe-commits at lists.llvm.org
Mon Mar 2 14:50:58 PST 2026
https://github.com/RiverDave updated https://github.com/llvm/llvm-project/pull/179082
>From d31f7d9cdc3d3a67686a28d9dd543b57288d02e9 Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Sat, 31 Jan 2026 20:32:24 -0500
Subject: [PATCH 1/5] [CIR] Address Space support for GlobalOps
---
.../CIR/Dialect/Builder/CIRBaseBuilder.h | 14 +++--
clang/include/clang/CIR/Dialect/IR/CIROps.td | 3 +
clang/lib/CIR/CodeGen/CIRGenBuilder.h | 7 ++-
clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 4 +-
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 62 +++++++++++++++----
clang/lib/CIR/CodeGen/CIRGenModule.h | 20 ++++--
clang/lib/CIR/CodeGen/TargetInfo.cpp | 9 +++
clang/lib/CIR/CodeGen/TargetInfo.h | 7 +++
clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 23 ++++++-
clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 15 +++++
.../Dialect/Transforms/LoweringPrepare.cpp | 9 ++-
.../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 12 ++--
clang/test/CIR/IR/address-space.cir | 30 +++++++++
clang/test/CIR/IR/invalid-addrspace.cir | 20 ++++++
.../CIR/Lowering/global-address-space.cir | 46 ++++++++++++++
15 files changed, 244 insertions(+), 37 deletions(-)
create mode 100644 clang/test/CIR/Lowering/global-address-space.cir
diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
index 32d0921d15363..298bca80f6398 100644
--- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
+++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
@@ -19,6 +19,7 @@
#include "llvm/IR/FPEnv.h"
#include "llvm/Support/ErrorHandling.h"
+#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/Location.h"
@@ -362,14 +363,15 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
return CIRBaseBuilderTy::createStore(loc, flag, dst);
}
- [[nodiscard]] cir::GlobalOp createGlobal(mlir::ModuleOp mlirModule,
- mlir::Location loc,
- mlir::StringRef name,
- mlir::Type type, bool isConstant,
- cir::GlobalLinkageKind linkage) {
+ [[nodiscard]] cir::GlobalOp
+ createGlobal(mlir::ModuleOp mlirModule, mlir::Location loc,
+ mlir::StringRef name, mlir::Type type, bool isConstant,
+ cir::GlobalLinkageKind linkage,
+ mlir::ptr::MemorySpaceAttrInterface addrSpace) {
mlir::OpBuilder::InsertionGuard guard(*this);
setInsertionPointToStart(mlirModule.getBody());
- return cir::GlobalOp::create(*this, loc, name, type, isConstant, linkage);
+ return cir::GlobalOp::create(*this, loc, name, type, isConstant, addrSpace,
+ linkage);
}
cir::GetMemberOp createGetMember(mlir::Location loc, mlir::Type resultTy,
diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td
index a5a5197cd3ea6..78db29ba772fc 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIROps.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td
@@ -2422,6 +2422,7 @@ def CIR_GlobalOp : CIR_Op<"global", [
OptionalAttr<StrAttr>:$sym_visibility,
TypeAttr:$sym_type,
CIR_GlobalLinkageKind:$linkage,
+ OptionalAttr<MemorySpaceAttrInterface>:$addr_space,
OptionalAttr<CIR_TLSModel>:$tls_model,
OptionalAttr<AnyAttr>:$initial_value,
UnitAttr:$comdat,
@@ -2443,6 +2444,7 @@ def CIR_GlobalOp : CIR_Op<"global", [
($tls_model^)?
(`dso_local` $dso_local^)?
(`static_local_guard` `` $static_local_guard^)?
+ (` ` custom<GlobalAddressSpaceValue>($addr_space)^ )?
$sym_name
custom<GlobalOpTypeAndInitialValue>($sym_type, $initial_value,
$ctorRegion, $dtorRegion)
@@ -2463,6 +2465,7 @@ def CIR_GlobalOp : CIR_Op<"global", [
"llvm::StringRef":$sym_name,
"mlir::Type":$sym_type,
CArg<"bool", "false">:$isConstant,
+ CArg<"mlir::ptr::MemorySpaceAttrInterface", "{}">:$addrSpace,
// CIR defaults to external linkage.
CArg<"cir::GlobalLinkageKind",
"cir::GlobalLinkageKind::ExternalLinkage">:$linkage,
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuilder.h b/clang/lib/CIR/CodeGen/CIRGenBuilder.h
index cd13c9578adf7..ac0bf3b7d2f98 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuilder.h
+++ b/clang/lib/CIR/CodeGen/CIRGenBuilder.h
@@ -12,6 +12,7 @@
#include "Address.h"
#include "CIRGenRecordLayout.h"
#include "CIRGenTypeCache.h"
+#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/Support/LLVM.h"
@@ -678,7 +679,8 @@ class CIRGenBuilderTy : public cir::CIRBaseBuilderTy {
[[nodiscard]] cir::GlobalOp
createVersionedGlobal(mlir::ModuleOp module, mlir::Location loc,
mlir::StringRef name, mlir::Type type, bool isConstant,
- cir::GlobalLinkageKind linkage) {
+ cir::GlobalLinkageKind linkage,
+ mlir::ptr::MemorySpaceAttrInterface addrSpace = {}) {
// Create a unique name if the given name is already taken.
std::string uniqueName;
if (unsigned version = globalsVersioning[name.str()]++)
@@ -686,7 +688,8 @@ class CIRGenBuilderTy : public cir::CIRBaseBuilderTy {
else
uniqueName = name.str();
- return createGlobal(module, loc, uniqueName, type, isConstant, linkage);
+ return createGlobal(module, loc, uniqueName, type, isConstant, linkage,
+ addrSpace);
}
cir::StackSaveOp createStackSave(mlir::Location loc, mlir::Type ty) {
diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
index 4eccf430cd622..2429dca077667 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
@@ -296,9 +296,9 @@ static LValue emitGlobalVarDeclLValue(CIRGenFunction &cgf, const Expr *e,
// as part of getAddrOfGlobalVar.
mlir::Value v = cgf.cgm.getAddrOfGlobalVar(vd);
- assert(!cir::MissingFeatures::addressSpace());
mlir::Type realVarTy = cgf.convertTypeForMem(vd->getType());
- cir::PointerType realPtrTy = cgf.getBuilder().getPointerTo(realVarTy);
+ cir::PointerType realPtrTy = cir::PointerType::get(
+ realVarTy, mlir::cast<cir::PointerType>(v.getType()).getAddrSpace());
if (realPtrTy != v.getType())
v = cgf.getBuilder().createBitcast(v.getLoc(), v, realPtrTy);
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 223b53731359a..23b91ddc6ece0 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -25,12 +25,14 @@
#include "clang/Basic/SourceManager.h"
#include "clang/CIR/Dialect/IR/CIRAttrs.h"
#include "clang/CIR/Dialect/IR/CIRDialect.h"
+#include "clang/CIR/Dialect/IR/CIROpsEnums.h"
#include "clang/CIR/Dialect/IR/CIRTypes.h"
#include "clang/CIR/Interfaces/CIROpInterfaces.h"
#include "clang/CIR/MissingFeatures.h"
#include "CIRGenFunctionInfo.h"
#include "TargetInfo.h"
+#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/Location.h"
#include "mlir/IR/MLIRContext.h"
@@ -590,10 +592,11 @@ mlir::Operation *CIRGenModule::getGlobalValue(StringRef name) {
return mlir::SymbolTable::lookupSymbolIn(theModule, name);
}
-cir::GlobalOp CIRGenModule::createGlobalOp(CIRGenModule &cgm,
- mlir::Location loc, StringRef name,
- mlir::Type t, bool isConstant,
- mlir::Operation *insertPoint) {
+cir::GlobalOp
+CIRGenModule::createGlobalOp(CIRGenModule &cgm, mlir::Location loc,
+ StringRef name, mlir::Type t, bool isConstant,
+ mlir::ptr::MemorySpaceAttrInterface addrSpace,
+ mlir::Operation *insertPoint) {
cir::GlobalOp g;
CIRGenBuilderTy &builder = cgm.getBuilder();
@@ -613,7 +616,7 @@ cir::GlobalOp CIRGenModule::createGlobalOp(CIRGenModule &cgm,
builder.setInsertionPointToStart(cgm.getModule().getBody());
}
- g = cir::GlobalOp::create(builder, loc, name, t, isConstant);
+ g = cir::GlobalOp::create(builder, loc, name, t, isConstant, addrSpace);
if (!insertPoint)
cgm.lastGlobalOp = g;
@@ -662,6 +665,39 @@ std::optional<cir::SourceLanguage> CIRGenModule::getCIRSourceLanguage() const {
return std::nullopt;
}
+LangAS CIRGenModule::getGlobalVarAddressSpace(const VarDecl *d) {
+ if (langOpts.OpenCL) {
+ LangAS as = d ? d->getType().getAddressSpace() : LangAS::opencl_global;
+ assert(as == LangAS::opencl_global || as == LangAS::opencl_global_device ||
+ as == LangAS::opencl_global_host || as == LangAS::opencl_constant ||
+ as == LangAS::opencl_local || as >= LangAS::FirstTargetAddressSpace);
+ return as;
+ }
+
+ if (langOpts.SYCLIsDevice &&
+ (!d || d->getType().getAddressSpace() == LangAS::Default))
+ llvm_unreachable("NYI");
+
+ if (langOpts.CUDA && langOpts.CUDAIsDevice) {
+ if (d) {
+ if (d->hasAttr<CUDAConstantAttr>())
+ return LangAS::cuda_constant;
+ if (d->hasAttr<CUDASharedAttr>())
+ return LangAS::cuda_shared;
+ if (d->hasAttr<CUDADeviceAttr>())
+ return LangAS::cuda_device;
+ if (d->getType().isConstQualified())
+ return LangAS::cuda_constant;
+ }
+ return LangAS::cuda_device;
+ }
+
+ if (langOpts.OpenMP)
+ llvm_unreachable("NYI");
+
+ return getTargetCIRGenInfo().getGlobalVarAddressSpace(*this, d);
+}
+
static void setLinkageForGV(cir::GlobalOp &gv, const NamedDecl *nd) {
// Set linkage and visibility in case we never see a definition.
LinkageInfo lv = nd->getLinkageAndVisibility();
@@ -700,7 +736,6 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty,
}
if (entry) {
- assert(!cir::MissingFeatures::addressSpace());
assert(!cir::MissingFeatures::opGlobalWeakRef());
assert(!cir::MissingFeatures::setDLLStorageClass());
@@ -728,6 +763,9 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty,
return entry;
}
+ mlir::ptr::MemorySpaceAttrInterface declCIRAS =
+ cir::toCIRAddressSpaceAttr(getMLIRContext(), getGlobalVarAddressSpace(d));
+
mlir::Location loc = getLoc(d->getSourceRange());
// Calculate constant storage flag before creating the global. This was moved
@@ -743,9 +781,9 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty,
// mlir::SymbolTable::Visibility::Public is the default, no need to explicitly
// mark it as such.
- cir::GlobalOp gv =
- CIRGenModule::createGlobalOp(*this, loc, mangledName, ty, isConstant,
- /*insertPoint=*/entry.getOperation());
+ cir::GlobalOp gv = CIRGenModule::createGlobalOp(
+ *this, loc, mangledName, ty, isConstant, declCIRAS,
+ /*insertPoint=*/entry.getOperation());
// This is the first use or definition of a mangled name. If there is a
// deferred decl with this name, remember that we need to emit it at the end
@@ -828,7 +866,7 @@ mlir::Value CIRGenModule::getAddrOfGlobalVar(const VarDecl *d, mlir::Type ty,
bool tlsAccess = d->getTLSKind() != VarDecl::TLS_None;
cir::GlobalOp g = getOrCreateCIRGlobal(d, ty, isForDefinition);
- mlir::Type ptrTy = builder.getPointerTo(g.getSymType());
+ mlir::Type ptrTy = builder.getPointerTo(g.getSymType(), g.getAddrSpaceAttr());
return cir::GetGlobalOp::create(
builder, getLoc(d->getSourceRange()), ptrTy, g.getSymNameAttr(),
tlsAccess,
@@ -840,8 +878,8 @@ cir::GlobalViewAttr CIRGenModule::getAddrOfGlobalVarAttr(const VarDecl *d) {
mlir::Type ty = getTypes().convertTypeForMem(d->getType());
cir::GlobalOp globalOp = getOrCreateCIRGlobal(d, ty, NotForDefinition);
- assert(!cir::MissingFeatures::addressSpace());
- cir::PointerType ptrTy = builder.getPointerTo(globalOp.getSymType());
+ cir::PointerType ptrTy =
+ builder.getPointerTo(globalOp.getSymType(), globalOp.getAddrSpaceAttr());
return builder.getGlobalViewAttr(ptrTy, globalOp);
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 52464a8bc30c4..aab8bdaf1c64f 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -26,6 +26,7 @@
#include "clang/CIR/Dialect/IR/CIRDialect.h"
#include "TargetInfo.h"
+#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/MLIRContext.h"
@@ -198,10 +199,11 @@ class CIRGenModule : public CIRGenTypeCache {
cir::GlobalOp getOrCreateCIRGlobal(const VarDecl *d, mlir::Type ty,
ForDefinition_t isForDefinition);
- static cir::GlobalOp createGlobalOp(CIRGenModule &cgm, mlir::Location loc,
- llvm::StringRef name, mlir::Type t,
- bool isConstant = false,
- mlir::Operation *insertPoint = nullptr);
+ static cir::GlobalOp
+ createGlobalOp(CIRGenModule &cgm, mlir::Location loc, llvm::StringRef name,
+ mlir::Type t, bool isConstant = false,
+ mlir::ptr::MemorySpaceAttrInterface addrSpace = {},
+ mlir::Operation *insertPoint = nullptr);
/// Add a global constructor or destructor to the module.
/// The priority is optional, if not specified, the default priority is used.
@@ -754,6 +756,16 @@ class CIRGenModule : public CIRGenTypeCache {
/// Map source language used to a CIR attribute.
std::optional<cir::SourceLanguage> getCIRSourceLanguage() const;
+
+ /// Return the AST address space of the underlying global variable for D, as
+ /// determined by its declaration. Normally this is the same as the address
+ /// space of D's type, but in CUDA, address spaces are associated with
+ /// declarations, not types. If D is nullptr, return the default address
+ /// space for global variable.
+ ///
+ /// For languages without explicit address spaces, if D has default address
+ /// space, target-specific global or constant address space may be returned.
+ LangAS getGlobalVarAddressSpace(const VarDecl *decl);
};
} // namespace CIRGen
diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp
index 2f3824d3d47a7..6570976e0dfeb 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp
@@ -91,3 +91,12 @@ bool TargetCIRGenInfo::isNoProtoCallVariadic(
// For everything else, we just prefer false unless we opt out.
return false;
}
+
+clang::LangAS
+TargetCIRGenInfo::getGlobalVarAddressSpace(CIRGenModule &CGM,
+ const clang::VarDecl *D) const {
+ assert(!CGM.getLangOpts().OpenCL &&
+ !(CGM.getLangOpts().CUDA && CGM.getLangOpts().CUDAIsDevice) &&
+ "Address space agnostic languages only");
+ return D ? D->getType().getAddressSpace() : LangAS::Default;
+}
diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h
index f4792d5309e36..161325c8668e8 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.h
+++ b/clang/lib/CIR/CodeGen/TargetInfo.h
@@ -49,6 +49,13 @@ class TargetCIRGenInfo {
/// Returns ABI info helper for the target.
const ABIInfo &getABIInfo() const { return *info; }
+ /// Get target favored AST address space of a global variable for languages
+ /// other than OpenCL and CUDA.
+ /// If \p D is nullptr, returns the default target favored address space
+ /// for global variable.
+ virtual clang::LangAS getGlobalVarAddressSpace(CIRGenModule &CGM,
+ const clang::VarDecl *D) const;
+
/// Get the address space for alloca.
virtual mlir::ptr::MemorySpaceAttrInterface getCIRAllocaAddressSpace() const {
return cir::LangAddressSpaceAttr::get(&info->cgt.getMLIRContext(),
diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
index d65aabf949f3f..f19db99ed4c79 100644
--- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
@@ -279,6 +279,13 @@ static void printOmittedTerminatorRegion(mlir::OpAsmPrinter &printer,
/*printBlockTerminators=*/!omitRegionTerm(region));
}
+mlir::OptionalParseResult
+parseGlobalAddressSpaceValue(mlir::AsmParser &p,
+ mlir::ptr::MemorySpaceAttrInterface &attr);
+
+void printGlobalAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp op,
+ mlir::ptr::MemorySpaceAttrInterface attr);
+
//===----------------------------------------------------------------------===//
// AllocaOp
//===----------------------------------------------------------------------===//
@@ -1709,7 +1716,9 @@ mlir::LogicalResult cir::GlobalOp::verify() {
void cir::GlobalOp::build(
OpBuilder &odsBuilder, OperationState &odsState, llvm::StringRef sym_name,
- mlir::Type sym_type, bool isConstant, cir::GlobalLinkageKind linkage,
+ mlir::Type sym_type, bool isConstant,
+ mlir::ptr::MemorySpaceAttrInterface addrSpace,
+ cir::GlobalLinkageKind linkage,
function_ref<void(OpBuilder &, Location)> ctorBuilder,
function_ref<void(OpBuilder &, Location)> dtorBuilder) {
odsState.addAttribute(getSymNameAttrName(odsState.name),
@@ -1719,6 +1728,8 @@ void cir::GlobalOp::build(
if (isConstant)
odsState.addAttribute(getConstantAttrName(odsState.name),
odsBuilder.getUnitAttr());
+ if (addrSpace)
+ odsState.addAttribute(getAddrSpaceAttrName(odsState.name), addrSpace);
cir::GlobalLinkageKindAttr linkageAttr =
cir::GlobalLinkageKindAttr::get(odsBuilder.getContext(), linkage);
@@ -1872,9 +1883,10 @@ cir::GetGlobalOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
<< "' does not reference a valid cir.global or cir.func";
mlir::Type symTy;
+ mlir::ptr::MemorySpaceAttrInterface symAddrSpaceAttr{};
if (auto g = dyn_cast<GlobalOp>(op)) {
symTy = g.getSymType();
- assert(!cir::MissingFeatures::addressSpace());
+ symAddrSpaceAttr = g.getAddrSpaceAttr();
// Verify that for thread local global access, the global needs to
// be marked with tls bits.
if (getTls() && !g.getTlsModel())
@@ -1900,6 +1912,13 @@ cir::GetGlobalOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
<< resultType.getPointee() << "' does not match type " << symTy
<< " of the global @" << getName();
+ if (symAddrSpaceAttr != resultType.getAddrSpace()) {
+ return emitOpError()
+ << "result type address space does not match the address "
+ "space of the global @"
+ << getName();
+ }
+
return success();
}
diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp
index 80dce3d3266b5..d96975b3e6aa7 100644
--- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp
@@ -1057,6 +1057,21 @@ void printAddressSpaceValue(mlir::AsmPrinter &p,
llvm_unreachable("unexpected address-space attribute kind");
}
+mlir::OptionalParseResult
+parseGlobalAddressSpaceValue(mlir::AsmParser &p,
+ mlir::ptr::MemorySpaceAttrInterface &attr) {
+
+ mlir::SMLoc loc = p.getCurrentLocation();
+ if (parseAddressSpaceValue(p, attr).failed())
+ return p.emitError(loc, "failed to parse Address Space Value for GlobalOp");
+ return mlir::success();
+}
+
+void printGlobalAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp,
+ mlir::ptr::MemorySpaceAttrInterface attr) {
+ printAddressSpaceValue(printer, attr);
+}
+
mlir::ptr::MemorySpaceAttrInterface cir::normalizeDefaultAddressSpace(
mlir::ptr::MemorySpaceAttrInterface addrSpace) {
if (auto langAS =
diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
index 3dbb78ba8a4cf..bbc8d5e9db409 100644
--- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
@@ -1511,9 +1511,12 @@ void LoweringPreparePass::lowerStoreOfConstAggregate(cir::StoreOp op) {
// constexpr locals as globals when their address is taken), reuse it.
if (!mlir::SymbolTable::lookupSymbolIn(
mlirModule, mlir::StringAttr::get(&getContext(), name))) {
- auto gv = cir::GlobalOp::create(builder, op.getLoc(), name, ty,
- /*isConstant=*/true,
- cir::GlobalLinkageKind::PrivateLinkage);
+ auto gv = cir::GlobalOp::create(
+ builder, op.getLoc(), name, ty,
+ /*isConstant=*/true,
+ cir::LangAddressSpaceAttr::get(&getContext(),
+ cir::LangAddressSpace::Default),
+ cir::GlobalLinkageKind::PrivateLinkage);
mlir::SymbolTable::setSymbolVisibility(
gv, mlir::SymbolTable::Visibility::Private);
gv.setInitialValueAttr(constant);
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index 03085ad29ab78..ca76c34176aad 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -2488,8 +2488,9 @@ void CIRToLLVMGlobalOpLowering::setupRegionInitializedLLVMGlobalOp(
// in CIRToLLVMGlobalOpLowering::matchAndRewrite() but that will go
// away when the placeholders are no longer needed.
const bool isConst = op.getConstant();
- assert(!cir::MissingFeatures::addressSpace());
- const unsigned addrSpace = 0;
+ unsigned addrSpace = 0;
+if(auto targetAS = mlir::dyn_cast_if_present<cir::TargetAddressSpaceAttr>(op.getAddrSpaceAttr()))
+ addrSpace = targetAS.getValue();
const bool isDsoLocal = op.getDsoLocal();
const bool isThreadLocal = (bool)op.getTlsModelAttr();
const uint64_t alignment = op.getAlignment().value_or(0);
@@ -2545,11 +2546,10 @@ mlir::LogicalResult CIRToLLVMGlobalOpLowering::matchAndRewrite(
// This is the LLVM dialect type.
const mlir::Type llvmType =
convertTypeForMemory(*getTypeConverter(), dataLayout, cirSymType);
- // FIXME: These default values are placeholders until the the equivalent
- // attributes are available on cir.global ops.
const bool isConst = op.getConstant();
- assert(!cir::MissingFeatures::addressSpace());
- const unsigned addrSpace = 0;
+ unsigned addrSpace = 0;
+if(auto targetAS = mlir::dyn_cast_if_present<cir::TargetAddressSpaceAttr>(op.getAddrSpaceAttr()))
+ addrSpace = targetAS.getValue();
const bool isDsoLocal = op.getDsoLocal();
const bool isThreadLocal = (bool)op.getTlsModelAttr();
const uint64_t alignment = op.getAlignment().value_or(0);
diff --git a/clang/test/CIR/IR/address-space.cir b/clang/test/CIR/IR/address-space.cir
index 9a729c934bc11..0afe840952046 100644
--- a/clang/test/CIR/IR/address-space.cir
+++ b/clang/test/CIR/IR/address-space.cir
@@ -3,6 +3,8 @@
!s32i = !cir.int<s, 32>
module {
+ // ---- PointerType with address space ----
+
cir.func @target_address_space_ptr(%p: !cir.ptr<!s32i, target_address_space(1)>) {
cir.return
}
@@ -30,6 +32,23 @@ module {
cir.func @default_address_space(%p: !cir.ptr<!s32i>) {
cir.return
}
+
+ // ---- GlobalOp with address space ----
+
+ cir.global external target_address_space(1) @global_target_as = #cir.int<42> : !s32i
+ cir.global "private" internal lang_address_space(offload_local) @global_lang_local : !s32i
+ cir.global external lang_address_space(offload_global) @global_lang_global = #cir.int<1> : !s32i
+ cir.global external lang_address_space(offload_constant) @global_lang_constant = #cir.int<2> : !s32i
+ cir.global external @global_default_as = #cir.int<0> : !s32i
+
+ // ---- GetGlobalOp with address space ----
+
+ cir.func @get_global_with_address_space() {
+ %0 = cir.get_global @global_target_as : !cir.ptr<!s32i, target_address_space(1)>
+ %1 = cir.get_global @global_lang_global : !cir.ptr<!s32i, lang_address_space(offload_global)>
+ %2 = cir.get_global @global_default_as : !cir.ptr<!s32i>
+ cir.return
+ }
}
// CHECK: cir.func @target_address_space_ptr(%arg0: !cir.ptr<!s32i, target_address_space(1)>)
@@ -39,3 +58,14 @@ module {
// CHECK: cir.func @lang_address_space_offload_private(%arg0: !cir.ptr<!s32i, lang_address_space(offload_private)>)
// CHECK: cir.func @lang_address_space_offload_generic(%arg0: !cir.ptr<!s32i, lang_address_space(offload_generic)>)
// CHECK: cir.func @default_address_space(%arg0: !cir.ptr<!s32i>)
+
+// CHECK: cir.global external target_address_space(1) @global_target_as = #cir.int<42> : !s32i
+// CHECK: cir.global "private" internal lang_address_space(offload_local) @global_lang_local : !s32i
+// CHECK: cir.global external lang_address_space(offload_global) @global_lang_global = #cir.int<1> : !s32i
+// CHECK: cir.global external lang_address_space(offload_constant) @global_lang_constant = #cir.int<2> : !s32i
+// CHECK: cir.global external @global_default_as = #cir.int<0> : !s32i
+
+// CHECK: cir.func @get_global_with_address_space()
+// CHECK: cir.get_global @global_target_as : !cir.ptr<!s32i, target_address_space(1)>
+// CHECK: cir.get_global @global_lang_global : !cir.ptr<!s32i, lang_address_space(offload_global)>
+// CHECK: cir.get_global @global_default_as : !cir.ptr<!s32i>
diff --git a/clang/test/CIR/IR/invalid-addrspace.cir b/clang/test/CIR/IR/invalid-addrspace.cir
index d38868f1febf0..882199afd6490 100644
--- a/clang/test/CIR/IR/invalid-addrspace.cir
+++ b/clang/test/CIR/IR/invalid-addrspace.cir
@@ -50,3 +50,23 @@ cir.func @lang_address_space_empty(%p : !cir.ptr<!u64i, lang_address_space()>) {
cir.func @lang_address_space_invalid(%p : !cir.ptr<!u64i, lang_address_space(foobar)>) {
cir.return
}
+
+// -----
+
+!s32i = !cir.int<s, 32>
+cir.global external target_address_space(1) @global_in_as1 = #cir.int<42> : !s32i
+cir.func @get_global_mismatched_address_space() {
+ // expected-error at +1 {{result type address space does not match the address space of the global @global_in_as1}}
+ %0 = cir.get_global @global_in_as1 : !cir.ptr<!s32i>
+ cir.return
+}
+
+// -----
+
+!s32i = !cir.int<s, 32>
+cir.global external @global_default_as = #cir.int<0> : !s32i
+cir.func @get_global_unexpected_address_space() {
+ // expected-error at +1 {{result type address space does not match the address space of the global @global_default_as}}
+ %0 = cir.get_global @global_default_as : !cir.ptr<!s32i, target_address_space(1)>
+ cir.return
+}
diff --git a/clang/test/CIR/Lowering/global-address-space.cir b/clang/test/CIR/Lowering/global-address-space.cir
new file mode 100644
index 0000000000000..c9f25e1126098
--- /dev/null
+++ b/clang/test/CIR/Lowering/global-address-space.cir
@@ -0,0 +1,46 @@
+// RUN: cir-opt %s -cir-to-llvm -o %t.mlir
+// RUN: FileCheck --input-file=%t.mlir %s
+
+!s32i = !cir.int<s, 32>
+
+module {
+ cir.global external target_address_space(1) @global_as1 = #cir.int<42> : !s32i
+ // CHECK: llvm.mlir.global external @global_as1(42 : i32) {addr_space = 1 : i32} : i32
+
+ cir.global external target_address_space(3) @global_as3 = #cir.int<100> : !s32i
+ // CHECK: llvm.mlir.global external @global_as3(100 : i32) {addr_space = 3 : i32} : i32
+
+ cir.global external @global_default = #cir.int<0> : !s32i
+ // CHECK: llvm.mlir.global external @global_default(0 : i32) {addr_space = 0 : i32} : i32
+
+ // Test cir.get_global with address space produces correct llvm.mlir.addressof type
+ // CHECK-LABEL: llvm.func @test_get_global_as1
+ cir.func @test_get_global_as1() -> !s32i {
+ // CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @global_as1 : !llvm.ptr<1>
+ // CHECK: %[[VAL:.*]] = llvm.load %[[ADDR]] {{.*}} : !llvm.ptr<1> -> i32
+ // CHECK: llvm.return %[[VAL]] : i32
+ %0 = cir.get_global @global_as1 : !cir.ptr<!s32i, target_address_space(1)>
+ %1 = cir.load %0 : !cir.ptr<!s32i, target_address_space(1)>, !s32i
+ cir.return %1 : !s32i
+ }
+
+ // CHECK-LABEL: llvm.func @test_get_global_as3
+ cir.func @test_get_global_as3() -> !s32i {
+ // CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @global_as3 : !llvm.ptr<3>
+ // CHECK: %[[VAL:.*]] = llvm.load %[[ADDR]] {{.*}} : !llvm.ptr<3> -> i32
+ // CHECK: llvm.return %[[VAL]] : i32
+ %0 = cir.get_global @global_as3 : !cir.ptr<!s32i, target_address_space(3)>
+ %1 = cir.load %0 : !cir.ptr<!s32i, target_address_space(3)>, !s32i
+ cir.return %1 : !s32i
+ }
+
+ // CHECK-LABEL: llvm.func @test_get_global_default
+ cir.func @test_get_global_default() -> !s32i {
+ // CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @global_default : !llvm.ptr
+ // CHECK: %[[VAL:.*]] = llvm.load %[[ADDR]] {{.*}} : !llvm.ptr -> i32
+ // CHECK: llvm.return %[[VAL]] : i32
+ %0 = cir.get_global @global_default : !cir.ptr<!s32i>
+ %1 = cir.load %0 : !cir.ptr<!s32i>, !s32i
+ cir.return %1 : !s32i
+ }
+}
>From ad6c3fda5b88a4b4f57a1f0744cb4b607b577bbb Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Thu, 26 Feb 2026 17:21:27 -0500
Subject: [PATCH 2/5] Global AS lowering For CUDA and CIRGen tests for target
AS
---
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 28 ++++++++++++++----
clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 2 ++
clang/test/CIR/CodeGen/address-space.c | 17 +++++++++++
clang/test/CIR/CodeGenCUDA/address-spaces.cu | 30 ++++++++++++++++++++
4 files changed, 72 insertions(+), 5 deletions(-)
create mode 100644 clang/test/CIR/CodeGenCUDA/address-spaces.cu
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 23b91ddc6ece0..2afbd32c7733b 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -763,9 +763,6 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty,
return entry;
}
- mlir::ptr::MemorySpaceAttrInterface declCIRAS =
- cir::toCIRAddressSpaceAttr(getMLIRContext(), getGlobalVarAddressSpace(d));
-
mlir::Location loc = getLoc(d->getSourceRange());
// Calculate constant storage flag before creating the global. This was moved
@@ -779,6 +776,9 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty,
astContext, /*ExcludeCtor=*/true, /*ExcludeDtor=*/!needsDtor);
}
+ mlir::ptr::MemorySpaceAttrInterface declCIRAS =
+ cir::toCIRAddressSpaceAttr(getMLIRContext(), getGlobalVarAddressSpace(d));
+
// mlir::SymbolTable::Visibility::Public is the default, no need to explicitly
// mark it as such.
cir::GlobalOp gv = CIRGenModule::createGlobalOp(
@@ -985,8 +985,26 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd,
errorNYI(vd->getSourceRange(), "annotate global variable");
}
- if (langOpts.CUDA) {
- errorNYI(vd->getSourceRange(), "CUDA global variable");
+ // CUDA E.2.4.1 "__shared__ variables cannot have an initialization
+ // as part of their declaration." Sema has already checked for
+ // error cases, so we just need to set Init to UndefValue.
+ bool isCUDASharedVar =
+ getLangOpts().CUDAIsDevice && vd->hasAttr<CUDASharedAttr>();
+ // Shadows of initialized device-side global variables are also left
+ // undefined.
+ // Managed Variables should be initialized on both host side and device side.
+ bool isCUDAShadowVar =
+ !getLangOpts().CUDAIsDevice && !vd->hasAttr<HIPManagedAttr>() &&
+ (vd->hasAttr<CUDAConstantAttr>() || vd->hasAttr<CUDADeviceAttr>() ||
+ vd->hasAttr<CUDASharedAttr>());
+ bool isCUDADeviceShadowVar =
+ getLangOpts().CUDAIsDevice && !vd->hasAttr<HIPManagedAttr>() &&
+ (vd->getType()->isCUDADeviceBuiltinSurfaceType() ||
+ vd->getType()->isCUDADeviceBuiltinTextureType());
+
+ if (getLangOpts().CUDA &&
+ (isCUDASharedVar || isCUDAShadowVar || isCUDADeviceShadowVar)) {
+ init = cir::PoisonAttr::get(convertType(vd->getType()));
}
// Set initializer and finalize emission
diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
index f19db99ed4c79..e16397b1a708f 100644
--- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
@@ -1728,6 +1728,8 @@ void cir::GlobalOp::build(
if (isConstant)
odsState.addAttribute(getConstantAttrName(odsState.name),
odsBuilder.getUnitAttr());
+
+ addrSpace = skipDefaultAddressSpace(addrSpace);
if (addrSpace)
odsState.addAttribute(getAddrSpaceAttrName(odsState.name), addrSpace);
diff --git a/clang/test/CIR/CodeGen/address-space.c b/clang/test/CIR/CodeGen/address-space.c
index 77404c9eab1ca..491d0218f288b 100644
--- a/clang/test/CIR/CodeGen/address-space.c
+++ b/clang/test/CIR/CodeGen/address-space.c
@@ -5,6 +5,12 @@
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o %t.ll
// RUN: FileCheck --input-file=%t.ll %s -check-prefix=OGCG
+// Test global variable with address space
+// CIR: cir.global external @gvar = #cir.ptr<null> : !cir.ptr<!s32i, target_address_space(1)>
+// LLVM: @gvar = global ptr addrspace(1) null
+// OGCG: @gvar = global ptr addrspace(1) null
+int __attribute__((address_space(1))) *gvar;
+
// Test address space 1
// CIR: cir.func {{.*}} @foo(%arg0: !cir.ptr<!s32i, target_address_space(1)>
// LLVM: define dso_local void @foo(ptr addrspace(1) noundef %0)
@@ -28,3 +34,14 @@ void bar(int __attribute__((address_space(0))) *arg) {
void baz(int *arg) {
return;
}
+
+// End to end function returning pointer to address space global
+// CIR: cir.func {{.*}} @get_gvar()
+// CIR: cir.get_global @gvar : !cir.ptr<!cir.ptr<!s32i, target_address_space(1)>>
+// LLVM: define dso_local ptr addrspace(1) @get_gvar()
+// LLVM: load ptr addrspace(1), ptr @gvar
+// OGCG: define dso_local ptr addrspace(1) @get_gvar()
+// OGCG: load ptr addrspace(1), ptr @gvar
+int __attribute__((address_space(1)))* get_gvar() {
+ return gvar;
+}
diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
new file mode 100644
index 0000000000000..8f95c1600878d
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu
@@ -0,0 +1,30 @@
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda \
+// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s
+
+// Verifies CIR emits correct address spaces for CUDA globals.
+
+#include "Inputs/cuda.h"
+
+// CHECK: cir.global external lang_address_space(offload_global) @i = #cir.int<0> : !s32i
+__device__ int i;
+
+// CHECK: cir.global constant external lang_address_space(offload_constant) @j = #cir.int<0> : !s32i
+__constant__ int j;
+
+// CHECK: cir.global external lang_address_space(offload_local) @k = #cir.poison : !s32i
+__shared__ int k;
+
+// CHECK: cir.global external lang_address_space(offload_local) @b = #cir.poison : !cir.float
+__shared__ float b;
+
+__device__ void foo() {
+ // CHECK: cir.get_global @i : !cir.ptr<!s32i, lang_address_space(offload_global)>
+ i++;
+
+ // CHECK: cir.get_global @j : !cir.ptr<!s32i, lang_address_space(offload_constant)>
+ j++;
+
+ // CHECK: cir.get_global @k : !cir.ptr<!s32i, lang_address_space(offload_local)>
+ k++;
+}
>From da9ba20bb4886716cbb8e2b2c7b874b9747a165c Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Thu, 26 Feb 2026 17:31:26 -0500
Subject: [PATCH 3/5] fix fmt
---
clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 13 +++++++++----
1 file changed, 9 insertions(+), 4 deletions(-)
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index ca76c34176aad..4319e1e4b1750 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -2489,8 +2489,9 @@ void CIRToLLVMGlobalOpLowering::setupRegionInitializedLLVMGlobalOp(
// away when the placeholders are no longer needed.
const bool isConst = op.getConstant();
unsigned addrSpace = 0;
-if(auto targetAS = mlir::dyn_cast_if_present<cir::TargetAddressSpaceAttr>(op.getAddrSpaceAttr()))
- addrSpace = targetAS.getValue();
+ if (auto targetAS = mlir::dyn_cast_if_present<cir::TargetAddressSpaceAttr>(
+ op.getAddrSpaceAttr()))
+ addrSpace = targetAS.getValue();
const bool isDsoLocal = op.getDsoLocal();
const bool isThreadLocal = (bool)op.getTlsModelAttr();
const uint64_t alignment = op.getAlignment().value_or(0);
@@ -2546,10 +2547,14 @@ mlir::LogicalResult CIRToLLVMGlobalOpLowering::matchAndRewrite(
// This is the LLVM dialect type.
const mlir::Type llvmType =
convertTypeForMemory(*getTypeConverter(), dataLayout, cirSymType);
+
+ // FIXME: These default values are placeholders until the the equivalent
+ // attributes are available on cir.global ops.
const bool isConst = op.getConstant();
unsigned addrSpace = 0;
-if(auto targetAS = mlir::dyn_cast_if_present<cir::TargetAddressSpaceAttr>(op.getAddrSpaceAttr()))
- addrSpace = targetAS.getValue();
+ if (auto targetAS = mlir::dyn_cast_if_present<cir::TargetAddressSpaceAttr>(
+ op.getAddrSpaceAttr()))
+ addrSpace = targetAS.getValue();
const bool isDsoLocal = op.getDsoLocal();
const bool isThreadLocal = (bool)op.getTlsModelAttr();
const uint64_t alignment = op.getAlignment().value_or(0);
>From b89c142ccc2fad1e6c9d0002ac2924d98f289290 Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Thu, 26 Feb 2026 17:37:11 -0500
Subject: [PATCH 4/5] more fmt yo
---
clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index 4319e1e4b1750..b5a181c198993 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -2549,7 +2549,7 @@ mlir::LogicalResult CIRToLLVMGlobalOpLowering::matchAndRewrite(
convertTypeForMemory(*getTypeConverter(), dataLayout, cirSymType);
// FIXME: These default values are placeholders until the the equivalent
- // attributes are available on cir.global ops.
+ // attributes are available on cir.global ops.
const bool isConst = op.getConstant();
unsigned addrSpace = 0;
if (auto targetAS = mlir::dyn_cast_if_present<cir::TargetAddressSpaceAttr>(
>From d54d5c13c58ebe6945f132afe5136cca415f92a4 Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Mon, 2 Mar 2026 17:50:41 -0500
Subject: [PATCH 5/5] fix comp err
---
clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
index e16397b1a708f..8da81b3570044 100644
--- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
@@ -1729,7 +1729,7 @@ void cir::GlobalOp::build(
odsState.addAttribute(getConstantAttrName(odsState.name),
odsBuilder.getUnitAttr());
- addrSpace = skipDefaultAddressSpace(addrSpace);
+ addrSpace = normalizeDefaultAddressSpace(addrSpace);
if (addrSpace)
odsState.addAttribute(getAddrSpaceAttrName(odsState.name), addrSpace);
More information about the cfe-commits
mailing list