From 229b5c54c2d7acfef95a373b8f1a0bb1451c3e04 Mon Sep 17 00:00:00 2001 From: David Rivera Date: Sat, 15 Nov 2025 14:25:30 -0500 Subject: [PATCH 01/13] [CIR][AddrSpace] Backport TargetAddressSpaceAttr and Support both language(clang) and target address-space attributes in pointer types --- .../CIR/Dialect/Builder/CIRBaseBuilder.h | 40 ++++- .../CIR/Dialect/IR/CIRAttrConstraints.td | 19 +++ .../include/clang/CIR/Dialect/IR/CIRAttrs.td | 31 ++++ clang/include/clang/CIR/Dialect/IR/CIRTypes.h | 9 ++ .../include/clang/CIR/Dialect/IR/CIRTypes.td | 29 +++- clang/lib/CIR/Dialect/IR/CIRAttrs.cpp | 14 +- clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 15 +- clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 144 ++++++++++++++---- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 60 +++++--- 9 files changed, 300 insertions(+), 61 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h index a533bea6f1e6..e24684cdccca 100644 --- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h +++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h @@ -11,6 +11,7 @@ #include "clang/AST/Decl.h" #include "clang/AST/Type.h" +#include "clang/Basic/AddressSpaces.h" #include "clang/CIR/Dialect/IR/CIRAttrs.h" #include "clang/CIR/Dialect/IR/CIRDialect.h" #include "clang/CIR/Dialect/IR/CIROpsEnums.h" @@ -105,20 +106,49 @@ class CIRBaseBuilderTy : public mlir::OpBuilder { return cir::PointerType::get(ty); } - cir::PointerType getPointerTo(mlir::Type ty, cir::AddressSpace as) { - return cir::PointerType::get(ty, as); + /// Create a pointer type with an address space attribute. + cir::PointerType getPointerTo(mlir::Type ty, mlir::Attribute addrSpaceAttr) { + return cir::PointerType::get(ty, addrSpaceAttr); } + /// Create a pointer type from a cir::AddressSpace enum. + /// This converts the enum to the appropriate attribute. + cir::PointerType getPointerTo(mlir::Type ty, cir::AddressSpace addrSpace) { + if (addrSpace == cir::AddressSpace::Default) + return getPointerTo(ty); + if (cir::isTargetAddressSpace(addrSpace)) { + unsigned targetAS = cir::getTargetAddressSpaceValue(addrSpace); + auto attr = cir::TargetAddressSpaceAttr::get(getContext(), targetAS); + return getPointerTo(ty, attr); + } + auto attr = cir::AddressSpaceAttr::get(getContext(), addrSpace); + return getPointerTo(ty, attr); + } + + /// Create a pointer type from a LangAS. + /// This converts the LangAS to the appropriate attribute (AddressSpaceAttr or TargetAddressSpaceAttr). cir::PointerType getPointerTo(mlir::Type ty, clang::LangAS langAS) { - return getPointerTo(ty, cir::toCIRAddressSpace(langAS)); + if(langAS == clang::LangAS::Default) + return getPointerTo(ty); + + mlir::Attribute addrSpaceAttr = cir::toCIRAddressSpaceAttr(getContext(), langAS); + return getPointerTo(ty, addrSpaceAttr); + } + + /// Create a pointer type with a target-specific address space value. + /// This is used for address spaces specified via __attribute__((address_space(N))). + cir::PointerType getPointerToWithTargetAddrSpace(mlir::Type ty, unsigned targetAS) { + assert(clang::isTargetAddressSpace(static_cast(targetAS))); + auto attr = cir::TargetAddressSpaceAttr::get(getContext(), targetAS); + return getPointerTo(ty, attr); } cir::PointerType getVoidPtrTy(clang::LangAS langAS = clang::LangAS::Default) { return getPointerTo(cir::VoidType::get(getContext()), langAS); } - cir::PointerType getVoidPtrTy(cir::AddressSpace as) { - return getPointerTo(cir::VoidType::get(getContext()), as); + cir::PointerType getVoidPtrTy(mlir::Attribute addrSpaceAttr) { + return getPointerTo(cir::VoidType::get(getContext()), addrSpaceAttr); } cir::MethodAttr getMethodAttr(cir::MethodType ty, cir::FuncOp methodFuncOp) { diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td index 5023d143da01..0118102ec607 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td @@ -89,4 +89,23 @@ def CIR_AnyTBAAAttr : AnyAttrOf<[ string cppType = "::mlir::Attribute"; } +//===----------------------------------------------------------------------===// +// AddressSpaceAttr constraints +//===----------------------------------------------------------------------===// + +// NOTE: We might end up using this only for GlobalOps, as we cannot apply constraints +// to types. +def CIR_AddressSpaceAttrConstraint + : CIR_AttrConstraint<"::cir::AddressSpaceAttr", "language address space attribute">; + +def CIR_TargetAddressSpaceAttrConstraint + : CIR_AttrConstraint<"::cir::TargetAddressSpaceAttr", "target address space attribute">; + +def CIR_AnyAddressSpaceAttr : AnyAttrOf<[ + CIR_AddressSpaceAttrConstraint, + CIR_TargetAddressSpaceAttrConstraint +]> { + string cppType = "::mlir::Attribute"; +} + #endif // CLANG_CIR_DIALECT_IR_CIRATTRCONSTRAINTS_TD diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td index c3281e30506c..f35cb27bdc89 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -1002,6 +1002,37 @@ def CIR_AddressSpaceAttr : CIR_EnumAttr { }]; } +//===----------------------------------------------------------------------===// +// TargetAddressSpaceAttr +//===----------------------------------------------------------------------===// + +def CIR_TargetAddressSpaceAttr : CIR_Attr<"TargetAddressSpace", "target_address_space"> { + let summary = "Target-specific numeric address space attribute"; + let description = [{ + The TargetAddressSpaceAttr represents a target-specific numeric address space, + corresponding to the LLVM IR `addressspace` qualifier and the clang + `address_space` attribute. + + A value of zero represents the default address space. The semantics of non-zero + address spaces are target-specific. + + Unlike `AddressSpaceAttr` which represents language-specific address spaces + (like OpenCL/CUDA address spaces), this attribute directly represents a + target-specific numeric address space value. + + Example: + ```mlir + !cir.ptr + ``` + }]; + + let parameters = (ins "unsigned":$value); + + let assemblyFormat = [{ + `<` $value `>` + }]; +} + //===----------------------------------------------------------------------===// // AST Wrappers //===----------------------------------------------------------------------===// diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h index b7cf95d08ca8..b5628d29cd17 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h @@ -36,6 +36,15 @@ bool isSized(mlir::Type ty); cir::AddressSpace toCIRAddressSpace(clang::LangAS langAS); +/// Convert a LangAS to the appropriate address space attribute. +/// Returns AddressSpaceAttr for language-specific address spaces, +/// or TargetAddressSpaceAttr for target-specific address spaces. +mlir::Attribute toCIRAddressSpaceAttr(mlir::MLIRContext *ctx, clang::LangAS langAS); + +/// Extract the AddressSpace enum from an address space attribute. +/// Returns Default if the attribute is null. +cir::AddressSpace getCIRAddressSpaceFromAttr(mlir::Attribute attr); + constexpr unsigned getAsUnsignedValue(cir::AddressSpace as) { return static_cast(as); } diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td index 33abc3ffc66f..bd078ab08222 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td @@ -16,11 +16,13 @@ include "clang/CIR/Dialect/IR/CIRDialect.td" include "clang/CIR/Dialect/IR/CIREnumAttr.td" include "clang/CIR/Dialect/IR/CIRTypeConstraints.td" +include "clang/CIR/Dialect/IR/CIRAttrConstraints.td" include "clang/CIR/Interfaces/ASTAttrInterfaces.td" include "clang/CIR/Interfaces/CIRTypeInterfaces.td" include "mlir/Interfaces/DataLayoutInterfaces.td" include "mlir/IR/AttrTypeBase.td" include "mlir/IR/EnumAttr.td" +include "mlir/IR/CommonAttrConstraints.td" //===----------------------------------------------------------------------===// // CIR Types @@ -231,24 +233,23 @@ def CIR_PointerType : CIR_Type<"Pointer", "ptr", [ ``` }]; + let genVerifyDecl = 1; + let parameters = (ins "mlir::Type":$pointee, - CIR_DefaultValuedEnumParameter< - CIR_AddressSpace, - "cir::AddressSpace::Default" - >:$addrSpace + OptionalParameter<"mlir::Attribute">:$addrSpace ); let skipDefaultBuilders = 1; let builders = [ TypeBuilderWithInferredContext<(ins "mlir::Type":$pointee, - CArg<"cir::AddressSpace", "cir::AddressSpace::Default">:$addrSpace), [{ + CArg<"mlir::Attribute", "nullptr">:$addrSpace), [{ return $_get(pointee.getContext(), pointee, addrSpace); }]>, TypeBuilder<(ins "mlir::Type":$pointee, - CArg<"cir::AddressSpace", "cir::AddressSpace::Default">:$addrSpace), [{ + CArg<"mlir::Attribute", "nullptr">:$addrSpace), [{ return $_get($_ctxt, pointee, addrSpace); }]> ]; @@ -256,7 +257,7 @@ def CIR_PointerType : CIR_Type<"Pointer", "ptr", [ let assemblyFormat = [{ `<` $pointee - ( `,` `addrspace` `(` `` custom($addrSpace)^ `)` )? + ( `,` ` ` custom($addrSpace)^ )? `>` }]; @@ -286,6 +287,20 @@ def CIR_PointerType : CIR_Type<"Pointer", "ptr", [ return ptrType.isPtrTo(type); return false; } + + /// Returns true if this pointer type uses a target address space. + bool hasTargetAddressSpace() const; + + /// Returns the target address space value if this is a target address space, + /// otherwise returns std::nullopt. + std::optional getTargetAddressSpaceValue() const; + + /// Returns true if this pointer type uses a language (logical) address space. + bool hasLogicalAddressSpace() const; + + /// Returns the logical CIR address space if present, otherwise returns + /// cir::AddressSpace::Default. + cir::AddressSpace getLogicalAddressSpace() const; }]; } diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp index af9483d47a0d..dc241106e1c7 100644 --- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp @@ -77,9 +77,19 @@ static void printConstPtr(mlir::AsmPrinter &p, mlir::IntegerAttr value); //===----------------------------------------------------------------------===// mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, - cir::AddressSpace &addrSpace); + cir::AddressSpace &addrSpace) { + llvm::SMLoc loc = p.getCurrentLocation(); + mlir::FailureOr result = + mlir::FieldParser::parse(p); + if (mlir::failed(result)) + return p.emitError(loc, "expected address space keyword"); + addrSpace = result.value(); + return mlir::success(); +} -void printAddressSpaceValue(mlir::AsmPrinter &p, cir::AddressSpace addrSpace); +void printAddressSpaceValue(mlir::AsmPrinter &p, cir::AddressSpace addrSpace) { + p << cir::stringifyEnum(addrSpace); +} //===----------------------------------------------------------------------===// // Tablegen defined attributes diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index 4387142ac8c5..a5fa898d019f 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -2477,10 +2477,19 @@ cir::GetGlobalOp::verifySymbolUses(SymbolTableCollection &symbolTable) { << "' does not reference a valid cir.global or cir.func"; mlir::Type symTy; - cir::AddressSpace symAddrSpace{}; + mlir::Attribute symAddrSpaceAttr{}; if (auto g = dyn_cast(op)) { symTy = g.getSymType(); - symAddrSpace = g.getAddrSpace(); + // Convert enum to attribute for comparison + cir::AddressSpace symAddrSpace = g.getAddrSpace(); + if (symAddrSpace == cir::AddressSpace::Default) { + symAddrSpaceAttr = nullptr; + } else if (cir::isTargetAddressSpace(symAddrSpace)) { + unsigned targetAS = cir::getTargetAddressSpaceValue(symAddrSpace); + symAddrSpaceAttr = cir::TargetAddressSpaceAttr::get(getContext(), targetAS); + } else { + symAddrSpaceAttr = cir::AddressSpaceAttr::get(getContext(), symAddrSpace); + } // Verify that for thread local global access, the global needs to // be marked with tls bits. if (getTls() && !g.getTlsModel()) @@ -2496,7 +2505,7 @@ cir::GetGlobalOp::verifySymbolUses(SymbolTableCollection &symbolTable) { << resultType.getPointee() << "' does not match type " << symTy << " of the global @" << getName(); - if (symAddrSpace != resultType.getAddrSpace()) { + if (symAddrSpaceAttr != resultType.getAddrSpace()) { return emitOpError() << "result type address space does not match the address " "space of the global @" diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index 65103b68b3ac..0ff9b07f568f 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -11,6 +11,7 @@ //===----------------------------------------------------------------------===// #include "clang/CIR/Dialect/IR/CIRTypes.h" +#include "clang/Basic/AddressSpaces.h" #include "clang/CIR/Dialect/IR/CIRAttrs.h" #include "clang/CIR/Dialect/IR/CIRDialect.h" #include "clang/CIR/Dialect/IR/CIROpsEnums.h" @@ -32,6 +33,7 @@ #include "llvm/ADT/STLExtras.h" #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/TypeSwitch.h" +#include "llvm/Support/Casting.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/MathExtras.h" #include @@ -68,9 +70,9 @@ static void printFuncTypeParams(mlir::AsmPrinter &p, //===----------------------------------------------------------------------===// mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, - cir::AddressSpace &addrSpace); + mlir::Attribute &attr); -void printAddressSpaceValue(mlir::AsmPrinter &p, cir::AddressSpace addrSpace); +void printAddressSpaceValue(mlir::AsmPrinter &printer, mlir::Attribute attr); //===----------------------------------------------------------------------===// // Get autogenerated stuff @@ -781,8 +783,8 @@ LongDoubleType::getTypeSizeInBits(const mlir::DataLayout &dataLayout, uint64_t LongDoubleType::getABIAlignment(const mlir::DataLayout &dataLayout, mlir::DataLayoutEntryListRef params) const { - return mlir::cast(getUnderlying()).getABIAlignment( - dataLayout, params); + return mlir::cast(getUnderlying()) + .getABIAlignment(dataLayout, params); } //===----------------------------------------------------------------------===// @@ -962,12 +964,23 @@ cir::AddressSpace cir::toCIRAddressSpace(clang::LangAS langAS) { } mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, - cir::AddressSpace &addrSpace) { + mlir::Attribute &attr) { + llvm::SMLoc loc = p.getCurrentLocation(); - mlir::FailureOr result = - mlir::FieldParser::parse(p); - if (mlir::failed(result)) - return p.emitError(loc, "expected address space keyword"); + + // Try to parse target address space first. + attr = nullptr; + if (p.parseOptionalKeyword("target_address_space").succeeded()) { + unsigned val; + if (p.parseLParen()) + p.emitError(loc, "expected '(' after target_address_space"); + + if(p.parseInteger(val) || p.parseRParen()) + return p.emitError(loc, "expected target_address_space value"); + + attr = cir::TargetAddressSpaceAttr::get(p.getContext(), val); + return mlir::success(); + } // Address space is either a target address space or a regular one. // - If it is a target address space, we expect a value to follow in the form @@ -977,31 +990,108 @@ mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, // `cir::getMaxEnumValForAddressSpace()`. This allows us to use // the same enum for both regular and target address spaces. // - Otherwise, we just use the parsed value. - if (cir::isTargetAddressSpace(result.value())) { - if (p.parseLess()) - return p.emitError(loc, "expected '<' after target address space"); - int64_t targetValue; - if (p.parseInteger(targetValue) || p.parseGreater()) - return p.emitError(loc, "expected target address space value"); + // Try to parse language specific address space. + if (p.parseOptionalKeyword("clang_address_space").succeeded()) { + if (p.parseLParen()) + return p.emitError(loc, "expected '(' after clang_address_space"); + mlir::FailureOr result = + mlir::FieldParser::parse(p); - addrSpace = cir::computeTargetAddressSpace(targetValue); - } else { - addrSpace = result.value(); + if (mlir::failed(result) || p.parseRParen()) + return p.emitError(loc, "expected clang_address_space value"); + + attr = cir::AddressSpaceAttr::get(p.getContext(), result.value()); + return mlir::success(); } return mlir::success(); } -// Prints the address space value in the form of: -// - `target` for target address spaces -// - or just the address space name for regular address spaces. -void printAddressSpaceValue(mlir::AsmPrinter &p, cir::AddressSpace addrSpace) { - if (cir::isTargetAddressSpace(addrSpace)) - p << cir::stringifyEnum(cir::AddressSpace::Target) << '<' - << cir::getTargetAddressSpaceValue(addrSpace) << '>'; - else - p << cir::stringifyEnum(addrSpace); +void printAddressSpaceValue(mlir::AsmPrinter &printer, mlir::Attribute attr) { + if (!attr) + return; + + if (auto logical = dyn_cast(attr)) { + printer << "clang_address_space(" + << cir::stringifyAddressSpace(logical.getValue()) << ')'; + ; + return; + } + + if (auto target = dyn_cast(attr)) { + printer << "target_address_space(" << target.getValue() << ')'; + return; + } + + llvm_unreachable("unexpected address-space attribute kind"); +} + +mlir::Attribute cir::toCIRAddressSpaceAttr(mlir::MLIRContext *ctx, + clang::LangAS langAS) { + using clang::LangAS; + + if (clang::isTargetAddressSpace(langAS)) { + unsigned targetAS = clang::toTargetAddressSpace(langAS); + return cir::TargetAddressSpaceAttr::get(ctx, targetAS); + } + + return cir::AddressSpaceAttr::get(ctx, toCIRAddressSpace(langAS)); +} + +cir::AddressSpace cir::getCIRAddressSpaceFromAttr(mlir::Attribute attr) { + if (!attr) + return AddressSpace::Default; + if (auto addrSpaceAttr = mlir::dyn_cast(attr)) + return addrSpaceAttr.getValue(); + if (auto targetAddrSpaceAttr = + mlir::dyn_cast(attr)) + return cir::computeTargetAddressSpace(targetAddrSpaceAttr.getValue()); + return AddressSpace::Default; +} + +//===----------------------------------------------------------------------===// +// PointerType Definitions +//===----------------------------------------------------------------------===// + +mlir::LogicalResult cir::PointerType::verify( + llvm::function_ref emitError, + mlir::Type pointee, mlir::Attribute addrSpace) { + if (auto as = addrSpace) { + if (!mlir::isa(as) && + !mlir::isa(as)) { + return emitError() << "pointer address space must be either " + "!cir.address_space or !cir.target_address_space"; + } + } + + return success(); +} + +bool PointerType::hasTargetAddressSpace() const { + mlir::Attribute addrSpace = getAddrSpace(); + if (!addrSpace) + return false; + return mlir::isa(addrSpace); +} + +std::optional PointerType::getTargetAddressSpaceValue() const { + if (auto targetAddrSpace = + mlir::dyn_cast(getAddrSpace())) + return targetAddrSpace.getValue(); + return std::nullopt; +} + +bool PointerType::hasLogicalAddressSpace() const { + auto as = getAddrSpace(); + return as && llvm::isa_and_nonnull(as); +} + +cir::AddressSpace PointerType::getLogicalAddressSpace() const { + auto as = getAddrSpace(); + if (auto logAS = llvm::dyn_cast_or_null(as)) + return logAS.getValue(); + return cir::AddressSpace::Default; } //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 70c74cd0bab7..b5f397e5f216 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -382,6 +382,44 @@ void lowerAnnotationValue( } } +/// Extract address space value from either AddressSpaceAttr or TargetAddressSpaceAttr. +/// Returns the CIR AddressSpace enum value for use with PointerType. +static cir::AddressSpace getAddressSpaceFromAttr(mlir::Attribute attr) { + if (auto addrSpaceAttr = mlir::dyn_cast(attr)) + return addrSpaceAttr.getValue(); + if (auto targetAddrSpaceAttr = mlir::dyn_cast(attr)) + return cir::computeTargetAddressSpace(targetAddrSpaceAttr.getValue()); + llvm_unreachable("Expected AddressSpaceAttr or TargetAddressSpaceAttr"); +} + +/// Convert a CIR address space (enum) to a target-specific LLVM address space value. +/// This function handles both: +/// 1. Target address spaces: Encoded as offsets in the enum, extracted directly +/// 2. Language address spaces: Mapped via TargetLoweringInfo (e.g., OpenCL/CUDA) +static unsigned +getTargetAddrSpaceFromCIRAddrSpace(cir::AddressSpace addrSpace, + cir::LowerModule *lowerModule) { + if (addrSpace == cir::AddressSpace::Default) + return 0; // Default address space is always 0 in LLVM. + + // Target address spaces are encoded as enum offsets and can be extracted directly + if (cir::isTargetAddressSpace(addrSpace)) + return cir::getTargetAddressSpaceValue(addrSpace); + + // Language address spaces (e.g., OpenCL, CUDA) need target-specific mapping + assert(lowerModule && "CIR AS map is not available"); + return lowerModule->getTargetLoweringInfo() + .getTargetAddrSpaceFromCIRAddrSpace(addrSpace); +} + +/// Convert a CIR address space attribute (AddressSpaceAttr or TargetAddressSpaceAttr) +/// directly to a target-specific LLVM address space value. +static unsigned +getTargetAddrSpaceFromAttr(mlir::Attribute attr, cir::LowerModule *lowerModule) { + cir::AddressSpace addrSpace = getAddressSpaceFromAttr(attr); + return getTargetAddrSpaceFromCIRAddrSpace(addrSpace, lowerModule); +} + // Get addrspace by converting a pointer type. // TODO: The approach here is a little hacky. We should access the target info // directly to convert the address space of global op, similar to what we do @@ -390,7 +428,7 @@ unsigned getGlobalOpTargetAddrSpace(mlir::ConversionPatternRewriter &rewriter, const mlir::TypeConverter *converter, cir::GlobalOp op) { auto tempPtrTy = cir::PointerType::get(rewriter.getContext(), op.getSymType(), - op.getAddrSpace()); + op.getAddrSpaceAttr()); return cast(converter->convertType(tempPtrTy)) .getAddressSpace(); } @@ -5000,20 +5038,6 @@ std::unique_ptr prepareLowerModule(mlir::ModuleOp module) { return cir::createLowerModule(module, rewriter); } -static unsigned -getTargetAddrSpaceFromCIRAddrSpace(cir::AddressSpace addrSpace, - cir::LowerModule *lowerModule) { - if (addrSpace == cir::AddressSpace::Default) - return 0; // Default address space is always 0 in LLVM. - - if (cir::isTargetAddressSpace(addrSpace)) - return cir::getTargetAddressSpaceValue(addrSpace); - - assert(lowerModule && "CIR AS map is not available"); - return lowerModule->getTargetLoweringInfo() - .getTargetAddrSpaceFromCIRAddrSpace(addrSpace); -} - // FIXME: change the type of lowerModule to `LowerModule &` to have better // lambda capturing experience. Also blocked by makeTripleAlwaysPresent. void prepareTypeConverter(mlir::LLVMTypeConverter &converter, @@ -5021,8 +5045,10 @@ void prepareTypeConverter(mlir::LLVMTypeConverter &converter, cir::LowerModule *lowerModule) { converter.addConversion([&, lowerModule](cir::PointerType type) -> mlir::Type { - unsigned addrSpace = - getTargetAddrSpaceFromCIRAddrSpace(type.getAddrSpace(), lowerModule); + mlir::Attribute addrSpaceAttr = type.getAddrSpace(); + unsigned addrSpace = addrSpaceAttr + ? getTargetAddrSpaceFromAttr(addrSpaceAttr, lowerModule) + : 0; // Default address space return mlir::LLVM::LLVMPointerType::get(type.getContext(), addrSpace); }); converter.addConversion([&](cir::VPtrType type) -> mlir::Type { From c9594487f6decfcc7d6b89a8a637aff745a77e4e Mon Sep 17 00:00:00 2001 From: David Rivera Date: Mon, 17 Nov 2025 09:37:53 -0500 Subject: [PATCH 02/13] Polish and remove redundant functions --- .../CIR/Dialect/Builder/CIRBaseBuilder.h | 19 ++---- .../include/clang/CIR/Dialect/IR/CIRAttrs.td | 4 +- clang/include/clang/CIR/Dialect/IR/CIRTypes.h | 5 +- .../include/clang/CIR/Dialect/IR/CIRTypes.td | 14 +--- clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 8 ++- clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 46 ++++--------- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 66 ++++++++----------- 7 files changed, 60 insertions(+), 102 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h index e24684cdccca..6b4328b9aac2 100644 --- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h +++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h @@ -111,13 +111,11 @@ class CIRBaseBuilderTy : public mlir::OpBuilder { return cir::PointerType::get(ty, addrSpaceAttr); } - /// Create a pointer type from a cir::AddressSpace enum. - /// This converts the enum to the appropriate attribute. cir::PointerType getPointerTo(mlir::Type ty, cir::AddressSpace addrSpace) { if (addrSpace == cir::AddressSpace::Default) return getPointerTo(ty); if (cir::isTargetAddressSpace(addrSpace)) { - unsigned targetAS = cir::getTargetAddressSpaceValue(addrSpace); + unsigned targetAS = cir::getTargetAddressSpaceValueFromCIRAS(addrSpace); auto attr = cir::TargetAddressSpaceAttr::get(getContext(), targetAS); return getPointerTo(ty, attr); } @@ -125,24 +123,15 @@ class CIRBaseBuilderTy : public mlir::OpBuilder { return getPointerTo(ty, attr); } - /// Create a pointer type from a LangAS. - /// This converts the LangAS to the appropriate attribute (AddressSpaceAttr or TargetAddressSpaceAttr). cir::PointerType getPointerTo(mlir::Type ty, clang::LangAS langAS) { - if(langAS == clang::LangAS::Default) + if (langAS == clang::LangAS::Default) return getPointerTo(ty); - mlir::Attribute addrSpaceAttr = cir::toCIRAddressSpaceAttr(getContext(), langAS); + mlir::Attribute addrSpaceAttr = + cir::toCIRAddressSpaceAttr(getContext(), langAS); return getPointerTo(ty, addrSpaceAttr); } - /// Create a pointer type with a target-specific address space value. - /// This is used for address spaces specified via __attribute__((address_space(N))). - cir::PointerType getPointerToWithTargetAddrSpace(mlir::Type ty, unsigned targetAS) { - assert(clang::isTargetAddressSpace(static_cast(targetAS))); - auto attr = cir::TargetAddressSpaceAttr::get(getContext(), targetAS); - return getPointerTo(ty, attr); - } - cir::PointerType getVoidPtrTy(clang::LangAS langAS = clang::LangAS::Default) { return getPointerTo(cir::VoidType::get(getContext()), langAS); } diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td index f35cb27bdc89..94deeea69722 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -997,7 +997,7 @@ def CIR_AddressSpaceAttr : CIR_EnumAttr { } unsigned $cppClass::getTargetValue() const { - return cir::getTargetAddressSpaceValue(getValue()); + return cir::getTargetAddressSpaceValueFromCIRAS(getValue()); } }]; } @@ -1011,7 +1011,7 @@ def CIR_TargetAddressSpaceAttr : CIR_Attr<"TargetAddressSpace", "target_address_ let description = [{ The TargetAddressSpaceAttr represents a target-specific numeric address space, corresponding to the LLVM IR `addressspace` qualifier and the clang - `address_space` attribute. + `target_address_space` attribute. A value of zero represents the default address space. The semantics of non-zero address spaces are target-specific. diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h index b5628d29cd17..b85a8acaa8d8 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h @@ -39,7 +39,8 @@ cir::AddressSpace toCIRAddressSpace(clang::LangAS langAS); /// Convert a LangAS to the appropriate address space attribute. /// Returns AddressSpaceAttr for language-specific address spaces, /// or TargetAddressSpaceAttr for target-specific address spaces. -mlir::Attribute toCIRAddressSpaceAttr(mlir::MLIRContext *ctx, clang::LangAS langAS); +mlir::Attribute toCIRAddressSpaceAttr(mlir::MLIRContext *ctx, + clang::LangAS langAS); /// Extract the AddressSpace enum from an address space attribute. /// Returns Default if the attribute is null. @@ -67,7 +68,7 @@ constexpr bool isLangAddressSpace(cir::AddressSpace as) { return !isTargetAddressSpace(as); } -constexpr unsigned getTargetAddressSpaceValue(cir::AddressSpace as) { +constexpr unsigned getTargetAddressSpaceValueFromCIRAS(cir::AddressSpace as) { assert(isTargetAddressSpace(as) && "expected target address space"); return getAsUnsignedValue(as) - TargetAddressSpaceOffset; } diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td index bd078ab08222..6dd385a2ef09 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td @@ -244,12 +244,12 @@ def CIR_PointerType : CIR_Type<"Pointer", "ptr", [ let builders = [ TypeBuilderWithInferredContext<(ins "mlir::Type":$pointee, - CArg<"mlir::Attribute", "nullptr">:$addrSpace), [{ + CArg<"mlir::Attribute", "{}">:$addrSpace), [{ return $_get(pointee.getContext(), pointee, addrSpace); }]>, TypeBuilder<(ins "mlir::Type":$pointee, - CArg<"mlir::Attribute", "nullptr">:$addrSpace), [{ + CArg<"mlir::Attribute", "{}">:$addrSpace), [{ return $_get($_ctxt, pointee, addrSpace); }]> ]; @@ -291,16 +291,8 @@ def CIR_PointerType : CIR_Type<"Pointer", "ptr", [ /// Returns true if this pointer type uses a target address space. bool hasTargetAddressSpace() const; - /// Returns the target address space value if this is a target address space, - /// otherwise returns std::nullopt. - std::optional getTargetAddressSpaceValue() const; - /// Returns true if this pointer type uses a language (logical) address space. - bool hasLogicalAddressSpace() const; - - /// Returns the logical CIR address space if present, otherwise returns - /// cir::AddressSpace::Default. - cir::AddressSpace getLogicalAddressSpace() const; + bool hasLanguageAddressSpace() const; }]; } diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index a5fa898d019f..48468a23226c 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -2483,10 +2483,12 @@ cir::GetGlobalOp::verifySymbolUses(SymbolTableCollection &symbolTable) { // Convert enum to attribute for comparison cir::AddressSpace symAddrSpace = g.getAddrSpace(); if (symAddrSpace == cir::AddressSpace::Default) { - symAddrSpaceAttr = nullptr; + symAddrSpaceAttr = {}; } else if (cir::isTargetAddressSpace(symAddrSpace)) { - unsigned targetAS = cir::getTargetAddressSpaceValue(symAddrSpace); - symAddrSpaceAttr = cir::TargetAddressSpaceAttr::get(getContext(), targetAS); + unsigned targetAS = + cir::getTargetAddressSpaceValueFromCIRAS(symAddrSpace); + symAddrSpaceAttr = + cir::TargetAddressSpaceAttr::get(getContext(), targetAS); } else { symAddrSpaceAttr = cir::AddressSpaceAttr::get(getContext(), symAddrSpace); } diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index 0ff9b07f568f..9ca9542ed96f 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -957,9 +957,9 @@ cir::AddressSpace cir::toCIRAddressSpace(clang::LangAS langAS) { case LangAS::wasm_funcref: llvm_unreachable("NYI"); default: - // Target address space offset arithmetics - return static_cast(clang::toTargetAddressSpace(langAS) + - cir::getMaxEnumValForAddressSpace()); + // NOTE: In theory with TargetAddressSpaceAttr, we don't care at all about + // representing target AS here. + llvm_unreachable("unknown/unsupported clang language address space"); } } @@ -975,9 +975,9 @@ mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, if (p.parseLParen()) p.emitError(loc, "expected '(' after target_address_space"); - if(p.parseInteger(val) || p.parseRParen()) - return p.emitError(loc, "expected target_address_space value"); - + if (p.parseInteger(val) || p.parseRParen()) + return p.emitError(loc, "expected target_address_space value"); + attr = cir::TargetAddressSpaceAttr::get(p.getContext(), val); return mlir::success(); } @@ -1008,19 +1008,19 @@ mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, return mlir::success(); } -void printAddressSpaceValue(mlir::AsmPrinter &printer, mlir::Attribute attr) { +void printAddressSpaceValue(mlir::AsmPrinter &p, mlir::Attribute attr) { if (!attr) return; if (auto logical = dyn_cast(attr)) { - printer << "clang_address_space(" - << cir::stringifyAddressSpace(logical.getValue()) << ')'; - ; + p << "clang_address_space(" + << cir::stringifyAddressSpace(logical.getValue()) << ')'; + return; } if (auto target = dyn_cast(attr)) { - printer << "target_address_space(" << target.getValue() << ')'; + p << "target_address_space(" << target.getValue() << ')'; return; } @@ -1069,29 +1069,11 @@ mlir::LogicalResult cir::PointerType::verify( } bool PointerType::hasTargetAddressSpace() const { - mlir::Attribute addrSpace = getAddrSpace(); - if (!addrSpace) - return false; - return mlir::isa(addrSpace); -} - -std::optional PointerType::getTargetAddressSpaceValue() const { - if (auto targetAddrSpace = - mlir::dyn_cast(getAddrSpace())) - return targetAddrSpace.getValue(); - return std::nullopt; -} - -bool PointerType::hasLogicalAddressSpace() const { - auto as = getAddrSpace(); - return as && llvm::isa_and_nonnull(as); + return mlir::isa_and_nonnull(getAddrSpace()); } -cir::AddressSpace PointerType::getLogicalAddressSpace() const { - auto as = getAddrSpace(); - if (auto logAS = llvm::dyn_cast_or_null(as)) - return logAS.getValue(); - return cir::AddressSpace::Default; +bool PointerType::hasLanguageAddressSpace() const { + return mlir::isa_and_nonnull(getAddrSpace()); } //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index b5f397e5f216..ef2d187d5da6 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -382,43 +382,7 @@ void lowerAnnotationValue( } } -/// Extract address space value from either AddressSpaceAttr or TargetAddressSpaceAttr. -/// Returns the CIR AddressSpace enum value for use with PointerType. -static cir::AddressSpace getAddressSpaceFromAttr(mlir::Attribute attr) { - if (auto addrSpaceAttr = mlir::dyn_cast(attr)) - return addrSpaceAttr.getValue(); - if (auto targetAddrSpaceAttr = mlir::dyn_cast(attr)) - return cir::computeTargetAddressSpace(targetAddrSpaceAttr.getValue()); - llvm_unreachable("Expected AddressSpaceAttr or TargetAddressSpaceAttr"); -} - -/// Convert a CIR address space (enum) to a target-specific LLVM address space value. -/// This function handles both: -/// 1. Target address spaces: Encoded as offsets in the enum, extracted directly -/// 2. Language address spaces: Mapped via TargetLoweringInfo (e.g., OpenCL/CUDA) -static unsigned -getTargetAddrSpaceFromCIRAddrSpace(cir::AddressSpace addrSpace, - cir::LowerModule *lowerModule) { - if (addrSpace == cir::AddressSpace::Default) - return 0; // Default address space is always 0 in LLVM. - - // Target address spaces are encoded as enum offsets and can be extracted directly - if (cir::isTargetAddressSpace(addrSpace)) - return cir::getTargetAddressSpaceValue(addrSpace); - - // Language address spaces (e.g., OpenCL, CUDA) need target-specific mapping - assert(lowerModule && "CIR AS map is not available"); - return lowerModule->getTargetLoweringInfo() - .getTargetAddrSpaceFromCIRAddrSpace(addrSpace); -} -/// Convert a CIR address space attribute (AddressSpaceAttr or TargetAddressSpaceAttr) -/// directly to a target-specific LLVM address space value. -static unsigned -getTargetAddrSpaceFromAttr(mlir::Attribute attr, cir::LowerModule *lowerModule) { - cir::AddressSpace addrSpace = getAddressSpaceFromAttr(attr); - return getTargetAddrSpaceFromCIRAddrSpace(addrSpace, lowerModule); -} // Get addrspace by converting a pointer type. // TODO: The approach here is a little hacky. We should access the target info @@ -5038,6 +5002,34 @@ std::unique_ptr prepareLowerModule(mlir::ModuleOp module) { return cir::createLowerModule(module, rewriter); } +static unsigned +getTargetAddrSpaceFromCIRAddrSpace(cir::AddressSpace addrSpace, + cir::LowerModule *lowerModule) { + if (addrSpace == cir::AddressSpace::Default) + return 0; // Default address space is always 0 in LLVM. + + if (cir::isTargetAddressSpace(addrSpace)) + return cir::getTargetAddressSpaceValueFromCIRAS(addrSpace); + + assert(lowerModule && "CIR AS map is not available"); + return lowerModule->getTargetLoweringInfo() + .getTargetAddrSpaceFromCIRAddrSpace(addrSpace); +} + +static unsigned getTargetAddrSpaceFromASAttr(mlir::Attribute attr, + cir::LowerModule *lowerModule) { + assert(mlir::isa_and_nonnull(attr) || + mlir::isa_and_nonnull(attr)); + + if (auto targetAddrSpaceAttr = + mlir::dyn_cast(attr)) + return targetAddrSpaceAttr.getValue(); + + auto addrSpaceAttr = mlir::dyn_cast(attr); + return getTargetAddrSpaceFromCIRAddrSpace(addrSpaceAttr.getValue(), + lowerModule); +} + // FIXME: change the type of lowerModule to `LowerModule &` to have better // lambda capturing experience. Also blocked by makeTripleAlwaysPresent. void prepareTypeConverter(mlir::LLVMTypeConverter &converter, @@ -5047,7 +5039,7 @@ void prepareTypeConverter(mlir::LLVMTypeConverter &converter, lowerModule](cir::PointerType type) -> mlir::Type { mlir::Attribute addrSpaceAttr = type.getAddrSpace(); unsigned addrSpace = addrSpaceAttr - ? getTargetAddrSpaceFromAttr(addrSpaceAttr, lowerModule) + ? getTargetAddrSpaceFromASAttr(addrSpaceAttr, lowerModule) : 0; // Default address space return mlir::LLVM::LLVMPointerType::get(type.getContext(), addrSpace); }); From a6c0beae2707ac105f6a42dec7c34c8b4e48b784 Mon Sep 17 00:00:00 2001 From: David Rivera Date: Mon, 17 Nov 2025 11:57:06 -0500 Subject: [PATCH 03/13] Handle hybrid AS approach with GlobalOps --- .../CIR/Dialect/IR/CIRAttrConstraints.td | 1 + clang/include/clang/CIR/Dialect/IR/CIROps.td | 6 +++--- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 4 ++-- clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 21 +++++++++---------- clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 5 +++++ 5 files changed, 21 insertions(+), 16 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td index 0118102ec607..a4dd4f6375b4 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td @@ -106,6 +106,7 @@ def CIR_AnyAddressSpaceAttr : AnyAttrOf<[ CIR_TargetAddressSpaceAttrConstraint ]> { string cppType = "::mlir::Attribute"; + let constBuilderCall = "nullptr"; } #endif // CLANG_CIR_DIALECT_IR_CIRATTRCONSTRAINTS_TD diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index 5a512b7645aa..7d0de699e15f 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -2617,8 +2617,8 @@ def CIR_GlobalOp : CIR_Op<"global", [ TypeAttr:$sym_type, CIR_GlobalLinkageKind:$linkage, DefaultValuedAttr< - CIR_AddressSpaceAttr, - "AddressSpace::Default" + CIR_AnyAddressSpaceAttr, + "{}" >:$addr_space, OptionalAttr:$tls_model, // Note this can also be a FlatSymbolRefAttr @@ -2642,7 +2642,7 @@ def CIR_GlobalOp : CIR_Op<"global", [ (`comdat` $comdat^)? ($tls_model^)? (`dso_local` $dso_local^)? - (`addrspace` `` $addr_space^)? + ( `,` ` ` custom($addr_space)^ )? $sym_name custom($sym_type, $initial_value, $ctorRegion, $dtorRegion) ($annotations^)? diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 56121a721ae1..dc8214876a0f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -1112,9 +1112,9 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty, entry = dyn_cast_or_null(v); } - cir::AddressSpace cirAS = cir::toCIRAddressSpace(langAS); + mlir::Attribute cirAS = cir::toCIRAddressSpaceAttr(&getMLIRContext(), langAS); if (entry) { - cir::AddressSpace entryCIRAS = entry.getAddrSpace(); + mlir::Attribute entryCIRAS = entry.getAddrSpace(); if (WeakRefReferences.erase(entry)) { if (d && !d->hasAttr()) { auto lt = cir::GlobalLinkageKind::ExternalLinkage; diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index 48468a23226c..f017a55b7c10 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -28,6 +28,7 @@ #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/LLVMIR/LLVMTypes.h" +#include "mlir/IR/Attributes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinAttributes.h" #include "mlir/IR/BuiltinTypes.h" @@ -55,6 +56,8 @@ using namespace mlir; #include "clang/CIR/Interfaces/ASTAttrInterfaces.h" #include "clang/CIR/Interfaces/CIROpInterfaces.h" #include +#include +#include //===----------------------------------------------------------------------===// // CIR Dialect @@ -304,6 +307,12 @@ static void printOmittedTerminatorRegion(mlir::OpAsmPrinter &printer, } } +mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, + mlir::Attribute &attr); + +void printAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp op, + mlir::Attribute attr); + //===----------------------------------------------------------------------===// // AllocaOp //===----------------------------------------------------------------------===// @@ -2481,17 +2490,7 @@ cir::GetGlobalOp::verifySymbolUses(SymbolTableCollection &symbolTable) { if (auto g = dyn_cast(op)) { symTy = g.getSymType(); // Convert enum to attribute for comparison - cir::AddressSpace symAddrSpace = g.getAddrSpace(); - if (symAddrSpace == cir::AddressSpace::Default) { - symAddrSpaceAttr = {}; - } else if (cir::isTargetAddressSpace(symAddrSpace)) { - unsigned targetAS = - cir::getTargetAddressSpaceValueFromCIRAS(symAddrSpace); - symAddrSpaceAttr = - cir::TargetAddressSpaceAttr::get(getContext(), targetAS); - } else { - symAddrSpaceAttr = cir::AddressSpaceAttr::get(getContext(), symAddrSpace); - } + symAddrSpaceAttr = g.getAddrSpace(); // Verify that for thread local global access, the global needs to // be marked with tls bits. if (getTls() && !g.getTlsModel()) diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index 9ca9542ed96f..26ad50e6b7eb 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -1027,6 +1027,11 @@ void printAddressSpaceValue(mlir::AsmPrinter &p, mlir::Attribute attr) { llvm_unreachable("unexpected address-space attribute kind"); } +void printAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp, + mlir::Attribute attr) { + printAddressSpaceValue(printer, attr); +} + mlir::Attribute cir::toCIRAddressSpaceAttr(mlir::MLIRContext *ctx, clang::LangAS langAS) { using clang::LangAS; From 4a4ea3a5a451b89f3a1bf578383e1e8002e9fa1c Mon Sep 17 00:00:00 2001 From: David Rivera Date: Tue, 18 Nov 2025 15:16:28 -0500 Subject: [PATCH 04/13] Correct comma on global asm format by providing an optional parser --- clang/include/clang/CIR/Dialect/IR/CIROps.td | 2 +- clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 8 ++++---- clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 17 ++++++++++++----- 3 files changed, 17 insertions(+), 10 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index 7d0de699e15f..5118e800d9b4 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -2642,7 +2642,7 @@ def CIR_GlobalOp : CIR_Op<"global", [ (`comdat` $comdat^)? ($tls_model^)? (`dso_local` $dso_local^)? - ( `,` ` ` custom($addr_space)^ )? + (` ` custom($addr_space)^ )? $sym_name custom($sym_type, $initial_value, $ctorRegion, $dtorRegion) ($annotations^)? diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index f017a55b7c10..145bbdd9dd01 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -307,11 +307,11 @@ static void printOmittedTerminatorRegion(mlir::OpAsmPrinter &printer, } } -mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, - mlir::Attribute &attr); +mlir::OptionalParseResult parseGlobalAddressSpaceValue(mlir::AsmParser &p, + mlir::Attribute &attr); -void printAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp op, - mlir::Attribute attr); +void printGlobalAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp op, + mlir::Attribute attr); //===----------------------------------------------------------------------===// // AllocaOp diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index 26ad50e6b7eb..2a9427717529 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -1005,7 +1005,7 @@ mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, return mlir::success(); } - return mlir::success(); + return mlir::failure(); } void printAddressSpaceValue(mlir::AsmPrinter &p, mlir::Attribute attr) { @@ -1027,7 +1027,14 @@ void printAddressSpaceValue(mlir::AsmPrinter &p, mlir::Attribute attr) { llvm_unreachable("unexpected address-space attribute kind"); } -void printAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp, +mlir::OptionalParseResult parseGlobalAddressSpaceValue(mlir::AsmParser &p, + mlir::Attribute &attr) { + if (!parseAddressSpaceValue(p, attr)) + return mlir::failure(); + return mlir::success(); +} + +void printGlobalAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp, mlir::Attribute attr) { printAddressSpaceValue(printer, attr); } @@ -1062,9 +1069,9 @@ cir::AddressSpace cir::getCIRAddressSpaceFromAttr(mlir::Attribute attr) { mlir::LogicalResult cir::PointerType::verify( llvm::function_ref emitError, mlir::Type pointee, mlir::Attribute addrSpace) { - if (auto as = addrSpace) { - if (!mlir::isa(as) && - !mlir::isa(as)) { + if (addrSpace) { + if (!mlir::isa(addrSpace) && + !mlir::isa(addrSpace)) { return emitError() << "pointer address space must be either " "!cir.address_space or !cir.target_address_space"; } From 8aff82051e519c86b31d315e700e7e1f28eb3415 Mon Sep 17 00:00:00 2001 From: David Rivera Date: Wed, 19 Nov 2025 09:31:11 -0500 Subject: [PATCH 05/13] Change Builder for globalOps to take an empty attr to represent default AS and Correct test asm format --- .../clang/CIR/Dialect/IR/CIREnumAttr.td | 2 +- clang/include/clang/CIR/Dialect/IR/CIROps.td | 2 +- clang/lib/CIR/CodeGen/CIRGenBuilder.h | 11 +++--- clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp | 2 +- clang/lib/CIR/CodeGen/CIRGenDecl.cpp | 3 +- clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 8 ++-- clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp | 9 +++-- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 10 ++--- clang/lib/CIR/CodeGen/CIRGenModule.h | 3 +- clang/lib/CIR/CodeGen/TargetInfo.cpp | 4 +- clang/lib/CIR/CodeGen/TargetInfo.h | 5 ++- clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 7 ++-- clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 30 +++++++------- clang/test/CIR/CodeGen/CUDA/address-spaces.cu | 6 +-- clang/test/CIR/CodeGen/CUDA/global-vars.cu | 6 +-- clang/test/CIR/CodeGen/CUDA/surface.cu | 2 +- clang/test/CIR/CodeGen/CUDA/texture.cu | 2 +- clang/test/CIR/CodeGen/HIP/address-spaces.cpp | 6 +-- .../CIR/CodeGen/OpenCL/addrspace-alloca.cl | 14 +++---- clang/test/CIR/CodeGen/OpenCL/array-decay.cl | 6 +-- clang/test/CIR/CodeGen/OpenCL/global.cl | 14 +++---- clang/test/CIR/CodeGen/OpenCL/printf.cl | 10 ++--- .../test/CIR/CodeGen/OpenCL/static-vardecl.cl | 12 +++--- clang/test/CIR/CodeGen/OpenCL/str_literals.cl | 10 ++--- .../CIR/CodeGen/address-space-conversion.cpp | 30 +++++++------- clang/test/CIR/CodeGen/address-space.c | 4 +- clang/test/CIR/IR/address-space.cir | 24 ++++++------ clang/test/CIR/IR/cast.cir | 4 +- clang/test/CIR/IR/global.cir | 12 +++--- clang/test/CIR/IR/invalid.cir | 39 ++++++++----------- clang/test/CIR/Lowering/address-space.cir | 24 ++++++------ clang/test/CIR/Transforms/merge-cleanups.cir | 10 ++--- 32 files changed, 162 insertions(+), 169 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td b/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td index 46c6039674b0..56e5a4465795 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td +++ b/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td @@ -42,7 +42,7 @@ class CIR_DefaultValuedEnumParameter } def CIR_AddressSpace : CIR_I32EnumAttr< - "AddressSpace", "address space kind", [ + "AddressSpace", "clang address space kind", [ I32EnumAttrCase<"Default", 0, "default">, I32EnumAttrCase<"OffloadPrivate", 1, "offload_private">, I32EnumAttrCase<"OffloadLocal", 2, "offload_local">, diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index 5118e800d9b4..921b8c008cdb 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -2666,7 +2666,7 @@ def CIR_GlobalOp : CIR_Op<"global", [ // CIR defaults to external linkage. CArg<"cir::GlobalLinkageKind", "cir::GlobalLinkageKind::ExternalLinkage">:$linkage, - CArg<"cir::AddressSpace", "cir::AddressSpace::Default">:$addrSpace, + CArg<"mlir::Attribute", "{}">:$addrSpace, CArg<"llvm::function_ref", "nullptr">:$ctorBuilder, CArg<"llvm::function_ref", diff --git a/clang/lib/CIR/CodeGen/CIRGenBuilder.h b/clang/lib/CIR/CodeGen/CIRGenBuilder.h index 5b8391c9d408..bb19efddc641 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuilder.h +++ b/clang/lib/CIR/CodeGen/CIRGenBuilder.h @@ -749,7 +749,7 @@ class CIRGenBuilderTy : public cir::CIRBaseBuilderTy { [[nodiscard]] cir::GlobalOp createGlobal(mlir::ModuleOp module, mlir::Location loc, mlir::StringRef name, mlir::Type type, bool isConst, cir::GlobalLinkageKind linkage, - cir::AddressSpace addrSpace = cir::AddressSpace::Default) { + mlir::Attribute addrSpace = {}) { mlir::OpBuilder::InsertionGuard guard(*this); setInsertionPointToStart(module.getBody()); return cir::GlobalOp::create(*this, loc, name, type, isConst, linkage, @@ -759,10 +759,11 @@ class CIRGenBuilderTy : public cir::CIRBaseBuilderTy { /// Creates a versioned global variable. If the symbol is already taken, an ID /// will be appended to the symbol. The returned global must always be queried /// for its name so it can be referenced correctly. - [[nodiscard]] cir::GlobalOp createVersionedGlobal( - mlir::ModuleOp module, mlir::Location loc, mlir::StringRef name, - mlir::Type type, bool isConst, cir::GlobalLinkageKind linkage, - cir::AddressSpace addrSpace = cir::AddressSpace::Default) { + [[nodiscard]] cir::GlobalOp + createVersionedGlobal(mlir::ModuleOp module, mlir::Location loc, + mlir::StringRef name, mlir::Type type, bool isConst, + cir::GlobalLinkageKind linkage, + mlir::Attribute addrSpace = {}) { // Create a unique name if the given name is already taken. std::string uniqueName; if (unsigned version = GlobalsVersioning[name.str()]++) diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp index 2b64997de866..99bf99e61595 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp @@ -329,7 +329,7 @@ mlir::Operation *CIRGenNVCUDARuntime::getKernelHandle(cir::FuncOp fn, return CIRGenModule::createGlobalOp( cgm, fn->getLoc(), globalName, builder.getPointerTo(fn.getFunctionType()), true, - cir::AddressSpace::Default, + /*addrSpace=*/{}, /*insertPoint=*/nullptr); }); diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp index 6b5c942fa193..80006e34d192 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp @@ -479,7 +479,8 @@ CIRGenModule::getOrCreateStaticVarDecl(const VarDecl &D, Name = getStaticDeclName(*this, D); mlir::Type LTy = getTypes().convertTypeForMem(Ty); - cir::AddressSpace AS = cir::toCIRAddressSpace(getGlobalVarAddressSpace(&D)); + mlir::Attribute AS = cir::toCIRAddressSpaceAttr(&getMLIRContext(), + getGlobalVarAddressSpace(&D)); // OpenCL variables in local address space and CUDA shared // variables cannot have an initializer. diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index 4980e4ce7d11..3e3369d677da 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -2080,10 +2080,10 @@ LValue CIRGenFunction::emitCastLValue(const CastExpr *E) { case CK_AddressSpaceConversion: { LValue LV = emitLValue(E->getSubExpr()); QualType DestTy = getContext().getPointerType(E->getType()); - cir::AddressSpace SrcAS = - cir::toCIRAddressSpace(E->getSubExpr()->getType().getAddressSpace()); - cir::AddressSpace DestAS = - cir::toCIRAddressSpace(E->getType().getAddressSpace()); + mlir::Attribute SrcAS = cir::toCIRAddressSpaceAttr( + &getMLIRContext(), E->getSubExpr()->getType().getAddressSpace()); + mlir::Attribute DestAS = cir::toCIRAddressSpaceAttr( + &getMLIRContext(), E->getType().getAddressSpace()); mlir::Value V = getTargetHooks().performAddrSpaceCast( *this, LV.getPointer(), SrcAS, DestAS, convertType(DestTy)); return makeAddrLValue(Address(V, convertTypeForMem(E->getType()), diff --git a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp index 98a721c3532c..199b6dfe1327 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp @@ -27,6 +27,7 @@ #include "llvm/Support/ErrorHandling.h" #include +#include "mlir/IR/Attributes.h" #include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/Value.h" @@ -1689,10 +1690,12 @@ mlir::Value ScalarExprEmitter::VisitCastExpr(CastExpr *CE) { } // Since target may map different address spaces in AST to the same address // space, an address space conversion may end up as a bitcast. - cir::AddressSpace SrcAS = cir::toCIRAddressSpace( + mlir::Attribute SrcAS = cir::toCIRAddressSpaceAttr( + &CGF.getMLIRContext(), E->getType()->getPointeeType().getAddressSpace()); - cir::AddressSpace DestAS = - cir::toCIRAddressSpace(DestTy->getPointeeType().getAddressSpace()); + mlir::Attribute DestAS = cir::toCIRAddressSpaceAttr( + &CGF.getMLIRContext(), DestTy->getPointeeType().getAddressSpace()); + return CGF.CGM.getTargetCIRGenInfo().performAddrSpaceCast( CGF, Visit(E), SrcAS, DestAS, convertType(DestTy)); } diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index dc8214876a0f..6987f5ae9507 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -829,7 +829,7 @@ mlir::Value CIRGenModule::getGlobalValue(const Decl *d) { cir::GlobalOp CIRGenModule::createGlobalOp(CIRGenModule &cgm, mlir::Location loc, StringRef name, mlir::Type t, bool isConstant, - cir::AddressSpace addrSpace, + mlir::Attribute addrSpace, mlir::Operation *insertPoint, cir::GlobalLinkageKind linkage) { cir::GlobalOp g; @@ -1168,7 +1168,7 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty, return entry; } - auto declCIRAS = cir::toCIRAddressSpace(getGlobalVarAddressSpace(d)); + mlir::Attribute declCIRAS = cir::toCIRAddressSpaceAttr(&getMLIRContext(), getGlobalVarAddressSpace(d)); // TODO(cir): do we need to strip pointer casts for Entry? auto loc = getLoc(d->getSourceRange()); @@ -1773,8 +1773,8 @@ static cir::GlobalOp generateStringLiteral(mlir::Location loc, mlir::TypedAttr c, cir::GlobalLinkageKind lt, CIRGenModule &cgm, StringRef globalName, CharUnits alignment) { - cir::AddressSpace addrSpace = - cir::toCIRAddressSpace(cgm.getGlobalConstantAddressSpace()); + mlir::Attribute addrSpace = cir::toCIRAddressSpaceAttr( + &cgm.getMLIRContext(), cgm.getGlobalConstantAddressSpace()); // Create a global variable for this string // FIXME(cir): check for insertion point in module level. @@ -1987,7 +1987,7 @@ CIRGenModule::getAddrOfGlobalTemporary(const MaterializeTemporaryExpr *expr, linkage = cir::GlobalLinkageKind::InternalLinkage; } } - cir::AddressSpace targetAS = cir::toCIRAddressSpace(addrSpace); + mlir::Attribute targetAS = cir::toCIRAddressSpaceAttr(&getMLIRContext(), addrSpace); auto loc = getLoc(expr->getSourceRange()); auto gv = createGlobalOp(*this, loc, name, type, isConstant, targetAS, diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index f130b007d2b2..48b7b4c61926 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -278,8 +278,7 @@ class CIRGenModule : public CIRGenTypeCache { static cir::GlobalOp createGlobalOp( CIRGenModule &cgm, mlir::Location loc, llvm::StringRef name, mlir::Type t, - bool isConstant = false, - cir::AddressSpace addrSpace = cir::AddressSpace::Default, + bool isConstant = false, mlir::Attribute addrSpace = {}, mlir::Operation *insertPoint = nullptr, cir::GlobalLinkageKind linkage = cir::GlobalLinkageKind::ExternalLinkage); diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp index 7230cf5edc7c..ab3609d286a7 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.cpp +++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp @@ -662,8 +662,8 @@ TargetCIRGenInfo::getGlobalVarAddressSpace(CIRGenModule &CGM, } mlir::Value TargetCIRGenInfo::performAddrSpaceCast( - CIRGenFunction &CGF, mlir::Value Src, cir::AddressSpace SrcAddr, - cir::AddressSpace DestAddr, mlir::Type DestTy, bool IsNonNull) const { + CIRGenFunction &CGF, mlir::Value Src, mlir::Attribute SrcAddr, + mlir::Attribute DestAddr, mlir::Type DestTy, bool IsNonNull) const { // Since target may map different address spaces in AST to the same address // space, an address space conversion may end up as a bitcast. if (auto globalOp = Src.getDefiningOp()) diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h index 21f3b0a0637d..38c302b672b7 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -16,6 +16,7 @@ #include "ABIInfo.h" #include "CIRGenValue.h" +#include "mlir/IR/Attributes.h" #include "mlir/IR/Types.h" #include "clang/CIR/Dialect/IR/CIRAttrs.h" #include "clang/CIR/Target/AArch64.h" @@ -112,8 +113,8 @@ class TargetCIRGenInfo { /// \param DestTy is the destination pointer type. /// \param IsNonNull is the flag indicating \p V is known to be non null. virtual mlir::Value performAddrSpaceCast(CIRGenFunction &CGF, mlir::Value V, - cir::AddressSpace SrcAddr, - cir::AddressSpace DestAddr, + mlir::Attribute SrcAddr, + mlir::Attribute DestAddr, mlir::Type DestTy, bool IsNonNull = false) const; diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index 145bbdd9dd01..cafb43e9f98d 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -2405,7 +2405,7 @@ LogicalResult cir::GlobalOp::verify() { void cir::GlobalOp::build( OpBuilder &odsBuilder, OperationState &odsState, llvm::StringRef sym_name, Type sym_type, bool isConstant, cir::GlobalLinkageKind linkage, - cir::AddressSpace addrSpace, + mlir::Attribute addrSpace, function_ref ctorBuilder, function_ref dtorBuilder) { odsState.addAttribute(getSymNameAttrName(odsState.name), @@ -2420,9 +2420,8 @@ void cir::GlobalOp::build( cir::GlobalLinkageKindAttr::get(odsBuilder.getContext(), linkage); odsState.addAttribute(getLinkageAttrName(odsState.name), linkageAttr); - odsState.addAttribute( - getAddrSpaceAttrName(odsState.name), - cir::AddressSpaceAttr::get(odsBuilder.getContext(), addrSpace)); + if (addrSpace) + odsState.addAttribute(getAddrSpaceAttrName(odsState.name), addrSpace); Region *ctorRegion = odsState.addRegion(); if (ctorBuilder) { diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index 2a9427717529..e47d44c4cbd3 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -973,39 +973,30 @@ mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, if (p.parseOptionalKeyword("target_address_space").succeeded()) { unsigned val; if (p.parseLParen()) - p.emitError(loc, "expected '(' after target_address_space"); + p.emitError(loc, "expected '(' after target address space"); if (p.parseInteger(val) || p.parseRParen()) - return p.emitError(loc, "expected target_address_space value"); + return p.emitError(loc, "expected target address space value"); attr = cir::TargetAddressSpaceAttr::get(p.getContext(), val); return mlir::success(); } - // Address space is either a target address space or a regular one. - // - If it is a target address space, we expect a value to follow in the form - // of ``, where value is an integer that represents the target address - // space value. This value is kept in the address space enum as an offset - // from the maximum address space value, which is defined in - // `cir::getMaxEnumValForAddressSpace()`. This allows us to use - // the same enum for both regular and target address spaces. - // - Otherwise, we just use the parsed value. - // Try to parse language specific address space. if (p.parseOptionalKeyword("clang_address_space").succeeded()) { if (p.parseLParen()) - return p.emitError(loc, "expected '(' after clang_address_space"); + return p.emitError(loc, "expected '(' after clang address space"); mlir::FailureOr result = mlir::FieldParser::parse(p); if (mlir::failed(result) || p.parseRParen()) - return p.emitError(loc, "expected clang_address_space value"); + return p.emitError(loc, "expected clang address space keyword"); attr = cir::AddressSpaceAttr::get(p.getContext(), result.value()); return mlir::success(); } - return mlir::failure(); + return mlir::success(); } void printAddressSpaceValue(mlir::AsmPrinter &p, mlir::Attribute attr) { @@ -1029,13 +1020,15 @@ void printAddressSpaceValue(mlir::AsmPrinter &p, mlir::Attribute attr) { mlir::OptionalParseResult parseGlobalAddressSpaceValue(mlir::AsmParser &p, mlir::Attribute &attr) { - if (!parseAddressSpaceValue(p, attr)) - return mlir::failure(); + + 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::Attribute attr) { + mlir::Attribute attr) { printAddressSpaceValue(printer, attr); } @@ -1043,6 +1036,9 @@ mlir::Attribute cir::toCIRAddressSpaceAttr(mlir::MLIRContext *ctx, clang::LangAS langAS) { using clang::LangAS; + if (langAS == LangAS::Default) + return {}; + if (clang::isTargetAddressSpace(langAS)) { unsigned targetAS = clang::toTargetAddressSpace(langAS); return cir::TargetAddressSpaceAttr::get(ctx, targetAS); diff --git a/clang/test/CIR/CodeGen/CUDA/address-spaces.cu b/clang/test/CIR/CodeGen/CUDA/address-spaces.cu index b25977f1305a..4e3ba6481bbb 100644 --- a/clang/test/CIR/CodeGen/CUDA/address-spaces.cu +++ b/clang/test/CIR/CodeGen/CUDA/address-spaces.cu @@ -11,9 +11,9 @@ __global__ void fn() { j = i; } -// CIR: cir.global "private" internal dso_local addrspace(offload_local) @_ZZ2fnvE1j : !s32i +// CIR: cir.global "private" internal dso_local clang_address_space(offload_local) @_ZZ2fnvE1j : !s32i // CIR: cir.func dso_local @_Z2fnv // CIR: [[Local:%[0-9]+]] = cir.alloca !s32i, !cir.ptr, ["i", init] -// CIR: [[Shared:%[0-9]+]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr +// CIR: [[Shared:%[0-9]+]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr // CIR: [[Tmp:%[0-9]+]] = cir.load {{.*}} [[Local]] : !cir.ptr, !s32i -// CIR: cir.store{{.*}} [[Tmp]], [[Shared]] : !s32i, !cir.ptr +// CIR: cir.store{{.*}} [[Tmp]], [[Shared]] : !s32i, !cir.ptr diff --git a/clang/test/CIR/CodeGen/CUDA/global-vars.cu b/clang/test/CIR/CodeGen/CUDA/global-vars.cu index 55d4c67967fa..80e5b6003179 100644 --- a/clang/test/CIR/CodeGen/CUDA/global-vars.cu +++ b/clang/test/CIR/CodeGen/CUDA/global-vars.cu @@ -16,15 +16,15 @@ // RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s __device__ int a; -// CIR-DEVICE: cir.global external addrspace(offload_global) @a = #cir.int<0> +// CIR-DEVICE: cir.global external clang_address_space(offload_global) @a = #cir.int<0> // LLVM-DEVICE: @a = addrspace(1) externally_initialized global i32 0, align 4 // CIR-HOST: {{.*}}cir.global external @a = #cir.undef : !s32i {alignment = 4 : i64, cu.shadow_name = #cir.cu.shadow_name}{{.*}} __shared__ int shared; -// CIR-DEVICE: cir.global external addrspace(offload_local) @shared = #cir.undef +// CIR-DEVICE: cir.global external clang_address_space(offload_local) @shared = #cir.undef // LLVM-DEVICE: @shared = addrspace(3) global i32 undef, align 4 __constant__ int b; -// CIR-DEVICE: cir.global constant external addrspace(offload_constant) @b = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} +// CIR-DEVICE: cir.global constant external clang_address_space(offload_constant) @b = #cir.int<0> : !s32i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} // LLVM-DEVICE: @b = addrspace(4) externally_initialized constant i32 0, align 4 // CIR-HOST: {{.*}}cir.global external @b = #cir.undef : !s32i {alignment = 4 : i64, cu.shadow_name = #cir.cu.shadow_name}{{.*}} diff --git a/clang/test/CIR/CodeGen/CUDA/surface.cu b/clang/test/CIR/CodeGen/CUDA/surface.cu index da085137f325..3b299d9d0f29 100644 --- a/clang/test/CIR/CodeGen/CUDA/surface.cu +++ b/clang/test/CIR/CodeGen/CUDA/surface.cu @@ -22,5 +22,5 @@ struct __attribute__((device_builtin_surface_type)) surface : public surface surf; // DEVICE-LLVM: @surf = addrspace(1) externally_initialized global i64 undef, align 4 -// DEVICE-CIR: cir.global external addrspace(offload_global) @surf = #cir.undef : !s64i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} +// DEVICE-CIR: cir.global external clang_address_space(offload_global) @surf = #cir.undef : !s64i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} // HOST: @surf = global %"struct.surface" zeroinitializer, align 4 \ No newline at end of file diff --git a/clang/test/CIR/CodeGen/CUDA/texture.cu b/clang/test/CIR/CodeGen/CUDA/texture.cu index db431b658bda..bf9e56be4712 100644 --- a/clang/test/CIR/CodeGen/CUDA/texture.cu +++ b/clang/test/CIR/CodeGen/CUDA/texture.cu @@ -21,4 +21,4 @@ struct __attribute__((device_builtin_texture_type)) texture : public textureRefe texture tex; // DEVICE-LLVM: @tex = addrspace(1) externally_initialized global i64 undef, align 4 -// DEVICE-CIR: cir.global external addrspace(offload_global) @tex = #cir.undef : !s64i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} +// DEVICE-CIR: cir.global external clang_address_space(offload_global) @tex = #cir.undef : !s64i {alignment = 4 : i64, cu.externally_initialized = #cir.cu.externally_initialized} diff --git a/clang/test/CIR/CodeGen/HIP/address-spaces.cpp b/clang/test/CIR/CodeGen/HIP/address-spaces.cpp index 3ac0c30e1fe1..78a2c64ac2ab 100644 --- a/clang/test/CIR/CodeGen/HIP/address-spaces.cpp +++ b/clang/test/CIR/CodeGen/HIP/address-spaces.cpp @@ -11,9 +11,9 @@ __global__ void fn() { j = i; } -// CIR: cir.global "private" internal dso_local addrspace(offload_local) @_ZZ2fnvE1j : !s32i +// CIR: cir.global "private" internal dso_local clang_address_space(offload_local) @_ZZ2fnvE1j : !s32i // CIR: cir.func dso_local @_Z2fnv // CIR: [[Local:%[0-9]+]] = cir.alloca !s32i, !cir.ptr, ["i", init] -// CIR: [[Shared:%[0-9]+]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr +// CIR: [[Shared:%[0-9]+]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr // CIR: [[Tmp:%[0-9]+]] = cir.load {{.*}} [[Local]] : !cir.ptr, !s32i -// CIR: cir.store{{.*}} [[Tmp]], [[Shared]] : !s32i, !cir.ptr +// CIR: cir.store{{.*}} [[Tmp]], [[Shared]] : !s32i, !cir.ptr \ No newline at end of file diff --git a/clang/test/CIR/CodeGen/OpenCL/addrspace-alloca.cl b/clang/test/CIR/CodeGen/OpenCL/addrspace-alloca.cl index a1e0eb0950db..e99c57a1bb38 100644 --- a/clang/test/CIR/CodeGen/OpenCL/addrspace-alloca.cl +++ b/clang/test/CIR/CodeGen/OpenCL/addrspace-alloca.cl @@ -4,30 +4,30 @@ // RUN: FileCheck --input-file=%t.ll %s --check-prefix=LLVM -// CIR: cir.func @func(%arg0: !cir.ptr +// CIR: cir.func @func(%arg0: !cir.ptr // LLVM: @func(ptr addrspace(3) kernel void func(local int *p) { - // CIR-NEXT: %[[#ALLOCA_P:]] = cir.alloca !cir.ptr, !cir.ptr, addrspace(offload_private)>, ["p", init] {alignment = 8 : i64} + // CIR-NEXT: %[[#ALLOCA_P:]] = cir.alloca !cir.ptr, !cir.ptr, clang_address_space(offload_private)>, ["p", init] {alignment = 8 : i64} // LLVM-NEXT: %[[#ALLOCA_P:]] = alloca ptr addrspace(3), i64 1, align 8 int x; - // CIR-NEXT: %[[#ALLOCA_X:]] = cir.alloca !s32i, !cir.ptr, ["x"] {alignment = 4 : i64} + // CIR-NEXT: %[[#ALLOCA_X:]] = cir.alloca !s32i, !cir.ptr, ["x"] {alignment = 4 : i64} // LLVM-NEXT: %[[#ALLOCA_X:]] = alloca i32, i64 1, align 4 global char *b; - // CIR-NEXT: %[[#ALLOCA_B:]] = cir.alloca !cir.ptr, !cir.ptr, addrspace(offload_private)>, ["b"] {alignment = 8 : i64} + // CIR-NEXT: %[[#ALLOCA_B:]] = cir.alloca !cir.ptr, !cir.ptr, clang_address_space(offload_private)>, ["b"] {alignment = 8 : i64} // LLVM-NEXT: %[[#ALLOCA_B:]] = alloca ptr addrspace(1), i64 1, align 8 private int *ptr; - // CIR-NEXT: %[[#ALLOCA_PTR:]] = cir.alloca !cir.ptr, !cir.ptr, addrspace(offload_private)>, ["ptr"] {alignment = 8 : i64} + // CIR-NEXT: %[[#ALLOCA_PTR:]] = cir.alloca !cir.ptr, !cir.ptr, clang_address_space(offload_private)>, ["ptr"] {alignment = 8 : i64} // LLVM-NEXT: %[[#ALLOCA_PTR:]] = alloca ptr, i64 1, align 8 // Store of the argument `p` - // CIR-NEXT: cir.store{{.*}} %arg0, %[[#ALLOCA_P]] : !cir.ptr, !cir.ptr, addrspace(offload_private)> + // CIR-NEXT: cir.store{{.*}} %arg0, %[[#ALLOCA_P]] : !cir.ptr, !cir.ptr, clang_address_space(offload_private)> // LLVM-NEXT: store ptr addrspace(3) %{{[0-9]+}}, ptr %[[#ALLOCA_P]], align 8 ptr = &x; - // CIR-NEXT: cir.store{{.*}} %[[#ALLOCA_X]], %[[#ALLOCA_PTR]] : !cir.ptr, !cir.ptr, addrspace(offload_private)> + // CIR-NEXT: cir.store{{.*}} %[[#ALLOCA_X]], %[[#ALLOCA_PTR]] : !cir.ptr, !cir.ptr, clang_address_space(offload_private)> // LLVM-NEXT: store ptr %[[#ALLOCA_X]], ptr %[[#ALLOCA_PTR]] return; diff --git a/clang/test/CIR/CodeGen/OpenCL/array-decay.cl b/clang/test/CIR/CodeGen/OpenCL/array-decay.cl index 9ba283587309..f70dad17c470 100644 --- a/clang/test/CIR/CodeGen/OpenCL/array-decay.cl +++ b/clang/test/CIR/CodeGen/OpenCL/array-decay.cl @@ -9,8 +9,8 @@ kernel void func1(global int *data) { local int arr[32]; local int *ptr = arr; - // CIR: cir.cast array_to_ptrdecay %{{[0-9]+}} : !cir.ptr, addrspace(offload_local)> -> !cir.ptr - // CIR-NEXT: cir.store{{.*}} %{{[0-9]+}}, %{{[0-9]+}} : !cir.ptr, !cir.ptr, addrspace(offload_private)> + // CIR: cir.cast array_to_ptrdecay %{{[0-9]+}} : !cir.ptr, clang_address_space(offload_local)> -> !cir.ptr + // CIR-NEXT: cir.store{{.*}} %{{[0-9]+}}, %{{[0-9]+}} : !cir.ptr, !cir.ptr, clang_address_space(offload_private)> // LLVM: store ptr addrspace(3) @func1.arr, ptr %{{[0-9]+}} } @@ -19,7 +19,7 @@ kernel void func1(global int *data) { // LLVM: @func2 kernel void func2(global int *data) { private int arr[32] = {data[2]}; - // CIR: %{{[0-9]+}} = cir.get_element %{{[0-9]+}}[%{{[0-9]+}}] : (!cir.ptr, addrspace(offload_private)>, !s32i) -> !cir.ptr + // CIR: %{{[0-9]+}} = cir.get_element %{{[0-9]+}}[%{{[0-9]+}}] : (!cir.ptr, clang_address_space(offload_private)>, !s32i) -> !cir.ptr // LLVM: %{{[0-9]+}} = getelementptr [32 x i32], ptr %3, i32 0, i64 0 } diff --git a/clang/test/CIR/CodeGen/OpenCL/global.cl b/clang/test/CIR/CodeGen/OpenCL/global.cl index ef5c63e58f83..007f27f23aaa 100644 --- a/clang/test/CIR/CodeGen/OpenCL/global.cl +++ b/clang/test/CIR/CodeGen/OpenCL/global.cl @@ -4,23 +4,23 @@ // RUN: FileCheck --input-file=%t.ll %s --check-prefix=LLVM global int a = 13; -// CIR-DAG: cir.global external addrspace(offload_global) @a = #cir.int<13> : !s32i +// CIR-DAG: cir.global external clang_address_space(offload_global) @a = #cir.int<13> : !s32i // LLVM-DAG: @a = addrspace(1) global i32 13 global int b = 15; -// CIR-DAG: cir.global external addrspace(offload_global) @b = #cir.int<15> : !s32i +// CIR-DAG: cir.global external clang_address_space(offload_global) @b = #cir.int<15> : !s32i // LLVM-DAG: @b = addrspace(1) global i32 15 constant int c[2] = {18, 21}; -// CIR-DAG: cir.global constant {{.*}}addrspace(offload_constant) {{.*}}@c +// CIR-DAG: cir.global constant {{.*}}clang_address_space(offload_constant) {{.*}}@c // LLVM-DAG: @c = addrspace(2) constant kernel void test_get_global() { a = b; - // CIR: %[[#ADDRB:]] = cir.get_global @b : !cir.ptr - // CIR-NEXT: %[[#LOADB:]] = cir.load{{.*}} %[[#ADDRB]] : !cir.ptr, !s32i - // CIR-NEXT: %[[#ADDRA:]] = cir.get_global @a : !cir.ptr - // CIR-NEXT: cir.store{{.*}} %[[#LOADB]], %[[#ADDRA]] : !s32i, !cir.ptr + // CIR: %[[#ADDRB:]] = cir.get_global @b : !cir.ptr + // CIR-NEXT: %[[#LOADB:]] = cir.load{{.*}} %[[#ADDRB]] : !cir.ptr, !s32i + // CIR-NEXT: %[[#ADDRA:]] = cir.get_global @a : !cir.ptr + // CIR-NEXT: cir.store{{.*}} %[[#LOADB]], %[[#ADDRA]] : !s32i, !cir.ptr // LLVM: %[[#LOADB:]] = load i32, ptr addrspace(1) @b, align 4 // LLVM-NEXT: store i32 %[[#LOADB]], ptr addrspace(1) @a, align 4 diff --git a/clang/test/CIR/CodeGen/OpenCL/printf.cl b/clang/test/CIR/CodeGen/OpenCL/printf.cl index b539fce01c2b..f2305f329943 100644 --- a/clang/test/CIR/CodeGen/OpenCL/printf.cl +++ b/clang/test/CIR/CodeGen/OpenCL/printf.cl @@ -28,8 +28,8 @@ kernel void test_printf_float2(float2 arg) { printf("%v2hlf", arg); } // CIR-ALL-LABEL: @test_printf_float2( -// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr, !cir.vector) -> !s32i cc(spir_function) -// CIR-NOFP64:%{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr, !cir.vector) -> !s32i cc(spir_function) +// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr, !cir.vector) -> !s32i cc(spir_function) +// CIR-NOFP64:%{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr, !cir.vector) -> !s32i cc(spir_function) // LLVM-ALL-LABEL: @test_printf_float2( // LLVM-FP64: %{{.+}} = call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str, <2 x float> %{{.*}}) // LLVM-NOFP64: call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str, <2 x float> %{{.*}}) @@ -38,8 +38,8 @@ kernel void test_printf_half2(half2 arg) { printf("%v2hf", arg); } // CIR-ALL-LABEL: @test_printf_half2( -// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr, !cir.vector) -> !s32i cc(spir_function) -// CIR-NOFP64:%{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr, !cir.vector) -> !s32i cc(spir_function) +// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr, !cir.vector) -> !s32i cc(spir_function) +// CIR-NOFP64:%{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr, !cir.vector) -> !s32i cc(spir_function) // LLVM-ALL-LABEL: @test_printf_half2( // LLVM-FP64: %{{.+}} = call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str.1, <2 x half> %{{.*}}) // LLVM-NOFP64: %{{.+}} = call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str.1, <2 x half> %{{.*}}) @@ -49,7 +49,7 @@ kernel void test_printf_double2(double2 arg) { printf("%v2lf", arg); } // CIR-FP64-LABEL: @test_printf_double2( -// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr, !cir.vector) -> !s32i cc(spir_function) +// CIR-FP64: %{{.+}} = cir.call @printf(%{{.+}}, %{{.+}}) : (!cir.ptr, !cir.vector) -> !s32i cc(spir_function) // LLVM-FP64-LABEL: @test_printf_double2( // LLVM-FP64: call spir_func i32 (ptr addrspace(2), ...) @{{.*}}printf{{.*}}(ptr addrspace(2) @.str.2, <2 x double> %{{.*}}) #endif diff --git a/clang/test/CIR/CodeGen/OpenCL/static-vardecl.cl b/clang/test/CIR/CodeGen/OpenCL/static-vardecl.cl index 0a8e03bbfc9d..579f96ff1b8a 100644 --- a/clang/test/CIR/CodeGen/OpenCL/static-vardecl.cl +++ b/clang/test/CIR/CodeGen/OpenCL/static-vardecl.cl @@ -5,19 +5,19 @@ kernel void test_static(int i) { static global int b = 15; - // CIR-DAG: cir.global "private" internal dso_local addrspace(offload_global) @test_static.b = #cir.int<15> : !s32i {alignment = 4 : i64} + // CIR-DAG: cir.global "private" internal dso_local clang_address_space(offload_global) @test_static.b = #cir.int<15> : !s32i {alignment = 4 : i64} // LLVM-DAG: @test_static.b = internal addrspace(1) global i32 15 local int c; - // CIR-DAG: cir.global "private" internal dso_local addrspace(offload_local) @test_static.c : !s32i {alignment = 4 : i64} + // CIR-DAG: cir.global "private" internal dso_local clang_address_space(offload_local) @test_static.c : !s32i {alignment = 4 : i64} // LLVM-DAG: @test_static.c = internal addrspace(3) global i32 undef - // CIR-DAG: %[[#ADDRB:]] = cir.get_global @test_static.b : !cir.ptr - // CIR-DAG: %[[#ADDRC:]] = cir.get_global @test_static.c : !cir.ptr + // CIR-DAG: %[[#ADDRB:]] = cir.get_global @test_static.b : !cir.ptr + // CIR-DAG: %[[#ADDRC:]] = cir.get_global @test_static.c : !cir.ptr c = b; - // CIR: %[[#LOADB:]] = cir.load{{.*}} %[[#ADDRB]] : !cir.ptr, !s32i - // CIR-NEXT: cir.store{{.*}} %[[#LOADB]], %[[#ADDRC]] : !s32i, !cir.ptr + // CIR: %[[#LOADB:]] = cir.load{{.*}} %[[#ADDRB]] : !cir.ptr, !s32i + // CIR-NEXT: cir.store{{.*}} %[[#LOADB]], %[[#ADDRC]] : !s32i, !cir.ptr // LLVM: %[[#LOADB:]] = load i32, ptr addrspace(1) @test_static.b, align 4 // LLVM-NEXT: store i32 %[[#LOADB]], ptr addrspace(3) @test_static.c, align 4 diff --git a/clang/test/CIR/CodeGen/OpenCL/str_literals.cl b/clang/test/CIR/CodeGen/OpenCL/str_literals.cl index 4f1842b3d152..3e0358677db4 100644 --- a/clang/test/CIR/CodeGen/OpenCL/str_literals.cl +++ b/clang/test/CIR/CodeGen/OpenCL/str_literals.cl @@ -6,10 +6,10 @@ __constant char *__constant x = "hello world"; __constant char *__constant y = "hello world"; -// CIR: cir.global{{.*}} constant {{.*}}addrspace(offload_constant) @".str" = #cir.const_array<"hello world\00" : !cir.array> : !cir.array -// CIR: cir.global{{.*}} constant {{.*}}addrspace(offload_constant) @x = #cir.global_view<@".str"> : !cir.ptr -// CIR: cir.global{{.*}} constant {{.*}}addrspace(offload_constant) @y = #cir.global_view<@".str"> : !cir.ptr -// CIR: cir.global{{.*}} constant {{.*}}addrspace(offload_constant) @".str.1" = #cir.const_array<"f\00" : !cir.array> : !cir.array +// CIR: cir.global{{.*}} constant {{.*}}clang_address_space(offload_constant) @".str" = #cir.const_array<"hello world\00" : !cir.array> : !cir.array +// CIR: cir.global{{.*}} constant {{.*}}clang_address_space(offload_constant) @x = #cir.global_view<@".str"> : !cir.ptr +// CIR: cir.global{{.*}} constant {{.*}}clang_address_space(offload_constant) @y = #cir.global_view<@".str"> : !cir.ptr +// CIR: cir.global{{.*}} constant {{.*}}clang_address_space(offload_constant) @".str.1" = #cir.const_array<"f\00" : !cir.array> : !cir.array // LLVM: addrspace(2) constant{{.*}}"hello world\00" // LLVM-NOT: addrspace(2) constant // LLVM: @x = {{(dso_local )?}}addrspace(2) constant ptr addrspace(2) @@ -17,7 +17,7 @@ __constant char *__constant y = "hello world"; // LLVM: addrspace(2) constant{{.*}}"f\00" void f() { - // CIR: cir.store{{.*}} %{{.*}}, %{{.*}} : !cir.ptr, !cir.ptr, addrspace(offload_private)> + // CIR: cir.store{{.*}} %{{.*}}, %{{.*}} : !cir.ptr, !cir.ptr, clang_address_space(offload_private)> // LLVM: store ptr addrspace(2) {{.*}}, ptr constant const char *f3 = __func__; } diff --git a/clang/test/CIR/CodeGen/address-space-conversion.cpp b/clang/test/CIR/CodeGen/address-space-conversion.cpp index ce26ef69ebf8..55c857eb08f4 100644 --- a/clang/test/CIR/CodeGen/address-space-conversion.cpp +++ b/clang/test/CIR/CodeGen/address-space-conversion.cpp @@ -14,9 +14,9 @@ using ri2_t = int __attribute__((address_space(2))) &; void test_ptr() { pi1_t ptr1; pi2_t ptr2 = (pi2_t)ptr1; - // CIR: %[[#PTR1:]] = cir.load{{.*}} %{{[0-9]+}} : !cir.ptr)>>, !cir.ptr)> - // CIR-NEXT: %[[#CAST:]] = cir.cast address_space %[[#PTR1]] : !cir.ptr)> -> !cir.ptr)> - // CIR-NEXT: cir.store{{.*}} %[[#CAST]], %{{[0-9]+}} : !cir.ptr)>, !cir.ptr)>> + // CIR: %[[#PTR1:]] = cir.load{{.*}} %{{[0-9]+}} : !cir.ptr>, !cir.ptr + // CIR-NEXT: %[[#CAST:]] = cir.cast address_space %[[#PTR1]] : !cir.ptr -> !cir.ptr + // CIR-NEXT: cir.store{{.*}} %[[#CAST]], %{{[0-9]+}} : !cir.ptr, !cir.ptr> // LLVM: %[[#PTR1:]] = load ptr addrspace(1), ptr %{{[0-9]+}}, align 8 // LLVM-NEXT: %[[#CAST:]] = addrspacecast ptr addrspace(1) %[[#PTR1]] to ptr addrspace(2) @@ -29,11 +29,11 @@ void test_ref() { pi1_t ptr; ri1_t ref1 = *ptr; ri2_t ref2 = (ri2_t)ref1; - // CIR: %[[#DEREF:]] = cir.load deref{{.*}} %{{[0-9]+}} : !cir.ptr)>>, !cir.ptr)> - // CIR-NEXT: cir.store{{.*}} %[[#DEREF]], %[[#ALLOCAREF1:]] : !cir.ptr)>, !cir.ptr)>> - // CIR-NEXT: %[[#REF1:]] = cir.load{{.*}} %[[#ALLOCAREF1]] : !cir.ptr)>>, !cir.ptr)> - // CIR-NEXT: %[[#CAST:]] = cir.cast address_space %[[#REF1]] : !cir.ptr)> -> !cir.ptr)> - // CIR-NEXT: cir.store{{.*}} %[[#CAST]], %{{[0-9]+}} : !cir.ptr)>, !cir.ptr)>> + // CIR: %[[#DEREF:]] = cir.load deref{{.*}} %{{[0-9]+}} : !cir.ptr>, !cir.ptr + // CIR-NEXT: cir.store{{.*}} %[[#DEREF]], %[[#ALLOCAREF1:]] : !cir.ptr, !cir.ptr> + // CIR-NEXT: %[[#REF1:]] = cir.load{{.*}} %[[#ALLOCAREF1]] : !cir.ptr>, !cir.ptr + // CIR-NEXT: %[[#CAST:]] = cir.cast address_space %[[#REF1]] : !cir.ptr -> !cir.ptr + // CIR-NEXT: cir.store{{.*}} %[[#CAST]], %{{[0-9]+}} : !cir.ptr, !cir.ptr> // LLVM: %[[#DEREF:]] = load ptr addrspace(1), ptr %{{[0-9]+}}, align 8 // LLVM-NEXT: store ptr addrspace(1) %[[#DEREF]], ptr %[[#ALLOCAREF1:]], align 8 @@ -47,10 +47,10 @@ void test_ref() { void test_nullptr() { constexpr pi1_t null1 = nullptr; pi2_t ptr = (pi2_t)null1; - // CIR: %[[#NULL1:]] = cir.const #cir.ptr : !cir.ptr)> - // CIR-NEXT: cir.store{{.*}} %[[#NULL1]], %{{[0-9]+}} : !cir.ptr)>, !cir.ptr)>> - // CIR-NEXT: %[[#NULL2:]] = cir.const #cir.ptr : !cir.ptr)> - // CIR-NEXT: cir.store{{.*}} %[[#NULL2]], %{{[0-9]+}} : !cir.ptr)>, !cir.ptr)>> + // CIR: %[[#NULL1:]] = cir.const #cir.ptr : !cir.ptr + // CIR-NEXT: cir.store{{.*}} %[[#NULL1]], %{{[0-9]+}} : !cir.ptr, !cir.ptr> + // CIR-NEXT: %[[#NULL2:]] = cir.const #cir.ptr : !cir.ptr + // CIR-NEXT: cir.store{{.*}} %[[#NULL2]], %{{[0-9]+}} : !cir.ptr, !cir.ptr> // LLVM: store ptr addrspace(1) null, ptr %{{[0-9]+}}, align 8 // LLVM-NEXT: store ptr addrspace(2) null, ptr %{{[0-9]+}}, align 8 @@ -58,9 +58,9 @@ void test_nullptr() { void test_side_effect(pi1_t b) { pi2_t p = (pi2_t)(*b++, (int*)0); - // CIR: %{{[0-9]+}} = cir.ptr_stride %{{[0-9]+}}, %{{[0-9]+}} : (!cir.ptr)>, !s32i) -> !cir.ptr)> - // CIR: %[[#CAST:]] = cir.const #cir.ptr : !cir.ptr)> - // CIR-NEXT: cir.store{{.*}} %[[#CAST]], %{{[0-9]+}} : !cir.ptr)>, !cir.ptr)>> + // CIR: %{{[0-9]+}} = cir.ptr_stride %{{[0-9]+}}, %{{[0-9]+}} : (!cir.ptr, !s32i) -> !cir.ptr + // CIR: %[[#CAST:]] = cir.const #cir.ptr : !cir.ptr + // CIR-NEXT: cir.store{{.*}} %[[#CAST]], %{{[0-9]+}} : !cir.ptr, !cir.ptr> // LLVM: %{{[0-9]+}} = getelementptr i32, ptr addrspace(1) %{{[0-9]+}}, i64 1 // LLVM: store ptr addrspace(2) null, ptr %{{[0-9]+}}, align 8 diff --git a/clang/test/CIR/CodeGen/address-space.c b/clang/test/CIR/CodeGen/address-space.c index d131fb84d98d..c2776a9ca5ab 100644 --- a/clang/test/CIR/CodeGen/address-space.c +++ b/clang/test/CIR/CodeGen/address-space.c @@ -3,13 +3,13 @@ // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm %s -o %t.ll // RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM -// CIR: cir.func dso_local {{@.*foo.*}}(%arg0: !cir.ptr)> +// CIR: cir.func dso_local {{@.*foo.*}}(%arg0: !cir.ptr // LLVM: define dso_local void @foo(ptr addrspace(1) %0) void foo(int __attribute__((address_space(1))) *arg) { return; } -// CIR: cir.func dso_local {{@.*bar.*}}(%arg0: !cir.ptr)> +// CIR: cir.func dso_local {{@.*bar.*}}(%arg0: !cir.ptr // LLVM: define dso_local void @bar(ptr %0) void bar(int __attribute__((address_space(0))) *arg) { return; diff --git a/clang/test/CIR/IR/address-space.cir b/clang/test/CIR/IR/address-space.cir index 234d03fa19b8..41f8ca9acfe4 100644 --- a/clang/test/CIR/IR/address-space.cir +++ b/clang/test/CIR/IR/address-space.cir @@ -3,8 +3,8 @@ !s32i = !cir.int module { - // CHECK: @test_format1(%arg0: !cir.ptr)>) - cir.func @test_format1(%arg0: !cir.ptr)>) { + // CHECK: @test_format1(%arg0: !cir.ptr) + cir.func @test_format1(%arg0: !cir.ptr) { cir.return } @@ -13,28 +13,28 @@ module { cir.return } - // CHECK: @test_format3(%arg0: !cir.ptr) - cir.func @test_format3(%arg0: !cir.ptr) { + // CHECK: @test_format3(%arg0: !cir.ptr) + cir.func @test_format3(%arg0: !cir.ptr) { cir.return } - // CHECK: @test_format4(%arg0: !cir.ptr) - cir.func @test_format4(%arg0: !cir.ptr) { + // CHECK: @test_format4(%arg0: !cir.ptr) + cir.func @test_format4(%arg0: !cir.ptr) { cir.return } - // CHECK: @test_format5(%arg0: !cir.ptr) - cir.func @test_format5(%arg0: !cir.ptr) { + // CHECK: @test_format5(%arg0: !cir.ptr) + cir.func @test_format5(%arg0: !cir.ptr) { cir.return } - // CHECK: @test_format6(%arg0: !cir.ptr) - cir.func @test_format6(%arg0: !cir.ptr) { + // CHECK: @test_format6(%arg0: !cir.ptr) + cir.func @test_format6(%arg0: !cir.ptr) { cir.return } - // CHECK: @test_format7(%arg0: !cir.ptr) - cir.func @test_format7(%arg0: !cir.ptr) { + // CHECK: @test_format7(%arg0: !cir.ptr) + cir.func @test_format7(%arg0: !cir.ptr) { cir.return } } diff --git a/clang/test/CIR/IR/cast.cir b/clang/test/CIR/IR/cast.cir index c3b26cf79756..f9adde6d969a 100644 --- a/clang/test/CIR/IR/cast.cir +++ b/clang/test/CIR/IR/cast.cir @@ -17,7 +17,7 @@ module { } cir.func @addrspace_cast(%arg0: !cir.ptr) { - %0 = cir.cast address_space %arg0 : !cir.ptr -> !cir.ptr)> + %0 = cir.cast address_space %arg0 : !cir.ptr -> !cir.ptr cir.return } } @@ -30,4 +30,4 @@ module { // CHECK: %0 = cir.cast bitcast %arg0 : !cir.ptr -> !cir.ptr // CHECK: cir.func @addrspace_cast -// CHECK: %0 = cir.cast address_space %arg0 : !cir.ptr -> !cir.ptr)> +// CHECK: %0 = cir.cast address_space %arg0 : !cir.ptr -> !cir.ptr diff --git a/clang/test/CIR/IR/global.cir b/clang/test/CIR/IR/global.cir index 4020ddcaadf4..d5f0e341cd72 100644 --- a/clang/test/CIR/IR/global.cir +++ b/clang/test/CIR/IR/global.cir @@ -69,9 +69,9 @@ module { cir.return } - cir.global external addrspace(offload_global) @addrspace1 = #cir.int<1> : !s32i - cir.global "private" internal addrspace(offload_local) @addrspace2 : !s32i - cir.global external addrspace(target<1>) @addrspace3 = #cir.int<3> : !s32i + cir.global external clang_address_space(offload_global) @addrspace1 = #cir.int<1> : !s32i + cir.global "private" internal clang_address_space(offload_local) @addrspace2 : !s32i + cir.global external target_address_space(1) @addrspace3 = #cir.int<3> : !s32i } // CHECK: cir.global external @a = #cir.int<3> : !s32i @@ -108,6 +108,6 @@ module { // CHECK: cir.return // CHECK: } -// CHECK: cir.global external addrspace(offload_global) @addrspace1 = #cir.int<1> : !s32i -// CHECK: cir.global "private" internal addrspace(offload_local) @addrspace2 : !s32i -// CHECK: cir.global external addrspace(target<1>) @addrspace3 = #cir.int<3> : !s32i +// CHECK: cir.global external clang_address_space(offload_global) @addrspace1 = #cir.int<1> : !s32i +// CHECK: cir.global "private" internal clang_address_space(offload_local) @addrspace2 : !s32i +// CHECK: cir.global external target_address_space(1) @addrspace3 = #cir.int<3> : !s32i diff --git a/clang/test/CIR/IR/invalid.cir b/clang/test/CIR/IR/invalid.cir index f9a7bb92c656..0b30ee2b3757 100644 --- a/clang/test/CIR/IR/invalid.cir +++ b/clang/test/CIR/IR/invalid.cir @@ -302,16 +302,16 @@ cir.func @cast24(%p : !u32i) { !u32i = !cir.int !u64i = !cir.int -cir.func @cast25(%p : !cir.ptr)>) { - %0 = cir.cast address_space %p : !cir.ptr)> -> !cir.ptr)> // expected-error {{requires two types differ in addrspace only}} +cir.func @cast25(%p : !cir.ptr) { + %0 = cir.cast address_space %p : !cir.ptr -> !cir.ptr // expected-error {{requires two types differ in addrspace only}} cir.return } // ----- !u64i = !cir.int -cir.func @cast26(%p : !cir.ptr)>) { - %0 = cir.cast address_space %p : !cir.ptr)> -> !u64i // expected-error {{requires !cir.ptr type for source and result}} +cir.func @cast26(%p : !cir.ptr) { + %0 = cir.cast address_space %p : !cir.ptr -> !u64i // expected-error {{requires !cir.ptr type for source and result}} cir.return } @@ -319,7 +319,7 @@ cir.func @cast26(%p : !cir.ptr)>) { !u64i = !cir.int cir.func @cast27(%p : !u64i) { - %0 = cir.cast address_space %p : !u64i -> !cir.ptr)> // expected-error {{requires !cir.ptr type for source and result}} + %0 = cir.cast address_space %p : !u64i -> !cir.ptr // expected-error {{requires !cir.ptr type for source and result}} cir.return } @@ -1232,9 +1232,9 @@ cir.func @bad_goto() -> () { // ----- !u64i = !cir.int -// expected-error@below {{expected address space keyword}} -// expected-error@below {{expected keyword for address space kind}} -cir.func @address_space1(%p : !cir.ptr) { +// expected-error@below {{expected clang address space keyword}} +// expected-error@below {{expected keyword for clang address space kind}} +cir.func @address_space1(%p : !cir.ptr) { cir.return } @@ -1243,24 +1243,17 @@ cir.func @address_space1(%p : !cir.ptr) { !u64i = !cir.int // expected-error@below {{expected target address space value}} // expected-error@below {{expected integer value}} -cir.func @address_space2(%p : !cir.ptr)>) { +cir.func @address_space2(%p : !cir.ptr) { cir.return } // ----- -!u64i = !cir.int -// expected-error@below {{expected '<'}} -cir.func @address_space3(%p : !cir.ptr) { - cir.return -} - -// ----- !u64i = !cir.int -// expected-error@below {{expected one of [default, offload_private, offload_local, offload_global, offload_constant, offload_generic, target] for address space kind, got: foobar}} -// expected-error@below {{expected address space keyword}} -cir.func @address_space4(%p : !cir.ptr) { +// expected-error@below {{expected one of [default, offload_private, offload_local, offload_global, offload_constant, offload_generic, target] for clang address space kind, got: foobar}} +// expected-error@below {{expected clang address space keyword}} +cir.func @address_space4(%p : !cir.ptr) { cir.return } @@ -1369,7 +1362,7 @@ module { !s32i = !cir.int module { - cir.global external addrspace(offload_global) @gv = #cir.int<0> : !s32i + cir.global external clang_address_space(offload_global) @gv = #cir.int<0> : !s32i cir.func @test_get_global() { // expected-error@+1 {{'cir.get_global' op result type address space does not match the address space of the global @gv}} @@ -1384,9 +1377,9 @@ module { module { cir.func @array_to_ptrdecay_addrspace() { - %0 = cir.alloca !cir.array, !cir.ptr, addrspace(offload_private)>, ["x", init] + %0 = cir.alloca !cir.array, !cir.ptr, clang_address_space(offload_private)>, ["x", init] // expected-error@+1 {{requires same address space for source and result}} - %1 = cir.cast array_to_ptrdecay %0 : !cir.ptr, addrspace(offload_private)> -> !cir.ptr + %1 = cir.cast array_to_ptrdecay %0 : !cir.ptr, clang_address_space(offload_private)> -> !cir.ptr cir.return } } @@ -1415,7 +1408,7 @@ module { cir.func @test_bitcast_addrspace() { %0 = cir.alloca !s32i, !cir.ptr, ["tmp"] {alignment = 4 : i64} // expected-error@+1 {{'cir.cast' op result type address space does not match the address space of the operand}} - %1 = cir.cast bitcast %0 : !cir.ptr -> !cir.ptr + %1 = cir.cast bitcast %0 : !cir.ptr -> !cir.ptr } } diff --git a/clang/test/CIR/Lowering/address-space.cir b/clang/test/CIR/Lowering/address-space.cir index abe693a1cf51..6bdb773a26bd 100644 --- a/clang/test/CIR/Lowering/address-space.cir +++ b/clang/test/CIR/Lowering/address-space.cir @@ -4,13 +4,13 @@ !s32i = !cir.int module { - cir.global external addrspace(offload_global) @addrspace1 = #cir.int<1> : !s32i + cir.global external clang_address_space(offload_global) @addrspace1 = #cir.int<1> : !s32i // LLVM: @addrspace1 = addrspace(1) global i32 - cir.global "private" internal addrspace(offload_local) @addrspace2 : !s32i + cir.global "private" internal clang_address_space(offload_local) @addrspace2 : !s32i // LLVM: @addrspace2 = internal addrspace(3) global i32 undef - cir.global external addrspace(target<7>) @addrspace3 = #cir.int<3> : !s32i + cir.global external target_address_space(7) @addrspace3 = #cir.int<3> : !s32i // LLVM: @addrspace3 = addrspace(7) global i32 // LLVM: define void @foo(ptr %0) @@ -21,30 +21,30 @@ module { } // LLVM: define void @bar(ptr addrspace(1) %0) - cir.func @bar(%arg0: !cir.ptr)>) { + cir.func @bar(%arg0: !cir.ptr) { // LLVM-NEXT: alloca ptr addrspace(1) - %0 = cir.alloca !cir.ptr)>, !cir.ptr)>>, ["arg", init] {alignment = 8 : i64} + %0 = cir.alloca !cir.ptr, !cir.ptr>, ["arg", init] {alignment = 8 : i64} cir.return } // LLVM: define void @baz(ptr %0) - cir.func @baz(%arg0: !cir.ptr)>) { + cir.func @baz(%arg0: !cir.ptr) { // LLVM-NEXT: alloca ptr, - %0 = cir.alloca !cir.ptr)>, !cir.ptr)>>, ["arg", init] {alignment = 8 : i64} + %0 = cir.alloca !cir.ptr, !cir.ptr>, ["arg", init] {alignment = 8 : i64} cir.return } // LLVM: define void @test_lower_offload_as() cir.func @test_lower_offload_as() { - %0 = cir.alloca !cir.ptr, !cir.ptr>, ["arg0", init] {alignment = 8 : i64} + %0 = cir.alloca !cir.ptr, !cir.ptr>, ["arg0", init] {alignment = 8 : i64} // LLVM-NEXT: alloca ptr, - %1 = cir.alloca !cir.ptr, !cir.ptr>, ["arg1", init] {alignment = 8 : i64} + %1 = cir.alloca !cir.ptr, !cir.ptr>, ["arg1", init] {alignment = 8 : i64} // LLVM-NEXT: alloca ptr addrspace(1), - %2 = cir.alloca !cir.ptr, !cir.ptr>, ["arg2", init] {alignment = 8 : i64} + %2 = cir.alloca !cir.ptr, !cir.ptr>, ["arg2", init] {alignment = 8 : i64} // LLVM-NEXT: alloca ptr addrspace(2), - %3 = cir.alloca !cir.ptr, !cir.ptr>, ["arg3", init] {alignment = 8 : i64} + %3 = cir.alloca !cir.ptr, !cir.ptr>, ["arg3", init] {alignment = 8 : i64} // LLVM-NEXT: alloca ptr addrspace(3), - %4 = cir.alloca !cir.ptr, !cir.ptr>, ["arg4", init] {alignment = 8 : i64} + %4 = cir.alloca !cir.ptr, !cir.ptr>, ["arg4", init] {alignment = 8 : i64} // LLVM-NEXT: alloca ptr addrspace(4), cir.return } diff --git a/clang/test/CIR/Transforms/merge-cleanups.cir b/clang/test/CIR/Transforms/merge-cleanups.cir index f7888189aee4..3e7fee9964e0 100644 --- a/clang/test/CIR/Transforms/merge-cleanups.cir +++ b/clang/test/CIR/Transforms/merge-cleanups.cir @@ -131,11 +131,11 @@ module { // Should remove redundant address space casts. // CHECK-LABEL: @addrspacecastfold - // CHECK: %[[ARG0:.+]]: !cir.ptr)> - // CHECK: cir.return %[[ARG0]] : !cir.ptr)> - cir.func @addrspacecastfold(%arg0: !cir.ptr)>) -> !cir.ptr)> { - %0 = cir.cast address_space %arg0 : !cir.ptr)> -> !cir.ptr)> - cir.return %0 : !cir.ptr)> + // CHECK: %[[ARG0:.+]]: !cir.ptr + // CHECK: cir.return %[[ARG0]] : !cir.ptr + cir.func @addrspacecastfold(%arg0: !cir.ptr) -> !cir.ptr { + %0 = cir.cast address_space %arg0 : !cir.ptr -> !cir.ptr + cir.return %0 : !cir.ptr } // Should remove scope with only yield From 06d4c1b485b2d87740a36dbe06b1e6f66eb35383 Mon Sep 17 00:00:00 2001 From: David Rivera Date: Wed, 19 Nov 2025 18:37:03 -0500 Subject: [PATCH 06/13] Remove Target AS handling from `cir::AddressSpace` --- .../CIR/Dialect/Builder/CIRBaseBuilder.h | 15 +++-------- .../include/clang/CIR/Dialect/IR/CIRAttrs.td | 12 --------- .../clang/CIR/Dialect/IR/CIREnumAttr.td | 3 +-- clang/include/clang/CIR/Dialect/IR/CIRTypes.h | 27 ------------------- clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 7 ++--- clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 3 ++- clang/lib/CIR/CodeGen/CIRGenTypeCache.h | 5 ++-- clang/lib/CIR/CodeGen/TargetInfo.cpp | 5 ++-- clang/lib/CIR/CodeGen/TargetInfo.h | 4 +-- clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 1 - clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 24 ++++++++--------- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 18 +++++-------- clang/test/CIR/IR/invalid.cir | 2 +- 13 files changed, 36 insertions(+), 90 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h index 6b4328b9aac2..adc5e8dd0944 100644 --- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h +++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h @@ -108,21 +108,12 @@ class CIRBaseBuilderTy : public mlir::OpBuilder { /// Create a pointer type with an address space attribute. cir::PointerType getPointerTo(mlir::Type ty, mlir::Attribute addrSpaceAttr) { + assert(mlir::isa(addrSpaceAttr) || + mlir::isa(addrSpaceAttr) && + "expected address space attribute"); return cir::PointerType::get(ty, addrSpaceAttr); } - cir::PointerType getPointerTo(mlir::Type ty, cir::AddressSpace addrSpace) { - if (addrSpace == cir::AddressSpace::Default) - return getPointerTo(ty); - if (cir::isTargetAddressSpace(addrSpace)) { - unsigned targetAS = cir::getTargetAddressSpaceValueFromCIRAS(addrSpace); - auto attr = cir::TargetAddressSpaceAttr::get(getContext(), targetAS); - return getPointerTo(ty, attr); - } - auto attr = cir::AddressSpaceAttr::get(getContext(), addrSpace); - return getPointerTo(ty, attr); - } - cir::PointerType getPointerTo(mlir::Type ty, clang::LangAS langAS) { if (langAS == clang::LangAS::Default) return getPointerTo(ty); diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td index 94deeea69722..6a85651b12b7 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -987,18 +987,6 @@ def CIR_AddressSpaceAttr : CIR_EnumAttr { unsigned $cppClass::getAsUnsignedValue() const { return static_cast(getValue()); } - - bool $cppClass::isLang() const { - return cir::isLangAddressSpace(getValue()); - } - - bool $cppClass::isTarget() const { - return cir::isTargetAddressSpace(getValue()); - } - - unsigned $cppClass::getTargetValue() const { - return cir::getTargetAddressSpaceValueFromCIRAS(getValue()); - } }]; } diff --git a/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td b/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td index 56e5a4465795..6465e60cb951 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td +++ b/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td @@ -48,8 +48,7 @@ def CIR_AddressSpace : CIR_I32EnumAttr< I32EnumAttrCase<"OffloadLocal", 2, "offload_local">, I32EnumAttrCase<"OffloadGlobal", 3, "offload_global">, I32EnumAttrCase<"OffloadConstant", 4, "offload_constant">, - I32EnumAttrCase<"OffloadGeneric", 5, "offload_generic">, - I32EnumAttrCase<"Target", 6, "target"> + I32EnumAttrCase<"OffloadGeneric", 5, "offload_generic"> ]> { let description = [{ The `address_space` attribute is used to represent address spaces for diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h index b85a8acaa8d8..0f5329286eb2 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h @@ -50,33 +50,6 @@ constexpr unsigned getAsUnsignedValue(cir::AddressSpace as) { return static_cast(as); } -inline constexpr unsigned TargetAddressSpaceOffset = - cir::getMaxEnumValForAddressSpace(); - -// Target address space is used for target-specific address spaces that are not -// part of the enum. Its value is represented as an offset from the maximum -// value of the enum. Make sure that it is always the last enum value. -static_assert(getAsUnsignedValue(cir::AddressSpace::Target) == - cir::getMaxEnumValForAddressSpace(), - "Target address space must be the last enum value"); - -constexpr bool isTargetAddressSpace(cir::AddressSpace as) { - return getAsUnsignedValue(as) >= cir::getMaxEnumValForAddressSpace(); -} - -constexpr bool isLangAddressSpace(cir::AddressSpace as) { - return !isTargetAddressSpace(as); -} - -constexpr unsigned getTargetAddressSpaceValueFromCIRAS(cir::AddressSpace as) { - assert(isTargetAddressSpace(as) && "expected target address space"); - return getAsUnsignedValue(as) - TargetAddressSpaceOffset; -} - -constexpr cir::AddressSpace computeTargetAddressSpace(unsigned v) { - return static_cast(v + TargetAddressSpaceOffset); -} - } // namespace cir //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 14b76bb8b06a..fe336df7bc8d 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -34,6 +34,7 @@ #include "clang/Frontend/FrontendDiagnostic.h" #include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/IR/Attributes.h" #include "mlir/IR/BuiltinAttributes.h" #include "mlir/IR/Value.h" #include "mlir/Support/LLVM.h" @@ -1727,9 +1728,9 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID, // the AST level this is handled within CreateTempAlloca et al., but for the // builtin / dynamic alloca we have to handle it here. assert(!cir::MissingFeatures::addressSpace()); - cir::AddressSpace AAS = getCIRAllocaAddressSpace(); - cir::AddressSpace EAS = cir::toCIRAddressSpace( - E->getType()->getPointeeType().getAddressSpace()); + mlir::Attribute AAS = getCIRAllocaAddressSpace(); + mlir::Attribute EAS = cir::toCIRAddressSpaceAttr( + &getMLIRContext(), E->getType()->getPointeeType().getAddressSpace()); if (EAS != AAS) { assert(false && "Non-default address space for alloca NYI"); } diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index 3e3369d677da..8a3156d4103f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -3137,7 +3137,8 @@ Address CIRGenFunction::CreateTempAlloca(mlir::Type Ty, CharUnits Align, // be different from the type defined by the language. For example, // in C++ the auto variables are in the default address space. Therefore // cast alloca to the default address space when necessary. - if (auto ASTAS = cir::toCIRAddressSpace(CGM.getLangTempAllocaAddressSpace()); + if (auto ASTAS = cir::toCIRAddressSpaceAttr( + &getMLIRContext(), CGM.getLangTempAllocaAddressSpace()); getCIRAllocaAddressSpace() != ASTAS) { llvm_unreachable("Requires address space cast which is NYI"); } diff --git a/clang/lib/CIR/CodeGen/CIRGenTypeCache.h b/clang/lib/CIR/CodeGen/CIRGenTypeCache.h index c83d60673f23..d8f29f2d69fc 100644 --- a/clang/lib/CIR/CodeGen/CIRGenTypeCache.h +++ b/clang/lib/CIR/CodeGen/CIRGenTypeCache.h @@ -13,6 +13,7 @@ #ifndef LLVM_CLANG_LIB_CIR_CODEGENTYPECACHE_H #define LLVM_CLANG_LIB_CIR_CODEGENTYPECACHE_H +#include "mlir/IR/Attributes.h" #include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/Types.h" #include "clang/AST/CharUnits.h" @@ -106,7 +107,7 @@ struct CIRGenTypeCache { unsigned char SizeAlignInBytes; }; - cir::AddressSpace CIRAllocaAddressSpace; + mlir::Attribute CIRAllocaAddressSpace; clang::CharUnits getSizeSize() const { return clang::CharUnits::fromQuantity(SizeSizeInBytes); @@ -121,7 +122,7 @@ struct CIRGenTypeCache { return clang::CharUnits::fromQuantity(PointerAlignInBytes); } - cir::AddressSpace getCIRAllocaAddressSpace() const { + mlir::Attribute getCIRAllocaAddressSpace() const { return CIRAllocaAddressSpace; } }; diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp index ab3609d286a7..7cb5de0f5444 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.cpp +++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp @@ -266,8 +266,9 @@ class CommonSPIRTargetCIRGenInfo : public TargetCIRGenInfo { CommonSPIRTargetCIRGenInfo(std::unique_ptr ABIInfo) : TargetCIRGenInfo(std::move(ABIInfo)) {} - cir::AddressSpace getCIRAllocaAddressSpace() const override { - return cir::AddressSpace::OffloadPrivate; + mlir::Attribute getCIRAllocaAddressSpace() const override { + return cir::AddressSpaceAttr::get(&getABIInfo().CGT.getMLIRContext(), + cir::AddressSpace::OffloadPrivate); } cir::CallingConv getOpenCLKernelCallingConv() const override { diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h index 38c302b672b7..5a6b296edd50 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -102,8 +102,8 @@ class TargetCIRGenInfo { const clang::VarDecl *D) const; /// Get the CIR address space for alloca. - virtual cir::AddressSpace getCIRAllocaAddressSpace() const { - return cir::AddressSpace::Default; + virtual mlir::Attribute getCIRAllocaAddressSpace() const { + return {}; // Empty attribute represents the Default address space } /// Perform address space cast of an expression of pointer type. diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index cafb43e9f98d..4256b89a211b 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -2488,7 +2488,6 @@ cir::GetGlobalOp::verifySymbolUses(SymbolTableCollection &symbolTable) { mlir::Attribute symAddrSpaceAttr{}; if (auto g = dyn_cast(op)) { symTy = g.getSymType(); - // Convert enum to attribute for comparison symAddrSpaceAttr = g.getAddrSpace(); // Verify that for thread local global access, the global needs to // be marked with tls bits. diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index e47d44c4cbd3..6eb5e2991835 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -957,8 +957,6 @@ cir::AddressSpace cir::toCIRAddressSpace(clang::LangAS langAS) { case LangAS::wasm_funcref: llvm_unreachable("NYI"); default: - // NOTE: In theory with TargetAddressSpaceAttr, we don't care at all about - // representing target AS here. llvm_unreachable("unknown/unsupported clang language address space"); } } @@ -1037,7 +1035,7 @@ mlir::Attribute cir::toCIRAddressSpaceAttr(mlir::MLIRContext *ctx, using clang::LangAS; if (langAS == LangAS::Default) - return {}; + return {}; // Default address space is represented as an empty attribute. if (clang::isTargetAddressSpace(langAS)) { unsigned targetAS = clang::toTargetAddressSpace(langAS); @@ -1047,16 +1045,16 @@ mlir::Attribute cir::toCIRAddressSpaceAttr(mlir::MLIRContext *ctx, return cir::AddressSpaceAttr::get(ctx, toCIRAddressSpace(langAS)); } -cir::AddressSpace cir::getCIRAddressSpaceFromAttr(mlir::Attribute attr) { - if (!attr) - return AddressSpace::Default; - if (auto addrSpaceAttr = mlir::dyn_cast(attr)) - return addrSpaceAttr.getValue(); - if (auto targetAddrSpaceAttr = - mlir::dyn_cast(attr)) - return cir::computeTargetAddressSpace(targetAddrSpaceAttr.getValue()); - return AddressSpace::Default; -} +// cir::AddressSpace cir::getCIRAddressSpaceFromAttr(mlir::Attribute attr) { +// if (!attr) +// return AddressSpace::Default; +// if (auto addrSpaceAttr = mlir::dyn_cast(attr)) +// return addrSpaceAttr.getValue(); +// if (auto targetAddrSpaceAttr = +// mlir::dyn_cast(attr)) +// return cir::computeTargetAddressSpace(targetAddrSpaceAttr.getValue()); +// return AddressSpace::Default; +// } //===----------------------------------------------------------------------===// // PointerType Definitions diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index ef2d187d5da6..86d7a39ed9d4 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -5003,17 +5003,11 @@ std::unique_ptr prepareLowerModule(mlir::ModuleOp module) { } static unsigned -getTargetAddrSpaceFromCIRAddrSpace(cir::AddressSpace addrSpace, - cir::LowerModule *lowerModule) { - if (addrSpace == cir::AddressSpace::Default) - return 0; // Default address space is always 0 in LLVM. - - if (cir::isTargetAddressSpace(addrSpace)) - return cir::getTargetAddressSpaceValueFromCIRAS(addrSpace); - +getTargetAddrSpaceFromCIRAddrSpaceAttr(cir::AddressSpaceAttr addrSpace, + cir::LowerModule *lowerModule) { assert(lowerModule && "CIR AS map is not available"); return lowerModule->getTargetLoweringInfo() - .getTargetAddrSpaceFromCIRAddrSpace(addrSpace); + .getTargetAddrSpaceFromCIRAddrSpace(addrSpace.getValue()); } static unsigned getTargetAddrSpaceFromASAttr(mlir::Attribute attr, @@ -5025,9 +5019,9 @@ static unsigned getTargetAddrSpaceFromASAttr(mlir::Attribute attr, mlir::dyn_cast(attr)) return targetAddrSpaceAttr.getValue(); - auto addrSpaceAttr = mlir::dyn_cast(attr); - return getTargetAddrSpaceFromCIRAddrSpace(addrSpaceAttr.getValue(), - lowerModule); + cir::AddressSpaceAttr addrSpaceAttr = + mlir::dyn_cast(attr); + return getTargetAddrSpaceFromCIRAddrSpaceAttr(addrSpaceAttr, lowerModule); } // FIXME: change the type of lowerModule to `LowerModule &` to have better diff --git a/clang/test/CIR/IR/invalid.cir b/clang/test/CIR/IR/invalid.cir index 0b30ee2b3757..22b2a9ce47ab 100644 --- a/clang/test/CIR/IR/invalid.cir +++ b/clang/test/CIR/IR/invalid.cir @@ -1251,7 +1251,7 @@ cir.func @address_space2(%p : !cir.ptr) { !u64i = !cir.int -// expected-error@below {{expected one of [default, offload_private, offload_local, offload_global, offload_constant, offload_generic, target] for clang address space kind, got: foobar}} +// expected-error@below {{expected one of [default, offload_private, offload_local, offload_global, offload_constant, offload_generic] for clang address space kind, got: foobar}} // expected-error@below {{expected clang address space keyword}} cir.func @address_space4(%p : !cir.ptr) { cir.return From 49f7c11230460e66ee3b23d5e1864938d96e8ab4 Mon Sep 17 00:00:00 2001 From: David Rivera Date: Thu, 20 Nov 2025 07:20:25 -0500 Subject: [PATCH 07/13] Rename AddressSpace to ClangAddressSpace for clarity --- .../CIR/Dialect/Builder/CIRBaseBuilder.h | 4 +- .../CIR/Dialect/IR/CIRAttrConstraints.td | 2 +- .../include/clang/CIR/Dialect/IR/CIRAttrs.td | 9 ++-- .../clang/CIR/Dialect/IR/CIREnumAttr.td | 4 +- clang/include/clang/CIR/Dialect/IR/CIRTypes.h | 8 ++-- clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 2 +- clang/lib/CIR/CodeGen/CIRGenDecl.cpp | 2 +- clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 6 +-- clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp | 4 +- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 8 ++-- clang/lib/CIR/CodeGen/TargetInfo.cpp | 4 +- clang/lib/CIR/Dialect/IR/CIRAttrs.cpp | 8 ++-- clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 45 +++++++------------ .../TargetLowering/TargetLoweringInfo.h | 2 +- .../TargetLowering/Targets/AArch64.cpp | 12 ++--- .../TargetLowering/Targets/AMDGPU.cpp | 12 ++--- .../TargetLowering/Targets/NVPTX.cpp | 12 ++--- .../TargetLowering/Targets/SPIR.cpp | 12 ++--- .../Transforms/TargetLowering/Targets/X86.cpp | 12 ++--- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 10 ++--- 20 files changed, 81 insertions(+), 97 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h index adc5e8dd0944..534b667f5d3d 100644 --- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h +++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h @@ -108,7 +108,7 @@ class CIRBaseBuilderTy : public mlir::OpBuilder { /// Create a pointer type with an address space attribute. cir::PointerType getPointerTo(mlir::Type ty, mlir::Attribute addrSpaceAttr) { - assert(mlir::isa(addrSpaceAttr) || + assert(mlir::isa(addrSpaceAttr) || mlir::isa(addrSpaceAttr) && "expected address space attribute"); return cir::PointerType::get(ty, addrSpaceAttr); @@ -119,7 +119,7 @@ class CIRBaseBuilderTy : public mlir::OpBuilder { return getPointerTo(ty); mlir::Attribute addrSpaceAttr = - cir::toCIRAddressSpaceAttr(getContext(), langAS); + cir::toCIRClangAddressSpaceAttr(getContext(), langAS); return getPointerTo(ty, addrSpaceAttr); } diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td index a4dd4f6375b4..2b8054036b66 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td @@ -96,7 +96,7 @@ def CIR_AnyTBAAAttr : AnyAttrOf<[ // NOTE: We might end up using this only for GlobalOps, as we cannot apply constraints // to types. def CIR_AddressSpaceAttrConstraint - : CIR_AttrConstraint<"::cir::AddressSpaceAttr", "language address space attribute">; + : CIR_AttrConstraint<"::cir::ClangAddressSpaceAttr", "clang address space attribute">; def CIR_TargetAddressSpaceAttrConstraint : CIR_AttrConstraint<"::cir::TargetAddressSpaceAttr", "target address space attribute">; diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td index 6a85651b12b7..c111f6f214d4 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -963,10 +963,10 @@ def CIR_DynamicCastInfoAttr : CIR_Attr<"DynamicCastInfo", "dyn_cast_info"> { // AddressSpaceAttr //===----------------------------------------------------------------------===// -def CIR_AddressSpaceAttr : CIR_EnumAttr { +def CIR_ClangAddressSpaceAttr : CIR_EnumAttr { let builders = [ AttrBuilder<(ins "clang::LangAS":$langAS), [{ - return $_get($_ctxt, cir::toCIRAddressSpace(langAS)); + return $_get($_ctxt, cir::toCIRClangAddressSpace(langAS)); }]> ]; @@ -974,12 +974,9 @@ def CIR_AddressSpaceAttr : CIR_EnumAttr { `(` custom($value) `)` }]; - let defaultValue = "cir::AddressSpace::Default"; + let defaultValue = "cir::ClangAddressSpace::Default"; let extraClassDeclaration = [{ - bool isLang() const; - bool isTarget() const; - unsigned getTargetValue() const; unsigned getAsUnsignedValue() const; }]; diff --git a/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td b/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td index 6465e60cb951..5e3901b699cb 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td +++ b/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td @@ -41,8 +41,8 @@ class CIR_DefaultValuedEnumParameter let defaultValue = value; } -def CIR_AddressSpace : CIR_I32EnumAttr< - "AddressSpace", "clang address space kind", [ +def CIR_ClangAddressSpace : CIR_I32EnumAttr< + "ClangAddressSpace", "clang address space kind", [ I32EnumAttrCase<"Default", 0, "default">, I32EnumAttrCase<"OffloadPrivate", 1, "offload_private">, I32EnumAttrCase<"OffloadLocal", 2, "offload_local">, diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h index 0f5329286eb2..9c669d46d6f7 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h @@ -34,19 +34,19 @@ bool isSized(mlir::Type ty); // AddressSpace helpers //===----------------------------------------------------------------------===// -cir::AddressSpace toCIRAddressSpace(clang::LangAS langAS); +cir::ClangAddressSpace toCIRClangAddressSpace(clang::LangAS langAS); /// Convert a LangAS to the appropriate address space attribute. /// Returns AddressSpaceAttr for language-specific address spaces, /// or TargetAddressSpaceAttr for target-specific address spaces. -mlir::Attribute toCIRAddressSpaceAttr(mlir::MLIRContext *ctx, +mlir::Attribute toCIRClangAddressSpaceAttr(mlir::MLIRContext *ctx, clang::LangAS langAS); /// Extract the AddressSpace enum from an address space attribute. /// Returns Default if the attribute is null. -cir::AddressSpace getCIRAddressSpaceFromAttr(mlir::Attribute attr); +cir::ClangAddressSpace getCIRClangAddressSpaceFromAttr(mlir::Attribute attr); -constexpr unsigned getAsUnsignedValue(cir::AddressSpace as) { +constexpr unsigned getAsUnsignedValue(cir::ClangAddressSpace as) { return static_cast(as); } diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index fe336df7bc8d..51e4efc2ab48 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -1729,7 +1729,7 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID, // builtin / dynamic alloca we have to handle it here. assert(!cir::MissingFeatures::addressSpace()); mlir::Attribute AAS = getCIRAllocaAddressSpace(); - mlir::Attribute EAS = cir::toCIRAddressSpaceAttr( + mlir::Attribute EAS = cir::toCIRClangAddressSpaceAttr( &getMLIRContext(), E->getType()->getPointeeType().getAddressSpace()); if (EAS != AAS) { assert(false && "Non-default address space for alloca NYI"); diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp index 80006e34d192..df8591d386bf 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp @@ -479,7 +479,7 @@ CIRGenModule::getOrCreateStaticVarDecl(const VarDecl &D, Name = getStaticDeclName(*this, D); mlir::Type LTy = getTypes().convertTypeForMem(Ty); - mlir::Attribute AS = cir::toCIRAddressSpaceAttr(&getMLIRContext(), + mlir::Attribute AS = cir::toCIRClangAddressSpaceAttr(&getMLIRContext(), getGlobalVarAddressSpace(&D)); // OpenCL variables in local address space and CUDA shared diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index 8a3156d4103f..d2ebe878d288 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -2080,9 +2080,9 @@ LValue CIRGenFunction::emitCastLValue(const CastExpr *E) { case CK_AddressSpaceConversion: { LValue LV = emitLValue(E->getSubExpr()); QualType DestTy = getContext().getPointerType(E->getType()); - mlir::Attribute SrcAS = cir::toCIRAddressSpaceAttr( + mlir::Attribute SrcAS = cir::toCIRClangAddressSpaceAttr( &getMLIRContext(), E->getSubExpr()->getType().getAddressSpace()); - mlir::Attribute DestAS = cir::toCIRAddressSpaceAttr( + mlir::Attribute DestAS = cir::toCIRClangAddressSpaceAttr( &getMLIRContext(), E->getType().getAddressSpace()); mlir::Value V = getTargetHooks().performAddrSpaceCast( *this, LV.getPointer(), SrcAS, DestAS, convertType(DestTy)); @@ -3137,7 +3137,7 @@ Address CIRGenFunction::CreateTempAlloca(mlir::Type Ty, CharUnits Align, // be different from the type defined by the language. For example, // in C++ the auto variables are in the default address space. Therefore // cast alloca to the default address space when necessary. - if (auto ASTAS = cir::toCIRAddressSpaceAttr( + if (auto ASTAS = cir::toCIRClangAddressSpaceAttr( &getMLIRContext(), CGM.getLangTempAllocaAddressSpace()); getCIRAllocaAddressSpace() != ASTAS) { llvm_unreachable("Requires address space cast which is NYI"); diff --git a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp index 199b6dfe1327..418d9a86def9 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp @@ -1690,10 +1690,10 @@ mlir::Value ScalarExprEmitter::VisitCastExpr(CastExpr *CE) { } // Since target may map different address spaces in AST to the same address // space, an address space conversion may end up as a bitcast. - mlir::Attribute SrcAS = cir::toCIRAddressSpaceAttr( + mlir::Attribute SrcAS = cir::toCIRClangAddressSpaceAttr( &CGF.getMLIRContext(), E->getType()->getPointeeType().getAddressSpace()); - mlir::Attribute DestAS = cir::toCIRAddressSpaceAttr( + mlir::Attribute DestAS = cir::toCIRClangAddressSpaceAttr( &CGF.getMLIRContext(), DestTy->getPointeeType().getAddressSpace()); return CGF.CGM.getTargetCIRGenInfo().performAddrSpaceCast( diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 6987f5ae9507..59299f34bff2 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -1112,7 +1112,7 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty, entry = dyn_cast_or_null(v); } - mlir::Attribute cirAS = cir::toCIRAddressSpaceAttr(&getMLIRContext(), langAS); + mlir::Attribute cirAS = cir::toCIRClangAddressSpaceAttr(&getMLIRContext(), langAS); if (entry) { mlir::Attribute entryCIRAS = entry.getAddrSpace(); if (WeakRefReferences.erase(entry)) { @@ -1168,7 +1168,7 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty, return entry; } - mlir::Attribute declCIRAS = cir::toCIRAddressSpaceAttr(&getMLIRContext(), getGlobalVarAddressSpace(d)); + mlir::Attribute declCIRAS = cir::toCIRClangAddressSpaceAttr(&getMLIRContext(), getGlobalVarAddressSpace(d)); // TODO(cir): do we need to strip pointer casts for Entry? auto loc = getLoc(d->getSourceRange()); @@ -1773,7 +1773,7 @@ static cir::GlobalOp generateStringLiteral(mlir::Location loc, mlir::TypedAttr c, cir::GlobalLinkageKind lt, CIRGenModule &cgm, StringRef globalName, CharUnits alignment) { - mlir::Attribute addrSpace = cir::toCIRAddressSpaceAttr( + mlir::Attribute addrSpace = cir::toCIRClangAddressSpaceAttr( &cgm.getMLIRContext(), cgm.getGlobalConstantAddressSpace()); // Create a global variable for this string @@ -1987,7 +1987,7 @@ CIRGenModule::getAddrOfGlobalTemporary(const MaterializeTemporaryExpr *expr, linkage = cir::GlobalLinkageKind::InternalLinkage; } } - mlir::Attribute targetAS = cir::toCIRAddressSpaceAttr(&getMLIRContext(), addrSpace); + mlir::Attribute targetAS = cir::toCIRClangAddressSpaceAttr(&getMLIRContext(), addrSpace); auto loc = getLoc(expr->getSourceRange()); auto gv = createGlobalOp(*this, loc, name, type, isConstant, targetAS, diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp index 7cb5de0f5444..a4ddfbb69ccf 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.cpp +++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp @@ -267,8 +267,8 @@ class CommonSPIRTargetCIRGenInfo : public TargetCIRGenInfo { : TargetCIRGenInfo(std::move(ABIInfo)) {} mlir::Attribute getCIRAllocaAddressSpace() const override { - return cir::AddressSpaceAttr::get(&getABIInfo().CGT.getMLIRContext(), - cir::AddressSpace::OffloadPrivate); + return cir::ClangAddressSpaceAttr::get(&getABIInfo().CGT.getMLIRContext(), + cir::ClangAddressSpace::OffloadPrivate); } cir::CallingConv getOpenCLKernelCallingConv() const override { diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp index dc241106e1c7..b2b83da98a7f 100644 --- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp @@ -77,17 +77,17 @@ static void printConstPtr(mlir::AsmPrinter &p, mlir::IntegerAttr value); //===----------------------------------------------------------------------===// mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, - cir::AddressSpace &addrSpace) { + cir::ClangAddressSpace &addrSpace) { llvm::SMLoc loc = p.getCurrentLocation(); - mlir::FailureOr result = - mlir::FieldParser::parse(p); + mlir::FailureOr result = + mlir::FieldParser::parse(p); if (mlir::failed(result)) return p.emitError(loc, "expected address space keyword"); addrSpace = result.value(); return mlir::success(); } -void printAddressSpaceValue(mlir::AsmPrinter &p, cir::AddressSpace addrSpace) { +void printAddressSpaceValue(mlir::AsmPrinter &p, cir::ClangAddressSpace addrSpace) { p << cir::stringifyEnum(addrSpace); } diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index 6eb5e2991835..576201de3b62 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -922,27 +922,27 @@ MethodType::getABIAlignment(const mlir::DataLayout &dataLayout, // AddressSpace definitions //===----------------------------------------------------------------------===// -cir::AddressSpace cir::toCIRAddressSpace(clang::LangAS langAS) { +cir::ClangAddressSpace cir::toCIRClangAddressSpace(clang::LangAS langAS) { using clang::LangAS; switch (langAS) { case LangAS::Default: - return AddressSpace::Default; + return ClangAddressSpace::Default; case LangAS::opencl_global: - return AddressSpace::OffloadGlobal; + return ClangAddressSpace::OffloadGlobal; case LangAS::opencl_local: case LangAS::cuda_shared: // Local means local among the work-group (OpenCL) or block (CUDA). // All threads inside the kernel can access local memory. - return AddressSpace::OffloadLocal; + return ClangAddressSpace::OffloadLocal; case LangAS::cuda_device: - return AddressSpace::OffloadGlobal; + return ClangAddressSpace::OffloadGlobal; case LangAS::opencl_constant: case LangAS::cuda_constant: - return AddressSpace::OffloadConstant; + return ClangAddressSpace::OffloadConstant; case LangAS::opencl_private: - return AddressSpace::OffloadPrivate; + return ClangAddressSpace::OffloadPrivate; case LangAS::opencl_generic: - return AddressSpace::OffloadGeneric; + return ClangAddressSpace::OffloadGeneric; case LangAS::opencl_global_device: case LangAS::opencl_global_host: case LangAS::sycl_global: @@ -984,13 +984,13 @@ mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, if (p.parseOptionalKeyword("clang_address_space").succeeded()) { if (p.parseLParen()) return p.emitError(loc, "expected '(' after clang address space"); - mlir::FailureOr result = - mlir::FieldParser::parse(p); + mlir::FailureOr result = + mlir::FieldParser::parse(p); if (mlir::failed(result) || p.parseRParen()) return p.emitError(loc, "expected clang address space keyword"); - attr = cir::AddressSpaceAttr::get(p.getContext(), result.value()); + attr = cir::ClangAddressSpaceAttr::get(p.getContext(), result.value()); return mlir::success(); } @@ -1001,9 +1001,9 @@ void printAddressSpaceValue(mlir::AsmPrinter &p, mlir::Attribute attr) { if (!attr) return; - if (auto logical = dyn_cast(attr)) { + if (auto logical = dyn_cast(attr)) { p << "clang_address_space(" - << cir::stringifyAddressSpace(logical.getValue()) << ')'; + << cir::stringifyClangAddressSpace(logical.getValue()) << ')'; return; } @@ -1030,7 +1030,7 @@ void printGlobalAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp, printAddressSpaceValue(printer, attr); } -mlir::Attribute cir::toCIRAddressSpaceAttr(mlir::MLIRContext *ctx, +mlir::Attribute cir::toCIRClangAddressSpaceAttr(mlir::MLIRContext *ctx, clang::LangAS langAS) { using clang::LangAS; @@ -1042,20 +1042,9 @@ mlir::Attribute cir::toCIRAddressSpaceAttr(mlir::MLIRContext *ctx, return cir::TargetAddressSpaceAttr::get(ctx, targetAS); } - return cir::AddressSpaceAttr::get(ctx, toCIRAddressSpace(langAS)); + return cir::ClangAddressSpaceAttr::get(ctx, toCIRClangAddressSpace(langAS)); } -// cir::AddressSpace cir::getCIRAddressSpaceFromAttr(mlir::Attribute attr) { -// if (!attr) -// return AddressSpace::Default; -// if (auto addrSpaceAttr = mlir::dyn_cast(attr)) -// return addrSpaceAttr.getValue(); -// if (auto targetAddrSpaceAttr = -// mlir::dyn_cast(attr)) -// return cir::computeTargetAddressSpace(targetAddrSpaceAttr.getValue()); -// return AddressSpace::Default; -// } - //===----------------------------------------------------------------------===// // PointerType Definitions //===----------------------------------------------------------------------===// @@ -1064,7 +1053,7 @@ mlir::LogicalResult cir::PointerType::verify( llvm::function_ref emitError, mlir::Type pointee, mlir::Attribute addrSpace) { if (addrSpace) { - if (!mlir::isa(addrSpace) && + if (!mlir::isa(addrSpace) && !mlir::isa(addrSpace)) { return emitError() << "pointer address space must be either " "!cir.address_space or !cir.target_address_space"; @@ -1079,7 +1068,7 @@ bool PointerType::hasTargetAddressSpace() const { } bool PointerType::hasLanguageAddressSpace() const { - return mlir::isa_and_nonnull(getAddrSpace()); + return mlir::isa_and_nonnull(getAddrSpace()); } //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h index 114d8cc0f697..d5c6111b7ae7 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h @@ -33,7 +33,7 @@ class TargetLoweringInfo { const ABIInfo &getABIInfo() const { return *Info; } virtual unsigned - getTargetAddrSpaceFromCIRAddrSpace(cir::AddressSpace addrSpace) const = 0; + getTargetAddrSpaceFromCIRAddrSpace(cir::ClangAddressSpace addrSpace) const = 0; }; } // namespace cir diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AArch64.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AArch64.cpp index 9a8edf39d738..405e3db8d0f2 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AArch64.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AArch64.cpp @@ -62,13 +62,13 @@ class AArch64TargetLoweringInfo : public TargetLoweringInfo { } unsigned getTargetAddrSpaceFromCIRAddrSpace( - cir::AddressSpace addrSpace) const override { + cir::ClangAddressSpace addrSpace) const override { switch (addrSpace) { - case cir::AddressSpace::OffloadPrivate: - case cir::AddressSpace::OffloadLocal: - case cir::AddressSpace::OffloadGlobal: - case cir::AddressSpace::OffloadConstant: - case cir::AddressSpace::OffloadGeneric: + case cir::ClangAddressSpace::OffloadPrivate: + case cir::ClangAddressSpace::OffloadLocal: + case cir::ClangAddressSpace::OffloadGlobal: + case cir::ClangAddressSpace::OffloadConstant: + case cir::ClangAddressSpace::OffloadGeneric: return 0; default: cir_cconv_unreachable("Unknown CIR address space for this target"); diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp index 7432972889ed..b0bb66b05cf5 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp @@ -44,17 +44,17 @@ class AMDGPUTargetLoweringInfo : public TargetLoweringInfo { : TargetLoweringInfo(std::make_unique(lt)) {} // Taken from here: https://llvm.org/docs/AMDGPUUsage.html#address-spaces unsigned getTargetAddrSpaceFromCIRAddrSpace( - cir::AddressSpace addrSpace) const override { + cir::ClangAddressSpace addrSpace) const override { switch (addrSpace) { - case cir::AddressSpace::OffloadPrivate: + case cir::ClangAddressSpace::OffloadPrivate: return 5; - case cir::AddressSpace::OffloadLocal: + case cir::ClangAddressSpace::OffloadLocal: return 3; - case cir::AddressSpace::OffloadGlobal: + case cir::ClangAddressSpace::OffloadGlobal: return 1; - case cir::AddressSpace::OffloadConstant: + case cir::ClangAddressSpace::OffloadConstant: return 4; - case cir::AddressSpace::OffloadGeneric: + case cir::ClangAddressSpace::OffloadGeneric: return 0; default: cir_cconv_unreachable("Unknown CIR address space for this target"); diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp index 00f961d38666..6dbab8439f85 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp @@ -45,17 +45,17 @@ class NVPTXTargetLoweringInfo : public TargetLoweringInfo { : TargetLoweringInfo(std::make_unique(lt)) {} unsigned getTargetAddrSpaceFromCIRAddrSpace( - cir::AddressSpace addrSpace) const override { + cir::ClangAddressSpace addrSpace) const override { switch (addrSpace) { - case cir::AddressSpace::OffloadPrivate: + case cir::ClangAddressSpace::OffloadPrivate: return 0; - case cir::AddressSpace::OffloadLocal: + case cir::ClangAddressSpace::OffloadLocal: return 3; - case cir::AddressSpace::OffloadGlobal: + case cir::ClangAddressSpace::OffloadGlobal: return 1; - case cir::AddressSpace::OffloadConstant: + case cir::ClangAddressSpace::OffloadConstant: return 4; - case cir::AddressSpace::OffloadGeneric: + case cir::ClangAddressSpace::OffloadGeneric: return 0; default: cir_cconv_unreachable("Unknown CIR address space for this target"); diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIR.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIR.cpp index 0a4dc640decd..9e0b98079e00 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIR.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIR.cpp @@ -42,17 +42,17 @@ class SPIRVTargetLoweringInfo : public TargetLoweringInfo { : TargetLoweringInfo(std::make_unique(LT)) {} unsigned getTargetAddrSpaceFromCIRAddrSpace( - cir::AddressSpace addrSpace) const override { + cir::ClangAddressSpace addrSpace) const override { switch (addrSpace) { - case cir::AddressSpace::OffloadPrivate: + case cir::ClangAddressSpace::OffloadPrivate: return 0; - case cir::AddressSpace::OffloadLocal: + case cir::ClangAddressSpace::OffloadLocal: return 3; - case cir::AddressSpace::OffloadGlobal: + case cir::ClangAddressSpace::OffloadGlobal: return 1; - case cir::AddressSpace::OffloadConstant: + case cir::ClangAddressSpace::OffloadConstant: return 2; - case cir::AddressSpace::OffloadGeneric: + case cir::ClangAddressSpace::OffloadGeneric: return 4; default: cir_cconv_unreachable("Unknown CIR address space for this target"); diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/X86.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/X86.cpp index ec8c880ef3ab..6f085e43bf35 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/X86.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/X86.cpp @@ -110,13 +110,13 @@ class X86_64TargetLoweringInfo : public TargetLoweringInfo { } unsigned getTargetAddrSpaceFromCIRAddrSpace( - cir::AddressSpace addrSpace) const override { + cir::ClangAddressSpace addrSpace) const override { switch (addrSpace) { - case cir::AddressSpace::OffloadPrivate: - case cir::AddressSpace::OffloadLocal: - case cir::AddressSpace::OffloadGlobal: - case cir::AddressSpace::OffloadConstant: - case cir::AddressSpace::OffloadGeneric: + case cir::ClangAddressSpace::OffloadPrivate: + case cir::ClangAddressSpace::OffloadLocal: + case cir::ClangAddressSpace::OffloadGlobal: + case cir::ClangAddressSpace::OffloadConstant: + case cir::ClangAddressSpace::OffloadGeneric: return 0; default: cir_cconv_unreachable("Unknown CIR address space for this target"); diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 86d7a39ed9d4..7bbd74e0d04e 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -382,8 +382,6 @@ void lowerAnnotationValue( } } - - // Get addrspace by converting a pointer type. // TODO: The approach here is a little hacky. We should access the target info // directly to convert the address space of global op, similar to what we do @@ -5003,7 +5001,7 @@ std::unique_ptr prepareLowerModule(mlir::ModuleOp module) { } static unsigned -getTargetAddrSpaceFromCIRAddrSpaceAttr(cir::AddressSpaceAttr addrSpace, +getTargetAddrSpaceFromCIRAddrSpaceAttr(cir::ClangAddressSpaceAttr addrSpace, cir::LowerModule *lowerModule) { assert(lowerModule && "CIR AS map is not available"); return lowerModule->getTargetLoweringInfo() @@ -5012,15 +5010,15 @@ getTargetAddrSpaceFromCIRAddrSpaceAttr(cir::AddressSpaceAttr addrSpace, static unsigned getTargetAddrSpaceFromASAttr(mlir::Attribute attr, cir::LowerModule *lowerModule) { - assert(mlir::isa_and_nonnull(attr) || + assert(mlir::isa_and_nonnull(attr) || mlir::isa_and_nonnull(attr)); if (auto targetAddrSpaceAttr = mlir::dyn_cast(attr)) return targetAddrSpaceAttr.getValue(); - cir::AddressSpaceAttr addrSpaceAttr = - mlir::dyn_cast(attr); + cir::ClangAddressSpaceAttr addrSpaceAttr = + mlir::dyn_cast(attr); return getTargetAddrSpaceFromCIRAddrSpaceAttr(addrSpaceAttr, lowerModule); } From 5a5d8d2648e29bd88580610bb494ca52755afdcc Mon Sep 17 00:00:00 2001 From: David Rivera Date: Thu, 20 Nov 2025 08:27:32 -0500 Subject: [PATCH 08/13] Fix crash on runnin isa on empty attr --- .../include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h index 534b667f5d3d..c6630f5172f8 100644 --- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h +++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h @@ -108,9 +108,13 @@ class CIRBaseBuilderTy : public mlir::OpBuilder { /// Create a pointer type with an address space attribute. cir::PointerType getPointerTo(mlir::Type ty, mlir::Attribute addrSpaceAttr) { - assert(mlir::isa(addrSpaceAttr) || - mlir::isa(addrSpaceAttr) && - "expected address space attribute"); + if (!addrSpaceAttr) + return cir::PointerType::get(ty); + + assert((mlir::isa(addrSpaceAttr) || + mlir::isa(addrSpaceAttr)) && + "expected address space attribute"); + return cir::PointerType::get(ty, addrSpaceAttr); } From 2fe4cfd1a179b703db3e31e8806dbecbd9346252 Mon Sep 17 00:00:00 2001 From: David Rivera Date: Thu, 20 Nov 2025 10:40:14 -0500 Subject: [PATCH 09/13] Polish comments and docs description --- .../CIR/Dialect/Builder/CIRBaseBuilder.h | 2 ++ .../include/clang/CIR/Dialect/IR/CIRAttrs.td | 36 ++++++++++++------- clang/include/clang/CIR/Dialect/IR/CIRTypes.h | 6 +--- .../include/clang/CIR/Dialect/IR/CIRTypes.td | 6 ---- clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 13 ++----- 5 files changed, 30 insertions(+), 33 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h index c6630f5172f8..6ba1bf585bf7 100644 --- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h +++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h @@ -107,6 +107,8 @@ class CIRBaseBuilderTy : public mlir::OpBuilder { } /// Create a pointer type with an address space attribute. + /// Either a cir::ClangAddressSpaceAttr or cir::TargetAddressSpaceAttr is + /// expected. cir::PointerType getPointerTo(mlir::Type ty, mlir::Attribute addrSpaceAttr) { if (!addrSpaceAttr) return cir::PointerType::get(ty); diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td index c111f6f214d4..e272f3a8da1c 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -963,7 +963,23 @@ def CIR_DynamicCastInfoAttr : CIR_Attr<"DynamicCastInfo", "dyn_cast_info"> { // AddressSpaceAttr //===----------------------------------------------------------------------===// -def CIR_ClangAddressSpaceAttr : CIR_EnumAttr { +def CIR_ClangAddressSpaceAttr : CIR_EnumAttr { + + let summary = "Represents a language/Clang-level address space"; + let description = [{ + Encodes the semantic address spaces defined by the front-end language + (e.g. `__shared__`, `__constant__`, `__local__`). Values are stored using the + `cir::ClangAddressSpace` enum, keeping the representation compact while and + preserves the qualifier until it is mapped onto target/LLVM address-space numbers. + + Example: + ``` mlir + !cir.ptr + cir.global constant external clang_address_space(offload_constant) + + ``` + }]; + let builders = [ AttrBuilder<(ins "clang::LangAS":$langAS), [{ return $_get($_ctxt, cir::toCIRClangAddressSpace(langAS)); @@ -991,31 +1007,27 @@ def CIR_ClangAddressSpaceAttr : CIR_EnumAttr { - let summary = "Target-specific numeric address space attribute"; +def CIR_TargetAddressSpaceAttr : CIR_Attr< "TargetAddressSpace", + "target_address_space"> { + let summary = "Represents a target-specific numeric address space"; let description = [{ The TargetAddressSpaceAttr represents a target-specific numeric address space, corresponding to the LLVM IR `addressspace` qualifier and the clang - `target_address_space` attribute. + `address_space` attribute. A value of zero represents the default address space. The semantics of non-zero address spaces are target-specific. - Unlike `AddressSpaceAttr` which represents language-specific address spaces - (like OpenCL/CUDA address spaces), this attribute directly represents a - target-specific numeric address space value. - Example: ```mlir + // Target-specific numeric address spaces + !cir.ptr !cir.ptr ``` }]; let parameters = (ins "unsigned":$value); - - let assemblyFormat = [{ - `<` $value `>` - }]; + let assemblyFormat = "`<` `target` `<` $value `>` `>`"; } //===----------------------------------------------------------------------===// diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h index 9c669d46d6f7..804dfe1900fd 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h @@ -37,15 +37,11 @@ bool isSized(mlir::Type ty); cir::ClangAddressSpace toCIRClangAddressSpace(clang::LangAS langAS); /// Convert a LangAS to the appropriate address space attribute. -/// Returns AddressSpaceAttr for language-specific address spaces, +/// Returns ClangAddressSpaceAttr for clang/language-specific address spaces, /// or TargetAddressSpaceAttr for target-specific address spaces. mlir::Attribute toCIRClangAddressSpaceAttr(mlir::MLIRContext *ctx, clang::LangAS langAS); -/// Extract the AddressSpace enum from an address space attribute. -/// Returns Default if the attribute is null. -cir::ClangAddressSpace getCIRClangAddressSpaceFromAttr(mlir::Attribute attr); - constexpr unsigned getAsUnsignedValue(cir::ClangAddressSpace as) { return static_cast(as); } diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td index 6dd385a2ef09..2f0814c94fc5 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td @@ -287,12 +287,6 @@ def CIR_PointerType : CIR_Type<"Pointer", "ptr", [ return ptrType.isPtrTo(type); return false; } - - /// Returns true if this pointer type uses a target address space. - bool hasTargetAddressSpace() const; - - /// Returns true if this pointer type uses a language (logical) address space. - bool hasLanguageAddressSpace() const; }]; } diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index 576201de3b62..78687c203c99 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -1055,22 +1055,15 @@ mlir::LogicalResult cir::PointerType::verify( if (addrSpace) { if (!mlir::isa(addrSpace) && !mlir::isa(addrSpace)) { - return emitError() << "pointer address space must be either " - "!cir.address_space or !cir.target_address_space"; + return emitError() + << "pointer address space must be either " + "clang_address_space or target_address_space attribute"; } } return success(); } -bool PointerType::hasTargetAddressSpace() const { - return mlir::isa_and_nonnull(getAddrSpace()); -} - -bool PointerType::hasLanguageAddressSpace() const { - return mlir::isa_and_nonnull(getAddrSpace()); -} - //===----------------------------------------------------------------------===// // CIR Dialect //===----------------------------------------------------------------------===// From 800bc75b72912a9015a46615f4723302d9dcdf87 Mon Sep 17 00:00:00 2001 From: David Rivera Date: Thu, 20 Nov 2025 10:45:58 -0500 Subject: [PATCH 10/13] Bring Assertion guards for data ptr size based on AS --- clang/include/clang/CIR/MissingFeatures.h | 2 ++ clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 1 + 2 files changed, 3 insertions(+) diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h index 86e0653b2f8c..5152ce6770c8 100644 --- a/clang/include/clang/CIR/MissingFeatures.h +++ b/clang/include/clang/CIR/MissingFeatures.h @@ -476,6 +476,8 @@ struct MissingFeatures { static bool mustProgress() { return false; } static bool skipTempCopy() { return false; } + + static bool dataLayoutPtrHandlingBasedOnLangAS() { return false; } }; } // namespace cir diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index 78687c203c99..73dc2c01cb9f 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -370,6 +370,7 @@ bool RecordType::isLayoutIdentical(const RecordType &other) { llvm::TypeSize BoolType::getTypeSizeInBits(const ::mlir::DataLayout &dataLayout, ::mlir::DataLayoutEntryListRef params) const { + assert(!cir::MissingFeatures::dataLayoutPtrHandlingBasedOnLangAS()); return llvm::TypeSize::getFixed(8); } From 0ead941adfb0978e5230f5e135cec4baf6575c50 Mon Sep 17 00:00:00 2001 From: David Rivera Date: Thu, 20 Nov 2025 10:52:38 -0500 Subject: [PATCH 11/13] Fix formatting --- clang/lib/CIR/CodeGen/CIRGenDecl.cpp | 4 ++-- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 9 ++++++--- clang/lib/CIR/CodeGen/TargetInfo.cpp | 5 +++-- clang/lib/CIR/Dialect/IR/CIRAttrs.cpp | 3 ++- clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 2 +- .../Transforms/TargetLowering/TargetLoweringInfo.h | 4 ++-- clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 6 +++--- 7 files changed, 19 insertions(+), 14 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp index df8591d386bf..a127489911b5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp @@ -479,8 +479,8 @@ CIRGenModule::getOrCreateStaticVarDecl(const VarDecl &D, Name = getStaticDeclName(*this, D); mlir::Type LTy = getTypes().convertTypeForMem(Ty); - mlir::Attribute AS = cir::toCIRClangAddressSpaceAttr(&getMLIRContext(), - getGlobalVarAddressSpace(&D)); + mlir::Attribute AS = cir::toCIRClangAddressSpaceAttr( + &getMLIRContext(), getGlobalVarAddressSpace(&D)); // OpenCL variables in local address space and CUDA shared // variables cannot have an initializer. diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 59299f34bff2..32c8ade68be6 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -1112,7 +1112,8 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty, entry = dyn_cast_or_null(v); } - mlir::Attribute cirAS = cir::toCIRClangAddressSpaceAttr(&getMLIRContext(), langAS); + mlir::Attribute cirAS = + cir::toCIRClangAddressSpaceAttr(&getMLIRContext(), langAS); if (entry) { mlir::Attribute entryCIRAS = entry.getAddrSpace(); if (WeakRefReferences.erase(entry)) { @@ -1168,7 +1169,8 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty, return entry; } - mlir::Attribute declCIRAS = cir::toCIRClangAddressSpaceAttr(&getMLIRContext(), getGlobalVarAddressSpace(d)); + mlir::Attribute declCIRAS = cir::toCIRClangAddressSpaceAttr( + &getMLIRContext(), getGlobalVarAddressSpace(d)); // TODO(cir): do we need to strip pointer casts for Entry? auto loc = getLoc(d->getSourceRange()); @@ -1987,7 +1989,8 @@ CIRGenModule::getAddrOfGlobalTemporary(const MaterializeTemporaryExpr *expr, linkage = cir::GlobalLinkageKind::InternalLinkage; } } - mlir::Attribute targetAS = cir::toCIRClangAddressSpaceAttr(&getMLIRContext(), addrSpace); + mlir::Attribute targetAS = + cir::toCIRClangAddressSpaceAttr(&getMLIRContext(), addrSpace); auto loc = getLoc(expr->getSourceRange()); auto gv = createGlobalOp(*this, loc, name, type, isConstant, targetAS, diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp index a4ddfbb69ccf..2fccb2a1754a 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.cpp +++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp @@ -267,8 +267,9 @@ class CommonSPIRTargetCIRGenInfo : public TargetCIRGenInfo { : TargetCIRGenInfo(std::move(ABIInfo)) {} mlir::Attribute getCIRAllocaAddressSpace() const override { - return cir::ClangAddressSpaceAttr::get(&getABIInfo().CGT.getMLIRContext(), - cir::ClangAddressSpace::OffloadPrivate); + return cir::ClangAddressSpaceAttr::get( + &getABIInfo().CGT.getMLIRContext(), + cir::ClangAddressSpace::OffloadPrivate); } cir::CallingConv getOpenCLKernelCallingConv() const override { diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp index b2b83da98a7f..c27604f2e39b 100644 --- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp @@ -87,7 +87,8 @@ mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, return mlir::success(); } -void printAddressSpaceValue(mlir::AsmPrinter &p, cir::ClangAddressSpace addrSpace) { +void printAddressSpaceValue(mlir::AsmPrinter &p, + cir::ClangAddressSpace addrSpace) { p << cir::stringifyEnum(addrSpace); } diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index 73dc2c01cb9f..abc9f1a8035f 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -1032,7 +1032,7 @@ void printGlobalAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp, } mlir::Attribute cir::toCIRClangAddressSpaceAttr(mlir::MLIRContext *ctx, - clang::LangAS langAS) { + clang::LangAS langAS) { using clang::LangAS; if (langAS == LangAS::Default) diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h index d5c6111b7ae7..4d8eb2f177c3 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h @@ -32,8 +32,8 @@ class TargetLoweringInfo { const ABIInfo &getABIInfo() const { return *Info; } - virtual unsigned - getTargetAddrSpaceFromCIRAddrSpace(cir::ClangAddressSpace addrSpace) const = 0; + virtual unsigned getTargetAddrSpaceFromCIRAddrSpace( + cir::ClangAddressSpace addrSpace) const = 0; }; } // namespace cir diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 7bbd74e0d04e..39c4253379a8 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -5030,9 +5030,9 @@ void prepareTypeConverter(mlir::LLVMTypeConverter &converter, converter.addConversion([&, lowerModule](cir::PointerType type) -> mlir::Type { mlir::Attribute addrSpaceAttr = type.getAddrSpace(); - unsigned addrSpace = addrSpaceAttr - ? getTargetAddrSpaceFromASAttr(addrSpaceAttr, lowerModule) - : 0; // Default address space + unsigned addrSpace = + addrSpaceAttr ? getTargetAddrSpaceFromASAttr(addrSpaceAttr, lowerModule) + : 0; // Default address space return mlir::LLVM::LLVMPointerType::get(type.getContext(), addrSpace); }); converter.addConversion([&](cir::VPtrType type) -> mlir::Type { From dfd52c218211198a2225a56ce915ba238a054cf9 Mon Sep 17 00:00:00 2001 From: David Rivera Date: Thu, 20 Nov 2025 11:04:48 -0500 Subject: [PATCH 12/13] Fix more formatting errors --- clang/include/clang/CIR/Dialect/IR/CIRTypes.h | 2 +- clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h index 804dfe1900fd..8cdff239538b 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h @@ -40,7 +40,7 @@ cir::ClangAddressSpace toCIRClangAddressSpace(clang::LangAS langAS); /// Returns ClangAddressSpaceAttr for clang/language-specific address spaces, /// or TargetAddressSpaceAttr for target-specific address spaces. mlir::Attribute toCIRClangAddressSpaceAttr(mlir::MLIRContext *ctx, - clang::LangAS langAS); + clang::LangAS langAS); constexpr unsigned getAsUnsignedValue(cir::ClangAddressSpace as) { return static_cast(as); diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index 4256b89a211b..eb8962d23e2c 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -55,9 +55,9 @@ using namespace mlir; #include "clang/CIR/Dialect/IR/CIROpsDialect.cpp.inc" #include "clang/CIR/Interfaces/ASTAttrInterfaces.h" #include "clang/CIR/Interfaces/CIROpInterfaces.h" -#include -#include #include +#include +#include //===----------------------------------------------------------------------===// // CIR Dialect From 0d731a7b239c0632e15a8297b854d49b9e4447d0 Mon Sep 17 00:00:00 2001 From: David Rivera Date: Thu, 20 Nov 2025 11:22:50 -0500 Subject: [PATCH 13/13] rename AS conversion fn for clarity --- clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 39c4253379a8..d6dcc3eb2bfa 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -5001,8 +5001,8 @@ std::unique_ptr prepareLowerModule(mlir::ModuleOp module) { } static unsigned -getTargetAddrSpaceFromCIRAddrSpaceAttr(cir::ClangAddressSpaceAttr addrSpace, - cir::LowerModule *lowerModule) { +convertCIRAddrSpaceToTarget(cir::ClangAddressSpaceAttr addrSpace, + cir::LowerModule *lowerModule) { assert(lowerModule && "CIR AS map is not available"); return lowerModule->getTargetLoweringInfo() .getTargetAddrSpaceFromCIRAddrSpace(addrSpace.getValue()); @@ -5019,7 +5019,7 @@ static unsigned getTargetAddrSpaceFromASAttr(mlir::Attribute attr, cir::ClangAddressSpaceAttr addrSpaceAttr = mlir::dyn_cast(attr); - return getTargetAddrSpaceFromCIRAddrSpaceAttr(addrSpaceAttr, lowerModule); + return convertCIRAddrSpaceToTarget(addrSpaceAttr, lowerModule); } // FIXME: change the type of lowerModule to `LowerModule &` to have better