shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

[CIR] Address Space support for GlobalOps (#179082)

Browse Source 
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.
This commit is contained in:
David Rivera 2026-03-20 14:01:20 -04:00 committed by GitHub
parent 4a5da64759
commit 0ec6e1d21e
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
15 changed files with 270 additions and 51 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
View File

@ -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,

3
clang/include/clang/CIR/Dialect/IR/CIROps.td
View File

@ -2834,6 +2834,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,
@ -2855,6 +2856,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)
@ -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,

7
clang/lib/CIR/CodeGen/CIRGenBuilder.h
View File

@ -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) {

4
clang/lib/CIR/CodeGen/CIRGenExpr.cpp
View File

@ -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<cir::PointerType>(v.getType()).getAddrSpace());
if (realPtrTy != v.getType())
v = cgf.getBuilder().createBitcast(v.getLoc(), v, realPtrTy);

89
clang/lib/CIR/CodeGen/CIRGenModule.cpp
View File

@ -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<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))
errorNYI("SYCL global address space");
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)
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<ConstantEmitter> 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<CUDASharedAttr>();
// 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<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) {
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<LoaderUninitializedAttr>()) {
errorNYI(vd->getSourceRange(),
"emitGlobalVarDefinition: loader uninitialized attribute");

20
clang/lib/CIR/CodeGen/CIRGenModule.h
View File

@ -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<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

25
clang/lib/CIR/Dialect/IR/CIRDialect.cpp
View File

@ -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<void(OpBuilder &, Location)> ctorBuilder,
function_ref<void(OpBuilder &, Location)> 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<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())
@ -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();
}

15
clang/lib/CIR/Dialect/IR/CIRTypes.cpp
View File

@ -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 =

9
clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
View File

@ -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);

13
clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
View File

@ -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<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);
@ -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<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);

17
clang/test/CIR/CodeGen/address-space.c
View File

@ -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;
}

9
clang/test/CIR/CodeGenCUDA/address-spaces.cu
View File

@ -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

30
clang/test/CIR/IR/address-space.cir
View File

@ -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>

20
clang/test/CIR/IR/invalid-addrspace.cir
View File

@ -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@+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@+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
}

46
clang/test/CIR/Lowering/global-address-space.cir Normal file
View File

@ -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
}
}