diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h index d76115e03d8b..512e33c567e2 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" @@ -18,6 +19,7 @@ #include "clang/CIR/Dialect/IR/FPEnv.h" #include "clang/CIR/MissingFeatures.h" +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinAttributes.h" @@ -25,6 +27,7 @@ #include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/Location.h" #include "mlir/IR/Types.h" +#include "mlir/Support/LLVM.h" #include "llvm/ADT/APSInt.h" #include "llvm/ADT/ArrayRef.h" #include "llvm/ADT/FloatingPointMode.h" @@ -105,20 +108,30 @@ 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::ptr::MemorySpaceAttrInterface memorySpace) { + if (!memorySpace) + return cir::PointerType::get(ty); + return cir::PointerType::get(ty, memorySpace); } cir::PointerType getPointerTo(mlir::Type ty, clang::LangAS langAS) { - return getPointerTo(ty, cir::toCIRAddressSpace(langAS)); + if (langAS == clang::LangAS::Default) + return getPointerTo(ty); + + mlir::ptr::MemorySpaceAttrInterface addrSpaceAttr = + cir::toCIRLangAddressSpaceAttr(getContext(), langAS); + return getPointerTo(ty, addrSpaceAttr); } 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::ptr::MemorySpaceAttrInterface memorySpace) { + return getPointerTo(cir::VoidType::get(getContext()), memorySpace); } cir::MethodAttr getMethodAttr(cir::MethodType ty, cir::FuncOp methodFuncOp) { @@ -419,7 +432,8 @@ class CIRBaseBuilderTy : public mlir::OpBuilder { mlir::Value createGetGlobal(mlir::Location loc, cir::GlobalOp global, bool threadLocal = false) { return cir::GetGlobalOp::create( - *this, loc, getPointerTo(global.getSymType(), global.getAddrSpace()), + *this, loc, + getPointerTo(global.getSymType(), global.getAddrSpaceAttr()), global.getName(), threadLocal); } diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.h b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.h index 7f1e206e0cb4..97c4932feaf2 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.h +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.h @@ -13,6 +13,7 @@ #ifndef CLANG_CIR_DIALECT_IR_CIRATTRS_H #define CLANG_CIR_DIALECT_IR_CIRATTRS_H +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/BuiltinAttributeInterfaces.h" diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td index 249d31487684..ff0bfe5605e9 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -14,6 +14,7 @@ #define CLANG_CIR_DIALECT_IR_CIRATTRS_TD include "mlir/IR/BuiltinAttributeInterfaces.td" +include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.td" include "clang/CIR/Dialect/IR/CIREnumAttr.td" include "clang/CIR/Dialect/IR/CIRDialect.td" @@ -963,10 +964,29 @@ def CIR_DynamicCastInfoAttr : CIR_Attr<"DynamicCastInfo", "dyn_cast_info"> { // AddressSpaceAttr //===----------------------------------------------------------------------===// -def CIR_AddressSpaceAttr : CIR_EnumAttr { +def CIR_LangAddressSpaceAttr : + CIR_EnumAttr + ]> { + + let summary = "Represents a language 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::LangAddressSpace` 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 lang_address_space(offload_constant) + + ``` + }]; + let builders = [ AttrBuilder<(ins "clang::LangAS":$langAS), [{ - return $_get($_ctxt, cir::toCIRAddressSpace(langAS)); + return $_get($_ctxt, cir::toCIRLangAddressSpace(langAS)); }]> ]; @@ -974,12 +994,9 @@ def CIR_AddressSpaceAttr : CIR_EnumAttr { `(` custom($value) `)` }]; - let defaultValue = "cir::AddressSpace::Default"; + let defaultValue = "cir::LangAddressSpace::Default"; let extraClassDeclaration = [{ - bool isLang() const; - bool isTarget() const; - unsigned getTargetValue() const; unsigned getAsUnsignedValue() const; }]; @@ -987,19 +1004,36 @@ def CIR_AddressSpaceAttr : CIR_EnumAttr { unsigned $cppClass::getAsUnsignedValue() const { return static_cast(getValue()); } + }]; +} - bool $cppClass::isLang() const { - return cir::isLangAddressSpace(getValue()); - } +//===----------------------------------------------------------------------===// +// TargetAddressSpaceAttr +//===----------------------------------------------------------------------===// - bool $cppClass::isTarget() const { - return cir::isTargetAddressSpace(getValue()); - } +def CIR_TargetAddressSpaceAttr : CIR_Attr< "TargetAddressSpace", + "target_address_space", [ + DeclareAttrInterfaceMethods + ]> { + 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 + `address_space` attribute. + + A value of zero represents the default address space. The semantics of non-zero + address spaces are target-specific. - unsigned $cppClass::getTargetValue() const { - return cir::getTargetAddressSpaceValue(getValue()); - } + Example: + ```mlir + // Target-specific numeric address spaces + !cir.ptr + !cir.ptr + ``` }]; + + let parameters = (ins "unsigned":$value); + let assemblyFormat = "`<` `target` `<` $value `>` `>`"; } //===----------------------------------------------------------------------===// diff --git a/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td b/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td index 46c6039674b0..da61f7be058e 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td +++ b/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td @@ -41,15 +41,14 @@ class CIR_DefaultValuedEnumParameter let defaultValue = value; } -def CIR_AddressSpace : CIR_I32EnumAttr< - "AddressSpace", "address space kind", [ +def CIR_LangAddressSpace : CIR_I32EnumAttr< + "LangAddressSpace", "language address space kind", [ I32EnumAttrCase<"Default", 0, "default">, I32EnumAttrCase<"OffloadPrivate", 1, "offload_private">, 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 @@ -58,7 +57,7 @@ def CIR_AddressSpace : CIR_I32EnumAttr< The `value` parameter is an extensible enum, which encodes target address space as an offset to the last language address space. For that reason, the - attribute is implemented as custom AddressSpaceAttr, which provides custom + attribute is implemented as custom LangAddressSpaceAttr, which provides custom printer and parser for the `value` parameter. }]; diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index 0c117bd6d636..7ffc3d54df2b 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -2618,10 +2618,7 @@ def CIR_GlobalOp : CIR_Op<"global", [ OptionalAttr:$sym_visibility, TypeAttr:$sym_type, CIR_GlobalLinkageKind:$linkage, - DefaultValuedAttr< - CIR_AddressSpaceAttr, - "AddressSpace::Default" - >:$addr_space, + OptionalAttr:$addr_space, OptionalAttr:$tls_model, // Note this can also be a FlatSymbolRefAttr OptionalAttr:$initial_value, @@ -2644,7 +2641,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^)? @@ -2668,7 +2665,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::ptr::MemorySpaceAttrInterface", "{}">:$addrSpace, CArg<"llvm::function_ref", "nullptr">:$ctorBuilder, CArg<"llvm::function_ref", diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h index b7cf95d08ca8..efcfeb7bf60f 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.h +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.h @@ -13,6 +13,7 @@ #ifndef CLANG_CIR_DIALECT_IR_CIRTYPES_H #define CLANG_CIR_DIALECT_IR_CIRTYPES_H +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" #include "mlir/IR/BuiltinAttributes.h" #include "mlir/IR/Types.h" #include "mlir/Interfaces/DataLayoutInterfaces.h" @@ -34,37 +35,18 @@ bool isSized(mlir::Type ty); // AddressSpace helpers //===----------------------------------------------------------------------===// -cir::AddressSpace toCIRAddressSpace(clang::LangAS langAS); +cir::LangAddressSpace toCIRLangAddressSpace(clang::LangAS langAS); -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(); -} +/// Convert a LangAS to the appropriate address space attribute interface. +/// Returns a MemorySpaceAttrInterface. +mlir::ptr::MemorySpaceAttrInterface +toCIRLangAddressSpaceAttr(mlir::MLIRContext *ctx, clang::LangAS langAS); -constexpr bool isLangAddressSpace(cir::AddressSpace as) { - return !isTargetAddressSpace(as); -} +bool isSupportedCIRMemorySpaceAttr( + mlir::ptr::MemorySpaceAttrInterface memorySpace); -constexpr unsigned getTargetAddressSpaceValue(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); +constexpr unsigned getAsUnsignedValue(cir::LangAddressSpace as) { + return static_cast(as); } } // namespace cir diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td index 33abc3ffc66f..6c0eaf919a8b 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::ptr::MemorySpaceAttrInterface">:$addrSpace ); let skipDefaultBuilders = 1; let builders = [ TypeBuilderWithInferredContext<(ins "mlir::Type":$pointee, - CArg<"cir::AddressSpace", "cir::AddressSpace::Default">:$addrSpace), [{ + CArg<"mlir::ptr::MemorySpaceAttrInterface", "{}">:$addrSpace), [{ return $_get(pointee.getContext(), pointee, addrSpace); }]>, TypeBuilder<(ins "mlir::Type":$pointee, - CArg<"cir::AddressSpace", "cir::AddressSpace::Default">:$addrSpace), [{ + CArg<"mlir::ptr::MemorySpaceAttrInterface", "{}">:$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)^ )? `>` }]; diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h index 2ffd2853ae22..2e8f04e7e16d 100644 --- a/clang/include/clang/CIR/MissingFeatures.h +++ b/clang/include/clang/CIR/MissingFeatures.h @@ -480,6 +480,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/CodeGen/CIRGenBuilder.h b/clang/lib/CIR/CodeGen/CIRGenBuilder.h index 5b8391c9d408..7a8c79bf1c0b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuilder.h +++ b/clang/lib/CIR/CodeGen/CIRGenBuilder.h @@ -25,6 +25,7 @@ #include "clang/CIR/Dialect/IR/CIRTypes.h" #include "clang/CIR/Dialect/IR/FPEnv.h" +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinAttributes.h" @@ -749,7 +750,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::ptr::MemorySpaceAttrInterface addrSpace = {}) { mlir::OpBuilder::InsertionGuard guard(*this); setInsertionPointToStart(module.getBody()); return cir::GlobalOp::create(*this, loc, name, type, isConst, linkage, @@ -759,10 +760,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::ptr::MemorySpaceAttrInterface 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/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 14b76bb8b06a..9c3b92aa5b8f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -34,6 +34,8 @@ #include "clang/Frontend/FrontendDiagnostic.h" #include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" +#include "mlir/IR/Attributes.h" #include "mlir/IR/BuiltinAttributes.h" #include "mlir/IR/Value.h" #include "mlir/Support/LLVM.h" @@ -1727,9 +1729,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::ptr::MemorySpaceAttrInterface AAS = getCIRAllocaAddressSpace(); + mlir::ptr::MemorySpaceAttrInterface EAS = cir::toCIRLangAddressSpaceAttr( + &getMLIRContext(), E->getType()->getPointeeType().getAddressSpace()); if (EAS != AAS) { assert(false && "Non-default address space for alloca NYI"); } diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp index 4377d27d438d..b4722c9ea814 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp @@ -335,7 +335,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 ec5124489109..7d7715241d7e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp @@ -15,10 +15,12 @@ #include "CIRGenFunction.h" #include "CIRGenOpenMPRuntime.h" #include "EHScopeStack.h" +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/BuiltinAttributeInterfaces.h" #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/SymbolTable.h" +#include "mlir/Support/LLVM.h" #include "clang/AST/Decl.h" #include "clang/AST/ExprCXX.h" @@ -481,7 +483,8 @@ CIRGenModule::getOrCreateStaticVarDecl(const VarDecl &D, Name = getStaticDeclName(*this, D); mlir::Type LTy = getTypes().convertTypeForMem(Ty); - cir::AddressSpace AS = cir::toCIRAddressSpace(getGlobalVarAddressSpace(&D)); + mlir::ptr::MemorySpaceAttrInterface AS = cir::toCIRLangAddressSpaceAttr( + &getMLIRContext(), getGlobalVarAddressSpace(&D)); // OpenCL variables in local address space and CUDA shared // variables cannot have an initializer. @@ -595,7 +598,7 @@ cir::GlobalOp CIRGenFunction::addInitializerToStaticVarDecl( // Given those constraints, thread in the GetGlobalOp and update it // directly. GVAddr.getAddr().setType( - getBuilder().getPointerTo(Init.getType(), GV.getAddrSpace())); + getBuilder().getPointerTo(Init.getType(), GV.getAddrSpaceAttr())); } bool NeedsDtor = diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index 4db465464e4a..194789192ed3 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -40,6 +40,7 @@ #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/IR/Operation.h" #include "mlir/IR/Value.h" +#include "mlir/Support/LLVM.h" #include @@ -2081,10 +2082,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::toCIRLangAddressSpaceAttr( + &getMLIRContext(), E->getSubExpr()->getType().getAddressSpace()); + mlir::Attribute DestAS = cir::toCIRLangAddressSpaceAttr( + &getMLIRContext(), E->getType().getAddressSpace()); mlir::Value V = getTargetHooks().performAddrSpaceCast( *this, LV.getPointer(), SrcAS, DestAS, convertType(DestTy)); return makeAddrLValue(Address(V, convertTypeForMem(E->getType()), @@ -3144,7 +3145,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::toCIRLangAddressSpaceAttr( + &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 0875cc11d05a..359176574200 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp @@ -28,6 +28,8 @@ #include "llvm/Support/ErrorHandling.h" #include +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" +#include "mlir/IR/Attributes.h" #include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/Value.h" @@ -1606,8 +1608,8 @@ mlir::Value ScalarExprEmitter::emitSub(const BinOpInfo &Ops) { cir::PointerType rhsPtrTy = mlir::dyn_cast(rhs.getType()); if (lhsPtrTy && rhsPtrTy) { - cir::AddressSpace lhsAS = lhsPtrTy.getAddrSpace(); - cir::AddressSpace rhsAS = rhsPtrTy.getAddrSpace(); + mlir::ptr::MemorySpaceAttrInterface lhsAS = lhsPtrTy.getAddrSpace(); + mlir::ptr::MemorySpaceAttrInterface rhsAS = rhsPtrTy.getAddrSpace(); if (lhsAS != rhsAS) { // Different address spaces → use addrspacecast @@ -1881,10 +1883,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::toCIRLangAddressSpaceAttr( + &CGF.getMLIRContext(), E->getType()->getPointeeType().getAddressSpace()); - cir::AddressSpace DestAS = - cir::toCIRAddressSpace(DestTy->getPointeeType().getAddressSpace()); + mlir::Attribute DestAS = cir::toCIRLangAddressSpaceAttr( + &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 a45e6820cbf9..56ba4b95114f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -21,6 +21,7 @@ #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinAttributeInterfaces.h" @@ -812,12 +813,10 @@ mlir::Value CIRGenModule::getGlobalValue(const Decl *d) { return CurCGF->symbolTable.lookup(d); } -cir::GlobalOp CIRGenModule::createGlobalOp(CIRGenModule &cgm, - mlir::Location loc, StringRef name, - mlir::Type t, bool isConstant, - cir::AddressSpace addrSpace, - mlir::Operation *insertPoint, - cir::GlobalLinkageKind linkage) { +cir::GlobalOp CIRGenModule::createGlobalOp( + CIRGenModule &cgm, mlir::Location loc, StringRef name, mlir::Type t, + bool isConstant, mlir::ptr::MemorySpaceAttrInterface addrSpace, + mlir::Operation *insertPoint, cir::GlobalLinkageKind linkage) { cir::GlobalOp g; auto &builder = cgm.getBuilder(); { @@ -1098,9 +1097,10 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty, entry = dyn_cast_or_null(v); } - cir::AddressSpace cirAS = cir::toCIRAddressSpace(langAS); + mlir::Attribute cirAS = + cir::toCIRLangAddressSpaceAttr(&getMLIRContext(), langAS); if (entry) { - cir::AddressSpace entryCIRAS = entry.getAddrSpace(); + mlir::ptr::MemorySpaceAttrInterface entryCIRAS = entry.getAddrSpaceAttr(); if (WeakRefReferences.erase(entry)) { if (d && !d->hasAttr()) { auto lt = cir::GlobalLinkageKind::ExternalLinkage; @@ -1154,7 +1154,9 @@ CIRGenModule::getOrCreateCIRGlobal(StringRef mangledName, mlir::Type ty, return entry; } - auto declCIRAS = cir::toCIRAddressSpace(getGlobalVarAddressSpace(d)); + mlir::ptr::MemorySpaceAttrInterface declCIRAS = + cir::toCIRLangAddressSpaceAttr(&getMLIRContext(), + getGlobalVarAddressSpace(d)); // TODO(cir): do we need to strip pointer casts for Entry? auto loc = getLoc(d->getSourceRange()); @@ -1275,7 +1277,7 @@ mlir::Value CIRGenModule::getAddrOfGlobalVar(const VarDecl *d, mlir::Type ty, bool tlsAccess = d->getTLSKind() != VarDecl::TLS_None; auto g = getOrCreateCIRGlobal(d, ty, isForDefinition); - auto ptrTy = builder.getPointerTo(g.getSymType(), g.getAddrSpace()); + auto ptrTy = builder.getPointerTo(g.getSymType(), g.getAddrSpaceAttr()); return cir::GetGlobalOp::create(builder, getLoc(d->getSourceRange()), ptrTy, g.getSymName(), tlsAccess); } @@ -1290,7 +1292,7 @@ CIRGenModule::getAddrOfGlobalVarAttr(const VarDecl *d, mlir::Type ty, auto globalOp = getOrCreateCIRGlobal(d, ty, isForDefinition); auto ptrTy = - builder.getPointerTo(globalOp.getSymType(), globalOp.getAddrSpace()); + builder.getPointerTo(globalOp.getSymType(), globalOp.getAddrSpaceAttr()); return builder.getGlobalViewAttr(ptrTy, globalOp); } @@ -1763,8 +1765,9 @@ 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::ptr::MemorySpaceAttrInterface addrSpace = + cir::toCIRLangAddressSpaceAttr(&cgm.getMLIRContext(), + cgm.getGlobalConstantAddressSpace()); // Create a global variable for this string // FIXME(cir): check for insertion point in module level. @@ -1869,8 +1872,8 @@ CIRGenModule::getAddrOfConstantStringFromLiteral(const StringLiteral *s, auto gv = getGlobalForStringLiteral(s, name); auto arrayTy = mlir::dyn_cast(gv.getSymType()); assert(arrayTy && "String literal must be array"); - auto ptrTy = - getBuilder().getPointerTo(arrayTy.getElementType(), gv.getAddrSpace()); + auto ptrTy = getBuilder().getPointerTo(arrayTy.getElementType(), + gv.getAddrSpaceAttr()); return builder.getGlobalViewAttr(ptrTy, gv); } @@ -1980,7 +1983,8 @@ CIRGenModule::getAddrOfGlobalTemporary(const MaterializeTemporaryExpr *expr, linkage = cir::GlobalLinkageKind::InternalLinkage; } } - cir::AddressSpace targetAS = cir::toCIRAddressSpace(addrSpace); + mlir::ptr::MemorySpaceAttrInterface targetAS = + cir::toCIRLangAddressSpaceAttr(&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 52282041a944..796d34c3c964 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -282,7 +282,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, + mlir::ptr::MemorySpaceAttrInterface addrSpace = {}, mlir::Operation *insertPoint = nullptr, cir::GlobalLinkageKind linkage = cir::GlobalLinkageKind::ExternalLinkage); diff --git a/clang/lib/CIR/CodeGen/CIRGenTypeCache.h b/clang/lib/CIR/CodeGen/CIRGenTypeCache.h index c83d60673f23..fb18ba1cd593 100644 --- a/clang/lib/CIR/CodeGen/CIRGenTypeCache.h +++ b/clang/lib/CIR/CodeGen/CIRGenTypeCache.h @@ -13,6 +13,8 @@ #ifndef LLVM_CLANG_LIB_CIR_CODEGENTYPECACHE_H #define LLVM_CLANG_LIB_CIR_CODEGENTYPECACHE_H +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" +#include "mlir/IR/Attributes.h" #include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/Types.h" #include "clang/AST/CharUnits.h" @@ -106,7 +108,7 @@ struct CIRGenTypeCache { unsigned char SizeAlignInBytes; }; - cir::AddressSpace CIRAllocaAddressSpace; + mlir::ptr::MemorySpaceAttrInterface CIRAllocaAddressSpace; clang::CharUnits getSizeSize() const { return clang::CharUnits::fromQuantity(SizeSizeInBytes); @@ -121,7 +123,7 @@ struct CIRGenTypeCache { return clang::CharUnits::fromQuantity(PointerAlignInBytes); } - cir::AddressSpace getCIRAllocaAddressSpace() const { + mlir::ptr::MemorySpaceAttrInterface getCIRAllocaAddressSpace() const { return CIRAllocaAddressSpace; } }; diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp index d7e57cfc5a87..a2974a9a6dbc 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.cpp +++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp @@ -4,6 +4,7 @@ #include "CIRGenCXXABI.h" #include "CIRGenFunctionInfo.h" #include "CIRGenTypes.h" +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" #include "clang/Basic/TargetInfo.h" #include "clang/CIR/ABIArgInfo.h" @@ -266,8 +267,11 @@ class CommonSPIRTargetCIRGenInfo : public TargetCIRGenInfo { CommonSPIRTargetCIRGenInfo(std::unique_ptr ABIInfo) : TargetCIRGenInfo(std::move(ABIInfo)) {} - cir::AddressSpace getCIRAllocaAddressSpace() const override { - return cir::AddressSpace::OffloadPrivate; + mlir::ptr::MemorySpaceAttrInterface + getCIRAllocaAddressSpace() const override { + return cir::LangAddressSpaceAttr::get( + &getABIInfo().CGT.getMLIRContext(), + cir::LangAddressSpace::OffloadPrivate); } cir::CallingConv getOpenCLKernelCallingConv() const override { @@ -662,8 +666,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..ab4a2362870d 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -16,6 +16,8 @@ #include "ABIInfo.h" #include "CIRGenValue.h" +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" +#include "mlir/IR/Attributes.h" #include "mlir/IR/Types.h" #include "clang/CIR/Dialect/IR/CIRAttrs.h" #include "clang/CIR/Target/AArch64.h" @@ -101,8 +103,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::ptr::MemorySpaceAttrInterface getCIRAllocaAddressSpace() const { + return {}; // Empty attribute represents the Default address space } /// Perform address space cast of an expression of pointer type. @@ -112,8 +114,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/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp index af9483d47a0d..99bcb3006c6f 100644 --- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp @@ -16,6 +16,7 @@ #include "clang/CIR/Dialect/IR/CIRTypes.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinAttributeInterfaces.h" @@ -77,9 +78,20 @@ static void printConstPtr(mlir::AsmPrinter &p, mlir::IntegerAttr value); //===----------------------------------------------------------------------===// mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, - cir::AddressSpace &addrSpace); + cir::LangAddressSpace &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::LangAddressSpace addrSpace) { + p << cir::stringifyEnum(addrSpace); +} //===----------------------------------------------------------------------===// // Tablegen defined attributes @@ -91,6 +103,110 @@ void printAddressSpaceValue(mlir::AsmPrinter &p, cir::AddressSpace addrSpace); using namespace mlir; using namespace cir; +//===----------------------------------------------------------------------===// +// MemorySpaceAttrInterface implementations for Clang and Target address space +// attributes +//===----------------------------------------------------------------------===// +namespace cir { + +bool LangAddressSpaceAttr::isValidLoad( + mlir::Type type, mlir::ptr::AtomicOrdering ordering, + std::optional alignment, const mlir::DataLayout *dataLayout, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +bool LangAddressSpaceAttr::isValidStore( + mlir::Type type, mlir::ptr::AtomicOrdering ordering, + std::optional alignment, const mlir::DataLayout *dataLayout, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +bool LangAddressSpaceAttr::isValidAtomicOp( + mlir::ptr::AtomicBinOp op, mlir::Type type, + mlir::ptr::AtomicOrdering ordering, std::optional alignment, + const mlir::DataLayout *dataLayout, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +bool LangAddressSpaceAttr::isValidAtomicXchg( + mlir::Type type, mlir::ptr::AtomicOrdering successOrdering, + mlir::ptr::AtomicOrdering failureOrdering, std::optional alignment, + const mlir::DataLayout *dataLayout, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +bool LangAddressSpaceAttr::isValidAddrSpaceCast( + mlir::Type tgt, mlir::Type src, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +bool LangAddressSpaceAttr::isValidPtrIntCast( + mlir::Type intLikeTy, mlir::Type ptrLikeTy, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +bool TargetAddressSpaceAttr::isValidLoad( + mlir::Type type, mlir::ptr::AtomicOrdering ordering, + std::optional alignment, const mlir::DataLayout *dataLayout, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +bool TargetAddressSpaceAttr::isValidStore( + mlir::Type type, mlir::ptr::AtomicOrdering ordering, + std::optional alignment, const mlir::DataLayout *dataLayout, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +bool TargetAddressSpaceAttr::isValidAtomicOp( + mlir::ptr::AtomicBinOp op, mlir::Type type, + mlir::ptr::AtomicOrdering ordering, std::optional alignment, + const mlir::DataLayout *dataLayout, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +bool TargetAddressSpaceAttr::isValidAtomicXchg( + mlir::Type type, mlir::ptr::AtomicOrdering successOrdering, + mlir::ptr::AtomicOrdering failureOrdering, std::optional alignment, + const mlir::DataLayout *dataLayout, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +bool TargetAddressSpaceAttr::isValidAddrSpaceCast( + mlir::Type tgt, mlir::Type src, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +bool TargetAddressSpaceAttr::isValidPtrIntCast( + mlir::Type intLikeTy, mlir::Type ptrLikeTy, + llvm::function_ref emitError) const { + assert(false && "NYI"); + return false; +} + +} // namespace cir + //===----------------------------------------------------------------------===// // CIR AST Attr helpers //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index 4a4d098e8663..74e97ced9e68 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -28,6 +28,8 @@ #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/LLVMIR/LLVMTypes.h" +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" +#include "mlir/IR/Attributes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinAttributes.h" #include "mlir/IR/BuiltinTypes.h" @@ -54,6 +56,8 @@ 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 //===----------------------------------------------------------------------===// @@ -304,6 +308,13 @@ static void printOmittedTerminatorRegion(mlir::OpAsmPrinter &printer, } } +mlir::OptionalParseResult +parseGlobalAddressSpaceValue(mlir::AsmParser &p, + mlir::ptr::MemorySpaceAttrInterface &attr); + +void printGlobalAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp op, + mlir::ptr::MemorySpaceAttrInterface attr); + //===----------------------------------------------------------------------===// // AllocaOp //===----------------------------------------------------------------------===// @@ -2414,7 +2425,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::ptr::MemorySpaceAttrInterface addrSpace, function_ref ctorBuilder, function_ref dtorBuilder) { odsState.addAttribute(getSymNameAttrName(odsState.name), @@ -2429,9 +2440,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) { @@ -2495,10 +2505,10 @@ cir::GetGlobalOp::verifySymbolUses(SymbolTableCollection &symbolTable) { << "' does not reference a valid cir.global or cir.func"; mlir::Type symTy; - cir::AddressSpace symAddrSpace{}; + mlir::ptr::MemorySpaceAttrInterface symAddrSpaceAttr{}; if (auto g = dyn_cast(op)) { symTy = g.getSymType(); - symAddrSpace = g.getAddrSpace(); + symAddrSpaceAttr = g.getAddrSpaceAttr(); // Verify that for thread local global access, the global needs to // be marked with tls bits. if (getTls() && !g.getTlsModel()) @@ -2514,7 +2524,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..1148f98054d9 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -11,12 +11,14 @@ //===----------------------------------------------------------------------===// #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" #include "clang/CIR/Dialect/IR/CIRTypesDetails.h" #include "clang/CIR/MissingFeatures.h" +#include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/BuiltinAttributes.h" #include "mlir/IR/BuiltinTypes.h" @@ -32,6 +34,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 @@ -67,10 +70,12 @@ static void printFuncTypeParams(mlir::AsmPrinter &p, // AddressSpace //===----------------------------------------------------------------------===// -mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, - cir::AddressSpace &addrSpace); +mlir::ParseResult +parseAddressSpaceValue(mlir::AsmParser &p, + mlir::ptr::MemorySpaceAttrInterface &attr); -void printAddressSpaceValue(mlir::AsmPrinter &p, cir::AddressSpace addrSpace); +void printAddressSpaceValue(mlir::AsmPrinter &printer, + mlir::ptr::MemorySpaceAttrInterface attr); //===----------------------------------------------------------------------===// // Get autogenerated stuff @@ -368,6 +373,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); } @@ -781,8 +787,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); } //===----------------------------------------------------------------------===// @@ -920,27 +926,33 @@ MethodType::getABIAlignment(const mlir::DataLayout &dataLayout, // AddressSpace definitions //===----------------------------------------------------------------------===// -cir::AddressSpace cir::toCIRAddressSpace(clang::LangAS langAS) { +bool cir::isSupportedCIRMemorySpaceAttr( + mlir::ptr::MemorySpaceAttrInterface memorySpace) { + return mlir::isa( + memorySpace); +} + +cir::LangAddressSpace cir::toCIRLangAddressSpace(clang::LangAS langAS) { using clang::LangAS; switch (langAS) { case LangAS::Default: - return AddressSpace::Default; + return LangAddressSpace::Default; case LangAS::opencl_global: - return AddressSpace::OffloadGlobal; + return LangAddressSpace::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 LangAddressSpace::OffloadLocal; case LangAS::cuda_device: - return AddressSpace::OffloadGlobal; + return LangAddressSpace::OffloadGlobal; case LangAS::opencl_constant: case LangAS::cuda_constant: - return AddressSpace::OffloadConstant; + return LangAddressSpace::OffloadConstant; case LangAS::opencl_private: - return AddressSpace::OffloadPrivate; + return LangAddressSpace::OffloadPrivate; case LangAS::opencl_generic: - return AddressSpace::OffloadGeneric; + return LangAddressSpace::OffloadGeneric; case LangAS::opencl_global_device: case LangAS::opencl_global_host: case LangAS::sycl_global: @@ -955,53 +967,112 @@ 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()); + llvm_unreachable("unknown/unsupported clang language address space"); } } -mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p, - cir::AddressSpace &addrSpace) { +mlir::ParseResult +parseAddressSpaceValue(mlir::AsmParser &p, + mlir::ptr::MemorySpaceAttrInterface &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"); - - // 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. - 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()) + + // 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"); - addrSpace = cir::computeTargetAddressSpace(targetValue); - } else { - addrSpace = result.value(); + attr = cir::TargetAddressSpaceAttr::get(p.getContext(), val); + return mlir::success(); + } + + // Try to parse language specific address space. + if (p.parseOptionalKeyword("lang_address_space").succeeded()) { + if (p.parseLParen()) + 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 language address space keyword"); + + attr = cir::LangAddressSpaceAttr::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 &p, + mlir::ptr::MemorySpaceAttrInterface attr) { + if (!attr) + return; + + if (auto logical = dyn_cast(attr)) { + p << "lang_address_space(" + << cir::stringifyLangAddressSpace(logical.getValue()) << ')'; + + return; + } + + if (auto target = dyn_cast(attr)) { + p << "target_address_space(" << target.getValue() << ')'; + return; + } + + llvm_unreachable("unexpected address-space attribute kind"); +} + +mlir::OptionalParseResult +parseGlobalAddressSpaceValue(mlir::AsmParser &p, + mlir::ptr::MemorySpaceAttrInterface &attr) { + + mlir::SMLoc loc = p.getCurrentLocation(); + if (parseAddressSpaceValue(p, attr).failed()) + return p.emitError(loc, "failed to parse Address Space Value for GlobalOp"); + return mlir::success(); +} + +void printGlobalAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp, + mlir::ptr::MemorySpaceAttrInterface attr) { + printAddressSpaceValue(printer, attr); +} + +mlir::ptr::MemorySpaceAttrInterface +cir::toCIRLangAddressSpaceAttr(mlir::MLIRContext *ctx, clang::LangAS langAS) { + using clang::LangAS; + + if (langAS == LangAS::Default) + return {}; // Default address space is represented as an empty attribute. + + if (clang::isTargetAddressSpace(langAS)) { + unsigned targetAS = clang::toTargetAddressSpace(langAS); + return cir::TargetAddressSpaceAttr::get(ctx, targetAS); + } + + return cir::LangAddressSpaceAttr::get(ctx, toCIRLangAddressSpace(langAS)); +} + +//===----------------------------------------------------------------------===// +// PointerType Definitions +//===----------------------------------------------------------------------===// + +mlir::LogicalResult cir::PointerType::verify( + llvm::function_ref emitError, + mlir::Type pointee, mlir::ptr::MemorySpaceAttrInterface addrSpace) { + if (addrSpace) { + if (!isSupportedCIRMemorySpaceAttr(addrSpace)) { + return emitError() << "CIR Address spaces must be either target_address " + "space or lang_address_space"; + } + } + + return success(); } //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h index 114d8cc0f697..99e464c6bbd6 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h @@ -19,6 +19,7 @@ #include #include "clang/CIR/Dialect/IR/CIRAttrs.h" +#include "clang/CIR/Dialect/IR/CIROpsEnums.h" namespace cir { @@ -33,7 +34,7 @@ class TargetLoweringInfo { const ABIInfo &getABIInfo() const { return *Info; } virtual unsigned - getTargetAddrSpaceFromCIRAddrSpace(cir::AddressSpace addrSpace) const = 0; + getTargetAddrSpaceFromCIRAddrSpace(cir::LangAddressSpace 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..8e7716e9069a 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::LangAddressSpace 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::LangAddressSpace::OffloadPrivate: + case cir::LangAddressSpace::OffloadLocal: + case cir::LangAddressSpace::OffloadGlobal: + case cir::LangAddressSpace::OffloadConstant: + case cir::LangAddressSpace::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..76d457f1607b 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp @@ -12,6 +12,7 @@ #include "TargetInfo.h" #include "TargetLoweringInfo.h" #include "clang/CIR/ABIArgInfo.h" +#include "clang/CIR/Dialect/IR/CIROpsEnums.h" #include "clang/CIR/Dialect/IR/CIRTypes.h" #include "clang/CIR/MissingFeatures.h" #include "llvm/Support/Casting.h" @@ -44,17 +45,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::LangAddressSpace addrSpace) const override { switch (addrSpace) { - case cir::AddressSpace::OffloadPrivate: + case cir::LangAddressSpace::OffloadPrivate: return 5; - case cir::AddressSpace::OffloadLocal: + case cir::LangAddressSpace::OffloadLocal: return 3; - case cir::AddressSpace::OffloadGlobal: + case cir::LangAddressSpace::OffloadGlobal: return 1; - case cir::AddressSpace::OffloadConstant: + case cir::LangAddressSpace::OffloadConstant: return 4; - case cir::AddressSpace::OffloadGeneric: + case cir::LangAddressSpace::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..66cbb1aad63c 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp @@ -12,6 +12,7 @@ #include "TargetInfo.h" #include "TargetLoweringInfo.h" #include "clang/CIR/ABIArgInfo.h" +#include "clang/CIR/Dialect/IR/CIROpsEnums.h" #include "clang/CIR/Dialect/IR/CIRTypes.h" #include "clang/CIR/MissingFeatures.h" #include "llvm/Support/Casting.h" @@ -45,17 +46,17 @@ class NVPTXTargetLoweringInfo : public TargetLoweringInfo { : TargetLoweringInfo(std::make_unique(lt)) {} unsigned getTargetAddrSpaceFromCIRAddrSpace( - cir::AddressSpace addrSpace) const override { + cir::LangAddressSpace addrSpace) const override { switch (addrSpace) { - case cir::AddressSpace::OffloadPrivate: + case cir::LangAddressSpace::OffloadPrivate: return 0; - case cir::AddressSpace::OffloadLocal: + case cir::LangAddressSpace::OffloadLocal: return 3; - case cir::AddressSpace::OffloadGlobal: + case cir::LangAddressSpace::OffloadGlobal: return 1; - case cir::AddressSpace::OffloadConstant: + case cir::LangAddressSpace::OffloadConstant: return 4; - case cir::AddressSpace::OffloadGeneric: + case cir::LangAddressSpace::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..2497e6768094 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIR.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIR.cpp @@ -12,6 +12,7 @@ #include "TargetInfo.h" #include "TargetLoweringInfo.h" #include "clang/CIR/ABIArgInfo.h" +#include "clang/CIR/Dialect/IR/CIROpsEnums.h" #include "clang/CIR/MissingFeatures.h" #include "llvm/Support/ErrorHandling.h" @@ -42,17 +43,17 @@ class SPIRVTargetLoweringInfo : public TargetLoweringInfo { : TargetLoweringInfo(std::make_unique(LT)) {} unsigned getTargetAddrSpaceFromCIRAddrSpace( - cir::AddressSpace addrSpace) const override { + cir::LangAddressSpace addrSpace) const override { switch (addrSpace) { - case cir::AddressSpace::OffloadPrivate: + case cir::LangAddressSpace::OffloadPrivate: return 0; - case cir::AddressSpace::OffloadLocal: + case cir::LangAddressSpace::OffloadLocal: return 3; - case cir::AddressSpace::OffloadGlobal: + case cir::LangAddressSpace::OffloadGlobal: return 1; - case cir::AddressSpace::OffloadConstant: + case cir::LangAddressSpace::OffloadConstant: return 2; - case cir::AddressSpace::OffloadGeneric: + case cir::LangAddressSpace::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..f2106da8109b 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::LangAddressSpace 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::LangAddressSpace::OffloadPrivate: + case cir::LangAddressSpace::OffloadLocal: + case cir::LangAddressSpace::OffloadGlobal: + case cir::LangAddressSpace::OffloadConstant: + case cir::LangAddressSpace::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 d2efc8075119..c94ae0f34bec 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -387,7 +387,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(); } @@ -5014,18 +5014,25 @@ 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); - +static unsigned convertCIRAddrSpaceToTarget(cir::LangAddressSpaceAttr 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, + cir::LowerModule *lowerModule) { + assert(mlir::isa_and_nonnull(attr) || + mlir::isa_and_nonnull(attr)); + + if (auto targetAddrSpaceAttr = + mlir::dyn_cast(attr)) + return targetAddrSpaceAttr.getValue(); + + cir::LangAddressSpaceAttr addrSpaceAttr = + mlir::dyn_cast(attr); + return convertCIRAddrSpaceToTarget(addrSpaceAttr, lowerModule); } // FIXME: change the type of lowerModule to `LowerModule &` to have better @@ -5035,8 +5042,10 @@ void prepareTypeConverter(mlir::LLVMTypeConverter &converter, cir::LowerModule *lowerModule) { converter.addConversion([&, lowerModule](cir::PointerType type) -> mlir::Type { + mlir::Attribute addrSpaceAttr = type.getAddrSpace(); unsigned addrSpace = - getTargetAddrSpaceFromCIRAddrSpace(type.getAddrSpace(), lowerModule); + addrSpaceAttr ? getTargetAddrSpaceFromASAttr(addrSpaceAttr, lowerModule) + : 0; // Default address space return mlir::LLVM::LLVMPointerType::get(type.getContext(), addrSpace); }); converter.addConversion([&](cir::VPtrType type) -> mlir::Type { diff --git a/clang/test/CIR/CodeGen/CUDA/address-spaces.cu b/clang/test/CIR/CodeGen/CUDA/address-spaces.cu index b25977f1305a..fc4fc831a8d5 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 lang_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..81ad8794dc4e 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 lang_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 lang_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 lang_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..61425a388099 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 lang_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..cb0b1a03dfef 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 lang_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..6caa76a8c795 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 lang_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/HIP/ptr-diff.cpp b/clang/test/CIR/CodeGen/HIP/ptr-diff.cpp index 10cb3832b00a..b647c40c3c0b 100644 --- a/clang/test/CIR/CodeGen/HIP/ptr-diff.cpp +++ b/clang/test/CIR/CodeGen/HIP/ptr-diff.cpp @@ -23,14 +23,14 @@ __device__ int ptr_diff() { // CIR-DEVICE: %[[#LenLocalAddr:]] = cir.alloca !cir.ptr, !cir.ptr>, ["len", init] -// CIR-DEVICE: %[[#GlobalPtr:]] = cir.get_global @_ZZ8ptr_diffvE5c_str : !cir.ptr, addrspace(offload_constant)> -// CIR-DEVICE: %[[#CastDecay:]] = cir.cast array_to_ptrdecay %[[#GlobalPtr]] : !cir.ptr, addrspace(offload_constant)> -// CIR-DEVICE: %[[#LenLocalAddrCast:]] = cir.cast bitcast %[[#LenLocalAddr]] : !cir.ptr> -> !cir.ptr> -// CIR-DEVICE: cir.store align(8) %[[#CastDecay]], %[[#LenLocalAddrCast]] : !cir.ptr, !cir.ptr> -// CIR-DEVICE: %[[#CStr:]] = cir.cast array_to_ptrdecay %[[#GlobalPtr]] : !cir.ptr, addrspace(offload_constant)> -> !cir.ptr +// CIR-DEVICE: %[[#GlobalPtr:]] = cir.get_global @_ZZ8ptr_diffvE5c_str : !cir.ptr, lang_address_space(offload_constant)> +// CIR-DEVICE: %[[#CastDecay:]] = cir.cast array_to_ptrdecay %[[#GlobalPtr]] : !cir.ptr, lang_address_space(offload_constant)> +// CIR-DEVICE: %[[#LenLocalAddrCast:]] = cir.cast bitcast %[[#LenLocalAddr]] : !cir.ptr> -> !cir.ptr> +// CIR-DEVICE: cir.store align(8) %[[#CastDecay]], %[[#LenLocalAddrCast]] : !cir.ptr, !cir.ptr> +// CIR-DEVICE: %[[#CStr:]] = cir.cast array_to_ptrdecay %[[#GlobalPtr]] : !cir.ptr, lang_address_space(offload_constant)> -> !cir.ptr // CIR-DEVICE: %[[#LoadedLenAddr:]] = cir.load align(8) %[[#LenLocalAddr]] : !cir.ptr>, !cir.ptr loc(#loc7) -// CIR-DEVICE: %[[#AddrCast:]] = cir.cast address_space %[[#LoadedLenAddr]] : !cir.ptr -> !cir.ptr -// CIR-DEVICE: %[[#DIFF:]] = cir.ptr_diff %[[#CStr]], %[[#AddrCast]] : !cir.ptr +// CIR-DEVICE: %[[#AddrCast:]] = cir.cast address_space %[[#LoadedLenAddr]] : !cir.ptr -> !cir.ptr +// CIR-DEVICE: %[[#DIFF:]] = cir.ptr_diff %[[#CStr]], %[[#AddrCast]] : !cir.ptr // LLVM-DEVICE: define dso_local i32 @_Z8ptr_diffv() // LLVM-DEVICE: %[[#GlobalPtrAddr:]] = alloca i32, i64 1, align 4, addrspace(5) diff --git a/clang/test/CIR/CodeGen/OpenCL/addrspace-alloca.cl b/clang/test/CIR/CodeGen/OpenCL/addrspace-alloca.cl index a1e0eb0950db..d4740560b497 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, lang_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, lang_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, lang_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, lang_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, lang_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..6cfcc83ec1d2 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, lang_address_space(offload_local)> -> !cir.ptr + // CIR-NEXT: cir.store{{.*}} %{{[0-9]+}}, %{{[0-9]+}} : !cir.ptr, !cir.ptr, lang_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, lang_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/as_type.cl b/clang/test/CIR/CodeGen/OpenCL/as_type.cl index 6fc8104e8a24..cb07dcdb1174 100644 --- a/clang/test/CIR/CodeGen/OpenCL/as_type.cl +++ b/clang/test/CIR/CodeGen/OpenCL/as_type.cl @@ -10,7 +10,7 @@ typedef __attribute__(( ext_vector_type(4) )) char char4; // CIR: cir.func @f4(%{{.*}}: !s32i loc({{.*}})) -> !cir.vector -// CIR: %[[x:.*]] = cir.load align(4) %{{.*}} : !cir.ptr +// CIR: %[[x:.*]] = cir.load align(4) %{{.*}} : !cir.ptr // CIR: cir.cast bitcast %[[x]] : !s32i -> !cir.vector // LLVM: define spir_func <4 x i8> @f4(i32 %[[x:.*]]) // LLVM: %[[astype:.*]] = bitcast i32 %[[x]] to <4 x i8> @@ -25,7 +25,7 @@ char4 f4(int x) { } // CIR: cir.func @f6(%{{.*}}: !cir.vector loc({{.*}})) -> !s32i -// CIR: %[[x:.*]] = cir.load align(4) %{{.*}} : !cir.ptr, addrspace(offload_private)>, !cir.vector +// CIR: %[[x:.*]] = cir.load align(4) %{{.*}} : !cir.ptr, lang_address_space(offload_private)>, !cir.vector // CIR: cir.cast bitcast %[[x]] : !cir.vector -> !s32i // LLVM: define{{.*}} spir_func i32 @f6(<4 x i8> %[[x:.*]]) // LLVM: %[[astype:.*]] = bitcast <4 x i8> %[[x]] to i32 @@ -39,9 +39,9 @@ int f6(char4 x) { return __builtin_astype(x, int); } -// CIR: cir.func @f4_ptr(%{{.*}}: !cir.ptr loc({{.*}})) -> !cir.ptr, addrspace(offload_local)> -// CIR: %[[x:.*]] = cir.load align(8) %{{.*}} : !cir.ptr, addrspace(offload_private)>, !cir.ptr -// CIR: cir.cast address_space %[[x]] : !cir.ptr -> !cir.ptr, addrspace(offload_local)> +// CIR: cir.func @f4_ptr(%{{.*}}: !cir.ptr loc({{.*}})) -> !cir.ptr, lang_address_space(offload_local)> +// CIR: %[[x:.*]] = cir.load align(8) %{{.*}} : !cir.ptr, lang_address_space(offload_private)>, !cir.ptr +// CIR: cir.cast address_space %[[x]] : !cir.ptr -> !cir.ptr, lang_address_space(offload_local)> // LLVM: define spir_func ptr addrspace(3) @f4_ptr(ptr addrspace(1) readnone captures(ret: address, provenance) %[[x:.*]]) // LLVM: %[[astype:.*]] = addrspacecast ptr addrspace(1) %[[x]] to ptr addrspace(3) // LLVM-NOT: shufflevector diff --git a/clang/test/CIR/CodeGen/OpenCL/global.cl b/clang/test/CIR/CodeGen/OpenCL/global.cl index ef5c63e58f83..bb5afa8b7e9d 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 lang_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 lang_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 {{.*}}lang_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..5e803b64d2fc 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..0e3885dc138f 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 lang_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 lang_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..81001b155d79 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 {{.*}}lang_address_space(offload_constant) @".str" = #cir.const_array<"hello world\00" : !cir.array> : !cir.array +// CIR: cir.global{{.*}} constant {{.*}}lang_address_space(offload_constant) @x = #cir.global_view<@".str"> : !cir.ptr +// CIR: cir.global{{.*}} constant {{.*}}lang_address_space(offload_constant) @y = #cir.global_view<@".str"> : !cir.ptr +// CIR: cir.global{{.*}} constant {{.*}}lang_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, lang_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..0be43e07fdcd 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..50fdcddf1610 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 lang_address_space(offload_global) @addrspace1 = #cir.int<1> : !s32i + cir.global "private" internal lang_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 lang_address_space(offload_global) @addrspace1 = #cir.int<1> : !s32i +// CHECK: cir.global "private" internal lang_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 70846ac264cd..2a22641496d3 100644 --- a/clang/test/CIR/IR/invalid.cir +++ b/clang/test/CIR/IR/invalid.cir @@ -301,8 +301,8 @@ cir.func @cast24(%p : !u32i) { // ----- !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 } @@ -310,7 +310,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 } @@ -1223,9 +1223,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 language address space keyword}} +// expected-error@below {{expected keyword for language address space kind}} +cir.func @address_space1(%p : !cir.ptr) { cir.return } @@ -1234,24 +1234,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] for language address space kind, got: foobar}} +// expected-error@below {{expected language address space keyword}} +cir.func @address_space4(%p : !cir.ptr) { cir.return } @@ -1360,7 +1353,7 @@ module { !s32i = !cir.int module { - cir.global external addrspace(offload_global) @gv = #cir.int<0> : !s32i + cir.global external lang_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}} @@ -1375,9 +1368,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, lang_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, lang_address_space(offload_private)> -> !cir.ptr cir.return } } @@ -1406,7 +1399,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..65874cf82dff 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 lang_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 lang_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