From 0ec6e1d21e83bedc8ca6a8f0b1c969197f1d1664 Mon Sep 17 00:00:00 2001 From: David Rivera Date: Fri, 20 Mar 2026 14:01:20 -0400 Subject: [PATCH] [CIR] Address Space support for GlobalOps (#179082) Related: https://github.com/llvm/llvm-project/issues/179278, https://github.com/llvm/llvm-project/issues/160386 Extends cir.global to accept address space attributes. Globals can now specify either `target_address_space(N)` or `lang_address_space(offload_*)`. Address spaces are also preserved throughout get_global ops. --- .../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 | 89 ++++++++++++++----- clang/lib/CIR/CodeGen/CIRGenModule.h | 20 ++++- clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 25 +++++- clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 15 ++++ .../Dialect/Transforms/LoweringPrepare.cpp | 9 +- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 13 ++- clang/test/CIR/CodeGen/address-space.c | 17 ++++ clang/test/CIR/CodeGenCUDA/address-spaces.cu | 9 +- 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, 270 insertions(+), 51 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 95785f33659f..a6ddfd5ca4d4 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" @@ -394,14 +395,15 @@ public: 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 c3b50ff0e731..41858a61480a 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -2834,6 +2834,7 @@ def CIR_GlobalOp : CIR_Op<"global", [ OptionalAttr:$sym_visibility, TypeAttr:$sym_type, CIR_GlobalLinkageKind:$linkage, + OptionalAttr:$addr_space, OptionalAttr:$tls_model, OptionalAttr:$initial_value, UnitAttr:$comdat, @@ -2855,6 +2856,7 @@ def CIR_GlobalOp : CIR_Op<"global", [ ($tls_model^)? (`dso_local` $dso_local^)? (`static_local_guard` `` $static_local_guard^)? + (` ` custom($addr_space)^ )? $sym_name custom($sym_type, $initial_value, $ctorRegion, $dtorRegion) @@ -2875,6 +2877,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 cbd1b83f49dc..6cf27126e3bc 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/Builders.h" #include "mlir/IR/BuiltinAttributes.h" @@ -691,7 +692,8 @@ public: [[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()]++) @@ -699,7 +701,8 @@ public: 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 5328bb0a812a..a204f63ca736 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -298,9 +298,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(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 54949159f27d..4c238779342a 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -27,12 +27,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" @@ -608,10 +610,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(); @@ -631,7 +634,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; @@ -680,6 +683,39 @@ std::optional 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)) + errorNYI("SYCL global address space"); + + if (langOpts.CUDA && langOpts.CUDAIsDevice) { + if (d) { + if (d->hasAttr()) + return LangAS::cuda_constant; + if (d->hasAttr()) + return LangAS::cuda_shared; + if (d->hasAttr()) + return LangAS::cuda_device; + if (d->getType().isConstQualified()) + return LangAS::cuda_constant; + } + return LangAS::cuda_device; + } + + if (langOpts.OpenMP) + errorNYI("OpenMP global address space"); + + 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(); @@ -857,13 +893,13 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty, } if (entry) { - assert(!cir::MissingFeatures::addressSpace()); assert(!cir::MissingFeatures::opGlobalWeakRef()); assert(!cir::MissingFeatures::setDLLStorageClass()); assert(!cir::MissingFeatures::openMP()); - if (entry.getSymType() == ty) + if (entry.getSymType() == ty && + (cir::isMatchingAddressSpace(entry.getAddrSpaceAttr(), langAS))) return entry; // If there are two attempts to define the same mangled name, issue an @@ -899,11 +935,14 @@ 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(*this, loc, mangledName, ty, isConstant, - /*insertPoint=*/entry.getOperation()); + cir::GlobalOp gv = CIRGenModule::createGlobalOp( + *this, loc, mangledName, ty, isConstant, declCIRAS, + /*insertPoint=*/entry.getOperation()); // If we already created a global with the same mangled name (but different // type) before, remove it from its parent. @@ -989,7 +1028,7 @@ CIRGenModule::getOrCreateCIRGlobal(const VarDecl *d, mlir::Type ty, ty = getTypes().convertTypeForMem(astTy); StringRef mangledName = getMangledName(d); - return getOrCreateCIRGlobal(mangledName, ty, astTy.getAddressSpace(), d, + return getOrCreateCIRGlobal(mangledName, ty, getGlobalVarAddressSpace(d), d, isForDefinition); } @@ -1008,7 +1047,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, @@ -1020,8 +1059,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); } @@ -1059,18 +1098,26 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd, std::optional emitter; - assert(!cir::MissingFeatures::cudaSupport()); - // 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. + // error cases, so we just need to set Init to PoisonValue. bool isCUDASharedVar = getLangOpts().CUDAIsDevice && vd->hasAttr(); - // TODO(cir): implement isCUDAShadowVar and isCUDADeviceShadowVar, reference: - // OGCG + // 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() && + (vd->hasAttr() || vd->hasAttr() || + vd->hasAttr()); + bool isCUDADeviceShadowVar = + getLangOpts().CUDAIsDevice && !vd->hasAttr() && + (vd->getType()->isCUDADeviceBuiltinSurfaceType() || + vd->getType()->isCUDADeviceBuiltinTextureType()); - if (getLangOpts().CUDA && isCUDASharedVar) { - init = cir::UndefAttr::get(&getMLIRContext(), convertType(vd->getType())); + if (getLangOpts().CUDA && + (isCUDASharedVar || isCUDAShadowVar || isCUDADeviceShadowVar)) { + init = cir::PoisonAttr::get(convertType(vd->getType())); } else if (vd->hasAttr()) { errorNYI(vd->getSourceRange(), "emitGlobalVarDefinition: loader uninitialized attribute"); diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index d8444fda7b40..e1bea451c859 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" @@ -206,10 +207,11 @@ public: 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. @@ -803,6 +805,16 @@ private: /// Map source language used to a CIR attribute. std::optional 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/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index 01a35534f9bf..bf369bfe6999 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -284,6 +284,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 //===----------------------------------------------------------------------===// @@ -1749,7 +1756,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 ctorBuilder, function_ref dtorBuilder) { odsState.addAttribute(getSymNameAttrName(odsState.name), @@ -1760,6 +1769,10 @@ void cir::GlobalOp::build( odsState.addAttribute(getConstantAttrName(odsState.name), odsBuilder.getUnitAttr()); + addrSpace = normalizeDefaultAddressSpace(addrSpace); + if (addrSpace) + odsState.addAttribute(getAddrSpaceAttrName(odsState.name), addrSpace); + cir::GlobalLinkageKindAttr linkageAttr = cir::GlobalLinkageKindAttr::get(odsBuilder.getContext(), linkage); odsState.addAttribute(getLinkageAttrName(odsState.name), linkageAttr); @@ -1912,9 +1925,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(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()) @@ -1940,6 +1954,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 80dce3d3266b..d96975b3e6aa 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 3a196a44d50b..756a45300165 100644 --- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp +++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp @@ -1564,9 +1564,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 6463c662ade6..9fa0e720e159 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -2569,8 +2569,10 @@ 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( + 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); @@ -2626,11 +2628,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(); - assert(!cir::MissingFeatures::addressSpace()); - const unsigned addrSpace = 0; + unsigned addrSpace = 0; + if (auto targetAS = mlir::dyn_cast_if_present( + 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/CodeGen/address-space.c b/clang/test/CIR/CodeGen/address-space.c index 77404c9eab1c..491d0218f288 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 : !cir.ptr +// 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 // 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> +// 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 index 166da94fa905..1ed52378b99a 100644 --- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu +++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu @@ -19,17 +19,12 @@ // LLVM-DEVICE: @_ZZ2fnvE1j = internal global i32 undef, align 4 __device__ int a; -// CIR-DEVICE: cir.global external @[[DEV:.*]] = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} +// CIR-DEVICE: cir.global external lang_address_space(offload_global) @[[DEV:.*]] = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} // LLVM-DEVICE: @[[DEV_LD:.*]] = externally_initialized global i32 0, align 4 // OGCG-DEVICE: @[[DEV_OD:.*]] = addrspace(1) externally_initialized global i32 0, align 4 -__shared__ int b; -// CIR-DEVICE: cir.global external @[[SHARED:.*]] = #cir.undef : !s32i {alignment = 4 : i64} -// LLVM-DEVICE: @[[SHARED_LL:.*]] = global i32 undef, align 4 -// OGCG-DEVICE: @[[SHARED_OD:.*]] = addrspace(3) global i32 undef, align 4 - __constant__ int c; -// CIR-DEVICE: cir.global constant external @[[CONST:.*]] = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} +// CIR-DEVICE: cir.global constant external lang_address_space(offload_constant) @[[CONST:.*]] = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} // LLVM-DEVICE: @[[CONST_LL:.*]] = externally_initialized constant i32 0, align 4 // OGCG-DEVICE: @[[CONST_OD:.*]] = addrspace(4) externally_initialized constant i32 0, align 4 diff --git a/clang/test/CIR/IR/address-space.cir b/clang/test/CIR/IR/address-space.cir index 9a729c934bc1..0afe84095204 100644 --- a/clang/test/CIR/IR/address-space.cir +++ b/clang/test/CIR/IR/address-space.cir @@ -3,6 +3,8 @@ !s32i = !cir.int module { + // ---- PointerType with address space ---- + cir.func @target_address_space_ptr(%p: !cir.ptr) { cir.return } @@ -30,6 +32,23 @@ module { cir.func @default_address_space(%p: !cir.ptr) { 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 + %1 = cir.get_global @global_lang_global : !cir.ptr + %2 = cir.get_global @global_default_as : !cir.ptr + cir.return + } } // CHECK: cir.func @target_address_space_ptr(%arg0: !cir.ptr) @@ -39,3 +58,14 @@ module { // CHECK: cir.func @lang_address_space_offload_private(%arg0: !cir.ptr) // CHECK: cir.func @lang_address_space_offload_generic(%arg0: !cir.ptr) // CHECK: cir.func @default_address_space(%arg0: !cir.ptr) + +// 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 +// CHECK: cir.get_global @global_lang_global : !cir.ptr +// CHECK: cir.get_global @global_default_as : !cir.ptr diff --git a/clang/test/CIR/IR/invalid-addrspace.cir b/clang/test/CIR/IR/invalid-addrspace.cir index d38868f1febf..882199afd649 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) { cir.func @lang_address_space_invalid(%p : !cir.ptr) { cir.return } + +// ----- + +!s32i = !cir.int +cir.global external target_address_space(1) @global_in_as1 = #cir.int<42> : !s32i +cir.func @get_global_mismatched_address_space() { + // expected-error@+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 + cir.return +} + +// ----- + +!s32i = !cir.int +cir.global external @global_default_as = #cir.int<0> : !s32i +cir.func @get_global_unexpected_address_space() { + // expected-error@+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 + 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 000000000000..c9f25e112609 --- /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 + +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 + %1 = cir.load %0 : !cir.ptr, !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 + %1 = cir.load %0 : !cir.ptr, !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 + %1 = cir.load %0 : !cir.ptr, !s32i + cir.return %1 : !s32i + } +}