diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/CMakeLists.txt b/mlir/include/mlir/Dialect/SPIRV/IR/CMakeLists.txt --- a/mlir/include/mlir/Dialect/SPIRV/IR/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/SPIRV/IR/CMakeLists.txt @@ -31,8 +31,8 @@ add_public_tablegen_target(MLIRSPIRVAttrUtilsGen) add_dependencies(mlir-headers MLIRSPIRVAttrUtilsGen) -set(LLVM_TARGET_DEFINITIONS TargetAndABI.td) -mlir_tablegen(TargetAndABI.h.inc -gen-struct-attr-decls) -mlir_tablegen(TargetAndABI.cpp.inc -gen-struct-attr-defs) -add_public_tablegen_target(MLIRSPIRVTargetAndABIIncGen) -add_dependencies(mlir-headers MLIRSPIRVTargetAndABIIncGen) +set(LLVM_TARGET_DEFINITIONS SPIRVAttributes.td) +mlir_tablegen(SPIRVAttributes.h.inc -gen-attrdef-decls) +mlir_tablegen(SPIRVAttributes.cpp.inc -gen-attrdef-defs) +add_public_tablegen_target(MLIRSPIRVAttributeIncGen) +add_dependencies(mlir-headers MLIRSPIRVAttributeIncGen) diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.h b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.h --- a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.h +++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.h @@ -18,7 +18,8 @@ #include "mlir/Support/LLVM.h" // Pull in TableGen'erated SPIR-V attribute definitions for target and ABI. -#include "mlir/Dialect/SPIRV/IR/TargetAndABI.h.inc" +#define GET_ATTRDEF_CLASSES +#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h.inc" namespace mlir { namespace spirv { @@ -139,7 +140,7 @@ /// Gets a TargetEnvAttr instance. static TargetEnvAttr get(VerCapExtAttr triple, Vendor vendorID, DeviceType deviceType, uint32_t deviceId, - DictionaryAttr limits); + ResourceLimitsAttr limits); /// Returns the attribute kind's name (without the 'spv.' prefix). static StringRef getKindName(); @@ -171,11 +172,6 @@ /// Returns the target resource limits. ResourceLimitsAttr getResourceLimits() const; - - static LogicalResult verify(function_ref emitError, - VerCapExtAttr triple, Vendor vendorID, - DeviceType deviceType, uint32_t deviceID, - DictionaryAttr limits); }; } // namespace spirv } // namespace mlir diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/TargetAndABI.td b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.td rename from mlir/include/mlir/Dialect/SPIRV/IR/TargetAndABI.td rename to mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.td --- a/mlir/include/mlir/Dialect/SPIRV/IR/TargetAndABI.td +++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.td @@ -26,9 +26,11 @@ // For entry functions, this attribute specifies information related to entry // points in the generated SPIR-V module: // 1) WorkGroup Size. -def SPV_EntryPointABIAttr : StructAttr<"EntryPointABIAttr", SPIRV_Dialect, [ - StructFieldAttr<"local_size", OptionalAttr> -]>; +def SPV_EntryPointABIAttr : AttrDef { + let parameters = (ins OptionalParameter<"DenseIntElementsAttr">:$local_size); + let mnemonic = "entry_point_abi"; + let assemblyFormat = "`<` struct(params) `>`"; +} def SPV_ExtensionArrayAttr : TypedArrayAttrBase< SPV_ExtensionAttr, "SPIR-V extension array attribute">; @@ -40,16 +42,20 @@ // target. Represents `VkCooperativeMatrixPropertiesNV`. See // https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkCooperativeMatrixPropertiesNV.html def SPV_CooperativeMatrixPropertiesNVAttr : - StructAttr<"CooperativeMatrixPropertiesNVAttr", SPIRV_Dialect, [ - StructFieldAttr<"m_size", I32Attr>, - StructFieldAttr<"n_size", I32Attr>, - StructFieldAttr<"k_size", I32Attr>, - StructFieldAttr<"a_type", TypeAttr>, - StructFieldAttr<"b_type", TypeAttr>, - StructFieldAttr<"c_type", TypeAttr>, - StructFieldAttr<"result_type", TypeAttr>, - StructFieldAttr<"scope", SPV_ScopeAttr> -]>; + AttrDef { + let parameters = (ins + "int":$m_size, + "int":$n_size, + "int":$k_size, + "mlir::Type":$a_type, + "mlir::Type":$b_type, + "mlir::Type":$c_type, + "mlir::Type":$result_type, + "mlir::spirv::ScopeAttr":$scope + ); + let mnemonic = "coop_matrix_props"; + let assemblyFormat = "`<` struct(params) `>`"; +} def SPV_CooperativeMatrixPropertiesNVArrayAttr : TypedArrayAttrBase { + let parameters = (ins // The maximum total storage size, in bytes, available for variables // declared with the Workgroup storage class. - StructFieldAttr<"max_compute_shared_memory_size", - DefaultValuedAttr>, + DefaultValuedParameter<"int", "16384">:$max_compute_shared_memory_size, // The maximum total number of compute shader invocations in a single local // workgroup. - StructFieldAttr<"max_compute_workgroup_invocations", - DefaultValuedAttr>, + DefaultValuedParameter<"int", "128">:$max_compute_workgroup_invocations, // The maximum size of a local compute workgroup, per dimension. - StructFieldAttr<"max_compute_workgroup_size", - DefaultValuedAttr>, + DefaultValuedParameter< + "ArrayAttr", + "$_builder.getI32ArrayAttr({128, 128, 64})" + >:$max_compute_workgroup_size, // The default number of invocations in each subgroup. - StructFieldAttr<"subgroup_size", DefaultValuedAttr>, + DefaultValuedParameter<"int", "32">:$subgroup_size, // The configurations of cooperative matrix operations // supported. Default is an empty list. - StructFieldAttr< - "cooperative_matrix_properties_nv", - DefaultValuedAttr> -]>; + DefaultValuedParameter< + "ArrayAttr", + "nullptr" + >:$cooperative_matrix_properties_nv + ); + let mnemonic = "resource_limits"; + let assemblyFormat = "`<` struct(params) `>`"; +} #endif // MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVBase.td b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVBase.td --- a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVBase.td +++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVBase.td @@ -48,7 +48,6 @@ let cppNamespace = "::mlir::spirv"; let useDefaultTypePrinterParser = 1; - let useDefaultAttributePrinterParser = 1; let hasConstantMaterializer = 1; let hasOperationAttrVerify = 1; let hasRegionArgAttrVerify = 1; @@ -65,6 +64,13 @@ /// Returns the attribute name to use when specifying decorations on results /// of operations. static std::string getAttributeName(Decoration decoration); + + /// Dialect attribute parsing hook. + Attribute parseAttribute( + DialectAsmParser &parser, Type type) const override; + /// Dialect attribute printing hook. + void printAttribute( + Attribute attr, DialectAsmPrinter &printer) const override; }]; } diff --git a/mlir/lib/Dialect/SPIRV/IR/CMakeLists.txt b/mlir/lib/Dialect/SPIRV/IR/CMakeLists.txt --- a/mlir/lib/Dialect/SPIRV/IR/CMakeLists.txt +++ b/mlir/lib/Dialect/SPIRV/IR/CMakeLists.txt @@ -16,13 +16,13 @@ ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/SPIRV DEPENDS + MLIRSPIRVAttributeIncGen MLIRSPIRVAttrUtilsGen MLIRSPIRVAvailabilityIncGen MLIRSPIRVCanonicalizationIncGen MLIRSPIRVEnumAvailabilityIncGen MLIRSPIRVEnumsIncGen MLIRSPIRVOpsIncGen - MLIRSPIRVTargetAndABIIncGen LINK_LIBS PUBLIC MLIRControlFlowInterfaces diff --git a/mlir/lib/Dialect/SPIRV/IR/SPIRVAttributes.cpp b/mlir/lib/Dialect/SPIRV/IR/SPIRVAttributes.cpp --- a/mlir/lib/Dialect/SPIRV/IR/SPIRVAttributes.cpp +++ b/mlir/lib/Dialect/SPIRV/IR/SPIRVAttributes.cpp @@ -10,8 +10,11 @@ #include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h" #include "mlir/Dialect/SPIRV/IR/SPIRVTypes.h" #include "mlir/IR/Builders.h" +#include "mlir/IR/DialectImplementation.h" +#include "llvm/ADT/TypeSwitch.h" using namespace mlir; +using namespace mlir::spirv; //===----------------------------------------------------------------------===// // TableGen'erated attribute utility functions @@ -21,15 +24,6 @@ namespace spirv { #include "mlir/Dialect/SPIRV/IR/SPIRVAttrUtils.inc" } // namespace spirv -} // namespace mlir - -//===----------------------------------------------------------------------===// -// DictionaryDict derived attributes -//===----------------------------------------------------------------------===// - -#include "mlir/Dialect/SPIRV/IR/TargetAndABI.cpp.inc" - -namespace mlir { //===----------------------------------------------------------------------===// // Attribute storage classes @@ -292,7 +286,7 @@ Vendor vendorID, DeviceType deviceType, uint32_t deviceID, - DictionaryAttr limits) { + ResourceLimitsAttr limits) { assert(triple && limits && "expected valid triple and limits"); MLIRContext *context = triple.getContext(); return Base::get(context, triple, vendorID, deviceType, deviceID, limits); @@ -340,16 +334,332 @@ return getImpl()->limits.cast(); } -LogicalResult -spirv::TargetEnvAttr::verify(function_ref emitError, - spirv::VerCapExtAttr /*triple*/, - spirv::Vendor /*vendorID*/, - spirv::DeviceType /*deviceType*/, - uint32_t /*deviceID*/, DictionaryAttr limits) { - if (!limits.isa()) - return emitError() << "expected spirv::ResourceLimitsAttr for limits"; +//===----------------------------------------------------------------------===// +// ODS Generated Attributes +//===----------------------------------------------------------------------===// - return success(); +#define GET_ATTRDEF_CLASSES +#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.cpp.inc" + +//===----------------------------------------------------------------------===// +// Attribute Parsing +//===----------------------------------------------------------------------===// + +/// Parses a comma-separated list of keywords, invokes `processKeyword` on each +/// of the parsed keyword, and returns failure if any error occurs. +static ParseResult +parseKeywordList(DialectAsmParser &parser, + function_ref processKeyword) { + if (parser.parseLSquare()) + return failure(); + + // Special case for empty list. + if (succeeded(parser.parseOptionalRSquare())) + return success(); + + // Keep parsing the keyword and an optional comma following it. If the comma + // is successfully parsed, then we have more keywords to parse. + if (failed(parser.parseCommaSeparatedList([&]() { + auto loc = parser.getCurrentLocation(); + StringRef keyword; + if (parser.parseKeyword(&keyword) || + failed(processKeyword(loc, keyword))) + return failure(); + return success(); + }))) + return failure(); + return parser.parseRSquare(); +} + +/// Parses a spirv::InterfaceVarABIAttr. +static Attribute parseInterfaceVarABIAttr(DialectAsmParser &parser) { + if (parser.parseLess()) + return {}; + + Builder &builder = parser.getBuilder(); + + if (parser.parseLParen()) + return {}; + + IntegerAttr descriptorSetAttr; + { + auto loc = parser.getCurrentLocation(); + uint32_t descriptorSet = 0; + auto descriptorSetParseResult = parser.parseOptionalInteger(descriptorSet); + + if (!descriptorSetParseResult.hasValue() || + failed(*descriptorSetParseResult)) { + parser.emitError(loc, "missing descriptor set"); + return {}; + } + descriptorSetAttr = builder.getI32IntegerAttr(descriptorSet); + } + + if (parser.parseComma()) + return {}; + + IntegerAttr bindingAttr; + { + auto loc = parser.getCurrentLocation(); + uint32_t binding = 0; + auto bindingParseResult = parser.parseOptionalInteger(binding); + + if (!bindingParseResult.hasValue() || failed(*bindingParseResult)) { + parser.emitError(loc, "missing binding"); + return {}; + } + bindingAttr = builder.getI32IntegerAttr(binding); + } + + if (parser.parseRParen()) + return {}; + + IntegerAttr storageClassAttr; + { + if (succeeded(parser.parseOptionalComma())) { + auto loc = parser.getCurrentLocation(); + StringRef storageClass; + if (parser.parseKeyword(&storageClass)) + return {}; + + if (auto storageClassSymbol = + spirv::symbolizeStorageClass(storageClass)) { + storageClassAttr = builder.getI32IntegerAttr( + static_cast(*storageClassSymbol)); + } else { + parser.emitError(loc, "unknown storage class: ") << storageClass; + return {}; + } + } + } + + if (parser.parseGreater()) + return {}; + + return spirv::InterfaceVarABIAttr::get(descriptorSetAttr, bindingAttr, + storageClassAttr); +} + +static Attribute parseVerCapExtAttr(DialectAsmParser &parser) { + if (parser.parseLess()) + return {}; + + Builder &builder = parser.getBuilder(); + + IntegerAttr versionAttr; + { + auto loc = parser.getCurrentLocation(); + StringRef version; + if (parser.parseKeyword(&version) || parser.parseComma()) + return {}; + + if (auto versionSymbol = spirv::symbolizeVersion(version)) { + versionAttr = + builder.getI32IntegerAttr(static_cast(*versionSymbol)); + } else { + parser.emitError(loc, "unknown version: ") << version; + return {}; + } + } + + ArrayAttr capabilitiesAttr; + { + SmallVector capabilities; + SMLoc errorloc; + StringRef errorKeyword; + + auto processCapability = [&](SMLoc loc, StringRef capability) { + if (auto capSymbol = spirv::symbolizeCapability(capability)) { + capabilities.push_back( + builder.getI32IntegerAttr(static_cast(*capSymbol))); + return success(); + } + return errorloc = loc, errorKeyword = capability, failure(); + }; + if (parseKeywordList(parser, processCapability) || parser.parseComma()) { + if (!errorKeyword.empty()) + parser.emitError(errorloc, "unknown capability: ") << errorKeyword; + return {}; + } + + capabilitiesAttr = builder.getArrayAttr(capabilities); + } + + ArrayAttr extensionsAttr; + { + SmallVector extensions; + SMLoc errorloc; + StringRef errorKeyword; + + auto processExtension = [&](SMLoc loc, StringRef extension) { + if (spirv::symbolizeExtension(extension)) { + extensions.push_back(builder.getStringAttr(extension)); + return success(); + } + return errorloc = loc, errorKeyword = extension, failure(); + }; + if (parseKeywordList(parser, processExtension)) { + if (!errorKeyword.empty()) + parser.emitError(errorloc, "unknown extension: ") << errorKeyword; + return {}; + } + + extensionsAttr = builder.getArrayAttr(extensions); + } + + if (parser.parseGreater()) + return {}; + + return spirv::VerCapExtAttr::get(versionAttr, capabilitiesAttr, + extensionsAttr); +} + +/// Parses a spirv::TargetEnvAttr. +static Attribute parseTargetEnvAttr(DialectAsmParser &parser) { + if (parser.parseLess()) + return {}; + + spirv::VerCapExtAttr tripleAttr; + if (parser.parseAttribute(tripleAttr) || parser.parseComma()) + return {}; + + // Parse [vendor[:device-type[:device-id]]] + Vendor vendorID = Vendor::Unknown; + DeviceType deviceType = DeviceType::Unknown; + uint32_t deviceID = spirv::TargetEnvAttr::kUnknownDeviceID; + { + auto loc = parser.getCurrentLocation(); + StringRef vendorStr; + if (succeeded(parser.parseOptionalKeyword(&vendorStr))) { + if (auto vendorSymbol = spirv::symbolizeVendor(vendorStr)) { + vendorID = *vendorSymbol; + } else { + parser.emitError(loc, "unknown vendor: ") << vendorStr; + } + + if (succeeded(parser.parseOptionalColon())) { + loc = parser.getCurrentLocation(); + StringRef deviceTypeStr; + if (parser.parseKeyword(&deviceTypeStr)) + return {}; + if (auto deviceTypeSymbol = spirv::symbolizeDeviceType(deviceTypeStr)) { + deviceType = *deviceTypeSymbol; + } else { + parser.emitError(loc, "unknown device type: ") << deviceTypeStr; + } + + if (succeeded(parser.parseOptionalColon())) { + loc = parser.getCurrentLocation(); + if (parser.parseInteger(deviceID)) + return {}; + } + } + if (parser.parseComma()) + return {}; + } + } + + ResourceLimitsAttr limitsAttr; + if (parser.parseAttribute(limitsAttr) || parser.parseGreater()) + return {}; + + return spirv::TargetEnvAttr::get(tripleAttr, vendorID, deviceType, deviceID, + limitsAttr); +} + +Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser, + Type type) const { + // SPIR-V attributes are dictionaries so they do not have type. + if (type) { + parser.emitError(parser.getNameLoc(), "unexpected type"); + return {}; + } + + // Parse the kind keyword first. + StringRef attrKind; + if (parser.parseKeyword(&attrKind)) + return {}; + + Attribute attr; + OptionalParseResult result = + generatedAttributeParser(parser, attrKind, type, attr); + if (result.hasValue()) { + if (failed(result.getValue())) + return {}; + return attr; + } + + if (attrKind == spirv::TargetEnvAttr::getKindName()) + return parseTargetEnvAttr(parser); + if (attrKind == spirv::VerCapExtAttr::getKindName()) + return parseVerCapExtAttr(parser); + if (attrKind == spirv::InterfaceVarABIAttr::getKindName()) + return parseInterfaceVarABIAttr(parser); + + parser.emitError(parser.getNameLoc(), "unknown SPIR-V attribute kind: ") + << attrKind; + return {}; +} + +//===----------------------------------------------------------------------===// +// Attribute Printing +//===----------------------------------------------------------------------===// + +static void print(spirv::VerCapExtAttr triple, DialectAsmPrinter &printer) { + auto &os = printer.getStream(); + printer << spirv::VerCapExtAttr::getKindName() << "<" + << spirv::stringifyVersion(triple.getVersion()) << ", ["; + llvm::interleaveComma( + triple.getCapabilities(), os, + [&](spirv::Capability cap) { os << spirv::stringifyCapability(cap); }); + printer << "], ["; + llvm::interleaveComma(triple.getExtensionsAttr(), os, [&](Attribute attr) { + os << attr.cast().getValue(); + }); + printer << "]>"; +} + +static void print(spirv::TargetEnvAttr targetEnv, DialectAsmPrinter &printer) { + printer << spirv::TargetEnvAttr::getKindName() << "<#spv."; + print(targetEnv.getTripleAttr(), printer); + spirv::Vendor vendorID = targetEnv.getVendorID(); + spirv::DeviceType deviceType = targetEnv.getDeviceType(); + uint32_t deviceID = targetEnv.getDeviceID(); + if (vendorID != spirv::Vendor::Unknown) { + printer << ", " << spirv::stringifyVendor(vendorID); + if (deviceType != spirv::DeviceType::Unknown) { + printer << ":" << spirv::stringifyDeviceType(deviceType); + if (deviceID != spirv::TargetEnvAttr::kUnknownDeviceID) + printer << ":" << deviceID; + } + } + printer << ", " << targetEnv.getResourceLimits() << ">"; +} + +static void print(spirv::InterfaceVarABIAttr interfaceVarABIAttr, + DialectAsmPrinter &printer) { + printer << spirv::InterfaceVarABIAttr::getKindName() << "<(" + << interfaceVarABIAttr.getDescriptorSet() << ", " + << interfaceVarABIAttr.getBinding() << ")"; + auto storageClass = interfaceVarABIAttr.getStorageClass(); + if (storageClass) + printer << ", " << spirv::stringifyStorageClass(*storageClass); + printer << ">"; +} + +void SPIRVDialect::printAttribute(Attribute attr, + DialectAsmPrinter &printer) const { + if (succeeded(generatedAttributePrinter(attr, printer))) + return; + + if (auto targetEnv = attr.dyn_cast()) + print(targetEnv, printer); + else if (auto vceAttr = attr.dyn_cast()) + print(vceAttr, printer); + else if (auto interfaceVarABIAttr = attr.dyn_cast()) + print(interfaceVarABIAttr, printer); + else + llvm_unreachable("unhandled SPIR-V attribute kind"); } //===----------------------------------------------------------------------===// @@ -358,4 +668,8 @@ void spirv::SPIRVDialect::registerAttributes() { addAttributes(); + addAttributes< +#define GET_ATTRDEF_LIST +#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.cpp.inc" + >(); } diff --git a/mlir/lib/Dialect/SPIRV/IR/SPIRVDialect.cpp b/mlir/lib/Dialect/SPIRV/IR/SPIRVDialect.cpp --- a/mlir/lib/Dialect/SPIRV/IR/SPIRVDialect.cpp +++ b/mlir/lib/Dialect/SPIRV/IR/SPIRVDialect.cpp @@ -870,330 +870,6 @@ .Default([](Type) { llvm_unreachable("unhandled SPIR-V type"); }); } -//===----------------------------------------------------------------------===// -// Attribute Parsing -//===----------------------------------------------------------------------===// - -/// Parses a comma-separated list of keywords, invokes `processKeyword` on each -/// of the parsed keyword, and returns failure if any error occurs. -static ParseResult parseKeywordList( - DialectAsmParser &parser, - function_ref processKeyword) { - if (parser.parseLSquare()) - return failure(); - - // Special case for empty list. - if (succeeded(parser.parseOptionalRSquare())) - return success(); - - // Keep parsing the keyword and an optional comma following it. If the comma - // is successfully parsed, then we have more keywords to parse. - if (failed(parser.parseCommaSeparatedList([&]() { - auto loc = parser.getCurrentLocation(); - StringRef keyword; - if (parser.parseKeyword(&keyword) || - failed(processKeyword(loc, keyword))) - return failure(); - return success(); - }))) - return failure(); - return parser.parseRSquare(); -} - -/// Parses a spirv::InterfaceVarABIAttr. -static Attribute parseInterfaceVarABIAttr(DialectAsmParser &parser) { - if (parser.parseLess()) - return {}; - - Builder &builder = parser.getBuilder(); - - if (parser.parseLParen()) - return {}; - - IntegerAttr descriptorSetAttr; - { - auto loc = parser.getCurrentLocation(); - uint32_t descriptorSet = 0; - auto descriptorSetParseResult = parser.parseOptionalInteger(descriptorSet); - - if (!descriptorSetParseResult.hasValue() || - failed(*descriptorSetParseResult)) { - parser.emitError(loc, "missing descriptor set"); - return {}; - } - descriptorSetAttr = builder.getI32IntegerAttr(descriptorSet); - } - - if (parser.parseComma()) - return {}; - - IntegerAttr bindingAttr; - { - auto loc = parser.getCurrentLocation(); - uint32_t binding = 0; - auto bindingParseResult = parser.parseOptionalInteger(binding); - - if (!bindingParseResult.hasValue() || failed(*bindingParseResult)) { - parser.emitError(loc, "missing binding"); - return {}; - } - bindingAttr = builder.getI32IntegerAttr(binding); - } - - if (parser.parseRParen()) - return {}; - - IntegerAttr storageClassAttr; - { - if (succeeded(parser.parseOptionalComma())) { - auto loc = parser.getCurrentLocation(); - StringRef storageClass; - if (parser.parseKeyword(&storageClass)) - return {}; - - if (auto storageClassSymbol = - spirv::symbolizeStorageClass(storageClass)) { - storageClassAttr = builder.getI32IntegerAttr( - static_cast(*storageClassSymbol)); - } else { - parser.emitError(loc, "unknown storage class: ") << storageClass; - return {}; - } - } - } - - if (parser.parseGreater()) - return {}; - - return spirv::InterfaceVarABIAttr::get(descriptorSetAttr, bindingAttr, - storageClassAttr); -} - -static Attribute parseVerCapExtAttr(DialectAsmParser &parser) { - if (parser.parseLess()) - return {}; - - Builder &builder = parser.getBuilder(); - - IntegerAttr versionAttr; - { - auto loc = parser.getCurrentLocation(); - StringRef version; - if (parser.parseKeyword(&version) || parser.parseComma()) - return {}; - - if (auto versionSymbol = spirv::symbolizeVersion(version)) { - versionAttr = - builder.getI32IntegerAttr(static_cast(*versionSymbol)); - } else { - parser.emitError(loc, "unknown version: ") << version; - return {}; - } - } - - ArrayAttr capabilitiesAttr; - { - SmallVector capabilities; - SMLoc errorloc; - StringRef errorKeyword; - - auto processCapability = [&](SMLoc loc, StringRef capability) { - if (auto capSymbol = spirv::symbolizeCapability(capability)) { - capabilities.push_back( - builder.getI32IntegerAttr(static_cast(*capSymbol))); - return success(); - } - return errorloc = loc, errorKeyword = capability, failure(); - }; - if (parseKeywordList(parser, processCapability) || parser.parseComma()) { - if (!errorKeyword.empty()) - parser.emitError(errorloc, "unknown capability: ") << errorKeyword; - return {}; - } - - capabilitiesAttr = builder.getArrayAttr(capabilities); - } - - ArrayAttr extensionsAttr; - { - SmallVector extensions; - SMLoc errorloc; - StringRef errorKeyword; - - auto processExtension = [&](SMLoc loc, StringRef extension) { - if (spirv::symbolizeExtension(extension)) { - extensions.push_back(builder.getStringAttr(extension)); - return success(); - } - return errorloc = loc, errorKeyword = extension, failure(); - }; - if (parseKeywordList(parser, processExtension)) { - if (!errorKeyword.empty()) - parser.emitError(errorloc, "unknown extension: ") << errorKeyword; - return {}; - } - - extensionsAttr = builder.getArrayAttr(extensions); - } - - if (parser.parseGreater()) - return {}; - - return spirv::VerCapExtAttr::get(versionAttr, capabilitiesAttr, - extensionsAttr); -} - -/// Parses a spirv::TargetEnvAttr. -static Attribute parseTargetEnvAttr(DialectAsmParser &parser) { - if (parser.parseLess()) - return {}; - - spirv::VerCapExtAttr tripleAttr; - if (parser.parseAttribute(tripleAttr) || parser.parseComma()) - return {}; - - // Parse [vendor[:device-type[:device-id]]] - Vendor vendorID = Vendor::Unknown; - DeviceType deviceType = DeviceType::Unknown; - uint32_t deviceID = spirv::TargetEnvAttr::kUnknownDeviceID; - { - auto loc = parser.getCurrentLocation(); - StringRef vendorStr; - if (succeeded(parser.parseOptionalKeyword(&vendorStr))) { - if (auto vendorSymbol = spirv::symbolizeVendor(vendorStr)) { - vendorID = *vendorSymbol; - } else { - parser.emitError(loc, "unknown vendor: ") << vendorStr; - } - - if (succeeded(parser.parseOptionalColon())) { - loc = parser.getCurrentLocation(); - StringRef deviceTypeStr; - if (parser.parseKeyword(&deviceTypeStr)) - return {}; - if (auto deviceTypeSymbol = spirv::symbolizeDeviceType(deviceTypeStr)) { - deviceType = *deviceTypeSymbol; - } else { - parser.emitError(loc, "unknown device type: ") << deviceTypeStr; - } - - if (succeeded(parser.parseOptionalColon())) { - loc = parser.getCurrentLocation(); - if (parser.parseInteger(deviceID)) - return {}; - } - } - if (parser.parseComma()) - return {}; - } - } - - DictionaryAttr limitsAttr; - { - auto loc = parser.getCurrentLocation(); - if (parser.parseAttribute(limitsAttr)) - return {}; - - if (!limitsAttr.isa()) { - parser.emitError( - loc, - "limits must be a dictionary attribute containing two 32-bit integer " - "attributes 'max_compute_workgroup_invocations' and " - "'max_compute_workgroup_size'"); - return {}; - } - } - - if (parser.parseGreater()) - return {}; - - return spirv::TargetEnvAttr::get(tripleAttr, vendorID, deviceType, deviceID, - limitsAttr); -} - -Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser, - Type type) const { - // SPIR-V attributes are dictionaries so they do not have type. - if (type) { - parser.emitError(parser.getNameLoc(), "unexpected type"); - return {}; - } - - // Parse the kind keyword first. - StringRef attrKind; - if (parser.parseKeyword(&attrKind)) - return {}; - - if (attrKind == spirv::TargetEnvAttr::getKindName()) - return parseTargetEnvAttr(parser); - if (attrKind == spirv::VerCapExtAttr::getKindName()) - return parseVerCapExtAttr(parser); - if (attrKind == spirv::InterfaceVarABIAttr::getKindName()) - return parseInterfaceVarABIAttr(parser); - - parser.emitError(parser.getNameLoc(), "unknown SPIR-V attribute kind: ") - << attrKind; - return {}; -} - -//===----------------------------------------------------------------------===// -// Attribute Printing -//===----------------------------------------------------------------------===// - -static void print(spirv::VerCapExtAttr triple, DialectAsmPrinter &printer) { - auto &os = printer.getStream(); - printer << spirv::VerCapExtAttr::getKindName() << "<" - << spirv::stringifyVersion(triple.getVersion()) << ", ["; - llvm::interleaveComma( - triple.getCapabilities(), os, - [&](spirv::Capability cap) { os << spirv::stringifyCapability(cap); }); - printer << "], ["; - llvm::interleaveComma(triple.getExtensionsAttr(), os, [&](Attribute attr) { - os << attr.cast().getValue(); - }); - printer << "]>"; -} - -static void print(spirv::TargetEnvAttr targetEnv, DialectAsmPrinter &printer) { - printer << spirv::TargetEnvAttr::getKindName() << "<#spv."; - print(targetEnv.getTripleAttr(), printer); - spirv::Vendor vendorID = targetEnv.getVendorID(); - spirv::DeviceType deviceType = targetEnv.getDeviceType(); - uint32_t deviceID = targetEnv.getDeviceID(); - if (vendorID != spirv::Vendor::Unknown) { - printer << ", " << spirv::stringifyVendor(vendorID); - if (deviceType != spirv::DeviceType::Unknown) { - printer << ":" << spirv::stringifyDeviceType(deviceType); - if (deviceID != spirv::TargetEnvAttr::kUnknownDeviceID) - printer << ":" << deviceID; - } - } - printer << ", " << targetEnv.getResourceLimits() << ">"; -} - -static void print(spirv::InterfaceVarABIAttr interfaceVarABIAttr, - DialectAsmPrinter &printer) { - printer << spirv::InterfaceVarABIAttr::getKindName() << "<(" - << interfaceVarABIAttr.getDescriptorSet() << ", " - << interfaceVarABIAttr.getBinding() << ")"; - auto storageClass = interfaceVarABIAttr.getStorageClass(); - if (storageClass) - printer << ", " << spirv::stringifyStorageClass(*storageClass); - printer << ">"; -} - -void SPIRVDialect::printAttribute(Attribute attr, - DialectAsmPrinter &printer) const { - if (auto targetEnv = attr.dyn_cast()) - print(targetEnv, printer); - else if (auto vceAttr = attr.dyn_cast()) - print(vceAttr, printer); - else if (auto interfaceVarABIAttr = attr.dyn_cast()) - print(interfaceVarABIAttr, printer); - else - llvm_unreachable("unhandled SPIR-V attribute kind"); -} - //===----------------------------------------------------------------------===// // Constant //===----------------------------------------------------------------------===// @@ -1216,14 +892,11 @@ StringRef symbol = attribute.getName().strref(); Attribute attr = attribute.getValue(); - // TODO: figure out a way to generate the description from the - // StructAttr definition. if (symbol == spirv::getEntryPointABIAttrName()) { - if (!attr.isa()) + if (!attr.isa()) { return op->emitError("'") - << symbol - << "' attribute must be a dictionary attribute containing one " - "32-bit integer elements attribute: 'local_size'"; + << symbol << "' attribute must be an entry point ABI attribute"; + } } else if (symbol == spirv::getTargetEnvAttrName()) { if (!attr.isa()) return op->emitError("'") << symbol << "' must be a spirv::TargetEnvAttr"; diff --git a/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp b/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp --- a/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp +++ b/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp @@ -121,14 +121,13 @@ spirv::EntryPointABIAttr spirv::getEntryPointABIAttr(ArrayRef localSize, MLIRContext *context) { if (localSize.empty()) - return spirv::EntryPointABIAttr::get(nullptr, context); + return spirv::EntryPointABIAttr::get(context, nullptr); assert(localSize.size() == 3); return spirv::EntryPointABIAttr::get( - DenseElementsAttr::get( - VectorType::get(3, IntegerType::get(context, 32)), localSize) - .cast(), - context); + context, DenseElementsAttr::get( + VectorType::get(3, IntegerType::get(context, 32)), localSize) + .cast()); } spirv::EntryPointABIAttr spirv::lookupEntryPointABI(Operation *op) { @@ -146,7 +145,7 @@ DenseIntElementsAttr spirv::lookupLocalWorkGroupSize(Operation *op) { if (auto entryPoint = spirv::lookupEntryPointABI(op)) - return entryPoint.local_size(); + return entryPoint.getLocal_size(); return {}; } @@ -155,12 +154,14 @@ spirv::getDefaultResourceLimits(MLIRContext *context) { // All the fields have default values. Here we just provide a nicer way to // construct a default resource limit attribute. - return spirv::ResourceLimitsAttr ::get( - /*max_compute_shared_memory_size=*/nullptr, - /*max_compute_workgroup_invocations=*/nullptr, - /*max_compute_workgroup_size=*/nullptr, - /*subgroup_size=*/nullptr, - /*cooperative_matrix_properties_nv=*/nullptr, context); + Builder b(context); + return spirv::ResourceLimitsAttr::get( + context, + /*max_compute_shared_memory_size=*/16384, + /*max_compute_workgroup_invocations=*/128, + /*max_compute_workgroup_size=*/b.getI32ArrayAttr({128, 128, 64}), + /*subgroup_size=*/32, + /*cooperative_matrix_properties_nv=*/ArrayAttr()); } StringRef spirv::getTargetEnvAttrName() { return "spv.target_env"; } diff --git a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp --- a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp @@ -135,7 +135,7 @@ funcOp.getLoc(), executionModel.getValue(), funcOp, interfaceVars); // Specifies the spv.ExecutionModeOp. - auto localSizeAttr = entryPointAttr.local_size(); + auto localSizeAttr = entryPointAttr.getLocal_size(); if (localSizeAttr) { auto values = localSizeAttr.getValues(); SmallVector localSize(values); diff --git a/mlir/test/Conversion/ArithmeticToSPIRV/arithmetic-to-spirv.mlir b/mlir/test/Conversion/ArithmeticToSPIRV/arithmetic-to-spirv.mlir --- a/mlir/test/Conversion/ArithmeticToSPIRV/arithmetic-to-spirv.mlir +++ b/mlir/test/Conversion/ArithmeticToSPIRV/arithmetic-to-spirv.mlir @@ -6,7 +6,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // Check integer operation conversions. @@ -154,7 +154,7 @@ // Check that types are converted to 32-bit when no special capabilities. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @int_vector23 @@ -182,7 +182,7 @@ // Check that types are converted to 32-bit when no special capabilities that // are not supported. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // expected-error @+1 {{failed to materialize conversion for block argument #0 that remained live after conversion, type was 'vector<4xi64>', with target type 'vector<4xi32>'}} @@ -202,7 +202,7 @@ //===----------------------------------------------------------------------===// module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @bitwise_scalar @@ -280,7 +280,7 @@ //===----------------------------------------------------------------------===// module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @cmpf @@ -318,7 +318,7 @@ // With Kernel capability, we can convert NaN check to spv.Ordered/spv.Unordered. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @cmpf @@ -336,7 +336,7 @@ // Without Kernel capability, we need to convert NaN check to spv.IsNan. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @cmpf @@ -364,7 +364,7 @@ //===----------------------------------------------------------------------===// module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @cmpi @@ -420,7 +420,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @constant @@ -493,7 +493,7 @@ // Check that constants are converted to 32-bit when no special capability. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @constant_16bit @@ -585,7 +585,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: index_cast1 @@ -833,7 +833,7 @@ // Checks that cast types will be adjusted when missing special capabilities for // certain non-32-bit scalar types. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @fpext1 @@ -859,7 +859,7 @@ // Checks that cast types will be adjusted when missing special capabilities for // certain non-32-bit scalar types. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @fptrunc1 @@ -892,7 +892,7 @@ // Check OpenCL lowering of arith.remsi module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @scalar_srem @@ -928,7 +928,7 @@ module attributes { spv.target_env = #spv.target_env< #spv.vce, {}> + [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>> } { // CHECK-LABEL: @select @@ -949,7 +949,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // Check integer operation conversions. @@ -1079,7 +1079,7 @@ // Check that types are converted to 32-bit when no special capabilities. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @int_vector23 @@ -1107,7 +1107,7 @@ // Check that types are converted to 32-bit when no special capabilities that // are not supported. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // expected-error@below {{failed to materialize conversion for block argument #0 that remained live after conversion}} @@ -1127,7 +1127,7 @@ //===----------------------------------------------------------------------===// module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @bitwise_scalar @@ -1205,7 +1205,7 @@ //===----------------------------------------------------------------------===// module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @cmpf @@ -1243,7 +1243,7 @@ // With Kernel capability, we can convert NaN check to spv.Ordered/spv.Unordered. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @cmpf @@ -1261,7 +1261,7 @@ // Without Kernel capability, we need to convert NaN check to spv.IsNan. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @cmpf @@ -1289,7 +1289,7 @@ //===----------------------------------------------------------------------===// module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @cmpi @@ -1345,7 +1345,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @constant @@ -1407,7 +1407,7 @@ // Check that constants are converted to 32-bit when no special capability. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @constant_16bit @@ -1490,7 +1490,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: index_cast1 @@ -1729,7 +1729,7 @@ // Checks that cast types will be adjusted when missing special capabilities for // certain non-32-bit scalar types. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @fpext1 @@ -1755,7 +1755,7 @@ // Checks that cast types will be adjusted when missing special capabilities for // certain non-32-bit scalar types. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @fptrunc1 diff --git a/mlir/test/Conversion/ControlFlowToSPIRV/cf-ops-to-spirv.mlir b/mlir/test/Conversion/ControlFlowToSPIRV/cf-ops-to-spirv.mlir --- a/mlir/test/Conversion/ControlFlowToSPIRV/cf-ops-to-spirv.mlir +++ b/mlir/test/Conversion/ControlFlowToSPIRV/cf-ops-to-spirv.mlir @@ -5,7 +5,7 @@ //===----------------------------------------------------------------------===// module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: func @simple_loop diff --git a/mlir/test/Conversion/FuncToSPIRV/func-ops-to-spirv.mlir b/mlir/test/Conversion/FuncToSPIRV/func-ops-to-spirv.mlir --- a/mlir/test/Conversion/FuncToSPIRV/func-ops-to-spirv.mlir +++ b/mlir/test/Conversion/FuncToSPIRV/func-ops-to-spirv.mlir @@ -5,7 +5,7 @@ //===----------------------------------------------------------------------===// module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @return_none_val diff --git a/mlir/test/Conversion/FuncToSPIRV/types-to-spirv.mlir b/mlir/test/Conversion/FuncToSPIRV/types-to-spirv.mlir --- a/mlir/test/Conversion/FuncToSPIRV/types-to-spirv.mlir +++ b/mlir/test/Conversion/FuncToSPIRV/types-to-spirv.mlir @@ -8,7 +8,7 @@ // Check that non-32-bit integer types are converted to 32-bit types if the // corresponding capabilities are not available. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @integer8 @@ -48,7 +48,7 @@ // Check that non-32-bit integer types are kept untouched if the corresponding // capabilities are available. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @integer8 @@ -87,7 +87,7 @@ // Check that weird bitwidths are not supported. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-NOT: spv.func @integer4 @@ -108,7 +108,7 @@ // The index type is always converted into i32. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @index_type @@ -126,7 +126,7 @@ // Check that non-32-bit float types are converted to 32-bit types if the // corresponding capabilities are not available. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @float16 @@ -148,7 +148,7 @@ // Check that non-32-bit float types are kept untouched if the corresponding // capabilities are available. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @float16 @@ -169,7 +169,7 @@ // Check that bf16 is not supported. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-NOT: spv.func @bf16_type @@ -186,7 +186,7 @@ // Check that capabilities for scalar types affects vector types too: no special // capabilities available means using turning element types to 32-bit. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @int_vector @@ -215,7 +215,7 @@ // special capabilities means keep vector types untouched. module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @int_vector @@ -246,7 +246,7 @@ // Check that > 4-element vectors are not supported. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-NOT: spv.func @large_vector @@ -263,7 +263,7 @@ // Check memory spaces. module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: func @memref_mem_space @@ -301,7 +301,7 @@ // requires special capability and extension: convert them to 32-bit if not // satisfied. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // An i1 is store in 8-bit, so 5xi1 has 40 bits, which is stored in 2xi32. @@ -399,7 +399,7 @@ module attributes { spv.target_env = #spv.target_env< #spv.vce, {}> + [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @memref_8bit_PushConstant @@ -440,7 +440,7 @@ module attributes { spv.target_env = #spv.target_env< #spv.vce, {}> + [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @memref_8bit_StorageBuffer @@ -481,7 +481,7 @@ module attributes { spv.target_env = #spv.target_env< #spv.vce, {}> + [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @memref_8bit_Uniform @@ -521,7 +521,7 @@ // and extension is available. module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @memref_16bit_Input @@ -565,7 +565,7 @@ // Check that memref offset and strides affect the array size. module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @memref_offset_strides @@ -599,7 +599,7 @@ // Dynamic shapes module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // Check that unranked shapes are not supported. @@ -677,7 +677,7 @@ // Vector types module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: func @memref_vector @@ -701,7 +701,7 @@ // Vector types, check that sizes not available in SPIR-V are not transformed. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: func @memref_vector_wrong_size @@ -721,7 +721,7 @@ // Check that tensor element types are kept untouched with proper capabilities. module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @int_tensor_types @@ -752,7 +752,7 @@ // Check that tensor element types are changed to 32-bit without capabilities. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.func @int_tensor_types @@ -783,7 +783,7 @@ // Check that dynamic shapes are not supported. module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: func @unranked_tensor diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir --- a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir @@ -12,7 +12,7 @@ // CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") gpu.module @kernels { gpu.func @builtin_workgroup_id_x() kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]] // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]] // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}} @@ -38,7 +38,7 @@ // CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") gpu.module @kernels { gpu.func @builtin_workgroup_id_y() kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]] // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]] // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}} @@ -62,7 +62,7 @@ // CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") gpu.module @kernels { gpu.func @builtin_workgroup_id_z() kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]] // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]] // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}} @@ -85,7 +85,7 @@ // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 gpu.module @kernels { gpu.func @builtin_workgroup_size_x() kernel - attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // The constant value is obtained from the spv.entry_point_abi. // Note that this ignores the workgroup size specification in gpu.launch. // We may want to define gpu.workgroup_size and convert it to the entry @@ -110,7 +110,7 @@ // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 gpu.module @kernels { gpu.func @builtin_workgroup_size_y() kernel - attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // The constant value is obtained from the spv.entry_point_abi. // CHECK: spv.Constant 4 : i32 %0 = gpu.block_dim y @@ -132,7 +132,7 @@ // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450 gpu.module @kernels { gpu.func @builtin_workgroup_size_z() kernel - attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // The constant value is obtained from the spv.entry_point_abi. // CHECK: spv.Constant 1 : i32 %0 = gpu.block_dim z @@ -155,7 +155,7 @@ // CHECK: spv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId") gpu.module @kernels { gpu.func @builtin_local_id_x() kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[LOCALINVOCATIONID]] // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]] // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}} @@ -179,7 +179,7 @@ // CHECK: spv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") gpu.module @kernels { gpu.func @builtin_num_workgroups_x() kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[NUMWORKGROUPS]] // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]] // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}} @@ -196,7 +196,7 @@ // CHECK: spv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId") gpu.module @kernels { gpu.func @builtin_subgroup_id() kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[SUBGROUPID]] // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]] %0 = gpu.subgroup_id : index @@ -212,7 +212,7 @@ // CHECK: spv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups") gpu.module @kernels { gpu.func @builtin_num_subgroups() kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[NUMSUBGROUPS]] // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]] %0 = gpu.num_subgroups : index @@ -235,7 +235,7 @@ // CHECK: spv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") gpu.module @kernels { gpu.func @builtin_workgroup_size_x() kernel - attributes {spv.entry_point_abi = {}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi<>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPSIZE]] // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]] // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}} @@ -259,7 +259,7 @@ // CHECK: spv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") gpu.module @kernels { gpu.func @builtin_workgroup_size_y() kernel - attributes {spv.entry_point_abi = {}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi<>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPSIZE]] // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]] // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}} @@ -283,7 +283,7 @@ // CHECK: spv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") gpu.module @kernels { gpu.func @builtin_workgroup_size_z() kernel - attributes {spv.entry_point_abi = {}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi<>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPSIZE]] // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]] // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}} @@ -307,7 +307,7 @@ // CHECK: spv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") gpu.module @kernels { gpu.func @builtin_global_id_x() kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[GLOBALINVOCATIONID]] // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]] // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}} @@ -331,7 +331,7 @@ // CHECK: spv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") gpu.module @kernels { gpu.func @builtin_global_id_y() kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[GLOBALINVOCATIONID]] // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]] // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}} @@ -355,7 +355,7 @@ // CHECK: spv.GlobalVariable [[GLOBALINVOCATIONID:@.*]] built_in("GlobalInvocationId") gpu.module @kernels { gpu.func @builtin_global_id_z() kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[GLOBALINVOCATIONID]] // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]] // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}} @@ -373,7 +373,7 @@ // CHECK: spv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize") gpu.module @kernels { gpu.func @builtin_subgroup_size() kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[SUBGROUPSIZE]] // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]] %0 = gpu.subgroup_size : index diff --git a/mlir/test/Conversion/GPUToSPIRV/entry-point.mlir b/mlir/test/Conversion/GPUToSPIRV/entry-point.mlir --- a/mlir/test/Conversion/GPUToSPIRV/entry-point.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/entry-point.mlir @@ -2,10 +2,10 @@ // RUN: mlir-opt -test-spirv-entry-point-abi="workgroup-size=32" %s | FileCheck %s -check-prefix=WG32 // DEFAULT: gpu.func @foo() -// DEFAULT-SAME: spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>} +// DEFAULT-SAME: spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>> // WG32: gpu.func @foo() -// WG32-SAME: spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>} +// WG32-SAME: spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>> gpu.module @kernels { gpu.func @foo() kernel { diff --git a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir --- a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir @@ -3,7 +3,7 @@ module attributes { gpu.container_module, spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { func.func @load_store(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>) { %c0 = arith.constant 0 : index @@ -36,7 +36,7 @@ // CHECK-SAME: %[[ARG5:.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 5), StorageBuffer>} // CHECK-SAME: %[[ARG6:.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 6), StorageBuffer>} gpu.func @load_store_kernel(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>, %arg3: index, %arg4: index, %arg5: index, %arg6: index) kernel - attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: %[[ADDRESSWORKGROUPID:.*]] = spv.mlir.addressof @[[$WORKGROUPIDVAR]] // CHECK: %[[WORKGROUPID:.*]] = spv.Load "Input" %[[ADDRESSWORKGROUPID]] // CHECK: %[[WORKGROUPIDX:.*]] = spv.CompositeExtract %[[WORKGROUPID]]{{\[}}0 : i32{{\]}} diff --git a/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir b/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir --- a/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir @@ -2,7 +2,7 @@ module attributes { gpu.container_module, - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { gpu.module @kernels { // CHECK-LABEL: spv.module @{{.*}} Physical64 OpenCL @@ -11,9 +11,9 @@ // CHECK-NOT: spv.interface_var_abi // CHECK-SAME: {{%.*}}: !spv.ptr)>, CrossWorkgroup> // CHECK-NOT: spv.interface_var_abi - // CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>} + // CHECK-SAME: spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>> gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, 11>) kernel - attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { gpu.return } } diff --git a/mlir/test/Conversion/GPUToSPIRV/simple.mlir b/mlir/test/Conversion/GPUToSPIRV/simple.mlir --- a/mlir/test/Conversion/GPUToSPIRV/simple.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/simple.mlir @@ -6,9 +6,9 @@ // CHECK-LABEL: spv.func @basic_module_structure // CHECK-SAME: {{%.*}}: f32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>} // CHECK-SAME: {{%.*}}: !spv.ptr [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>} - // CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>} + // CHECK-SAME: spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>> gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32>) kernel - attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: spv.Return gpu.return } @@ -35,14 +35,14 @@ // CHECK-SAME: spv.interface_var_abi = #spv.interface_var_abi<(1, 2), StorageBuffer> // CHECK-SAME: !spv.ptr [0])>, StorageBuffer> // CHECK-SAME: spv.interface_var_abi = #spv.interface_var_abi<(3, 0)> - // CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>} + // CHECK-SAME: spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>> gpu.func @basic_module_structure_preset_ABI( %arg0 : f32 {spv.interface_var_abi = #spv.interface_var_abi<(1, 2), StorageBuffer>}, %arg1 : memref<12xf32> {spv.interface_var_abi = #spv.interface_var_abi<(3, 0)>}) kernel attributes - {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { + {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: spv.Return gpu.return } @@ -82,7 +82,7 @@ {spv.interface_var_abi = #spv.interface_var_abi<(1, 2), StorageBuffer>}, %arg1 : memref<12xf32>) kernel attributes - {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { + {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { gpu.return } } @@ -99,7 +99,7 @@ %arg1 : memref<12xf32> {spv.interface_var_abi = #spv.interface_var_abi<(3, 0)>}) kernel attributes - {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { + {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { gpu.return } } @@ -111,7 +111,7 @@ gpu.module @kernels { // CHECK-LABEL: spv.func @barrier gpu.func @barrier(%arg0 : f32, %arg1 : memref<12xf32>) kernel - attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { // CHECK: spv.ControlBarrier Workgroup, Workgroup, "AcquireRelease|WorkgroupMemory" gpu.barrier gpu.return diff --git a/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir b/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir --- a/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir +++ b/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir @@ -14,7 +14,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK: spv.GlobalVariable @@ -45,7 +45,7 @@ // CHECK: spv.Return func.func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>) attributes { - spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>} + spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>> } { linalg.generic #single_workgroup_reduction_trait ins(%input : memref<16xi32>) @@ -72,7 +72,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { func.func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>) { // expected-error @+1 {{failed to legalize operation 'linalg.generic'}} @@ -101,10 +101,10 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { func.func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>) attributes { - spv.entry_point_abi = {local_size = dense<[32, 1, 1]>: vector<3xi32>} + spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>> } { // expected-error @+1 {{failed to legalize operation 'linalg.generic'}} linalg.generic #single_workgroup_reduction_trait @@ -132,10 +132,10 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { func.func @single_workgroup_reduction(%input: memref<16x8xi32>, %output: memref<16xi32>) attributes { - spv.entry_point_abi = {local_size = dense<[16, 8, 1]>: vector<3xi32>} + spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>> } { // expected-error @+1 {{failed to legalize operation 'linalg.generic'}} linalg.generic #single_workgroup_reduction_trait diff --git a/mlir/test/Conversion/MathToSPIRV/math-to-core-spirv.mlir b/mlir/test/Conversion/MathToSPIRV/math-to-core-spirv.mlir --- a/mlir/test/Conversion/MathToSPIRV/math-to-core-spirv.mlir +++ b/mlir/test/Conversion/MathToSPIRV/math-to-core-spirv.mlir @@ -19,7 +19,7 @@ // ----- -module attributes { spv.target_env = #spv.target_env<#spv.vce, {}> } { +module attributes { spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { func.func @copy_sign_vector(%value: vector<3xf16>, %sign: vector<3xf16>) -> vector<3xf16> { %0 = math.copysign %value, %sign : vector<3xf16> diff --git a/mlir/test/Conversion/MathToSPIRV/math-to-glsl-spirv.mlir b/mlir/test/Conversion/MathToSPIRV/math-to-glsl-spirv.mlir --- a/mlir/test/Conversion/MathToSPIRV/math-to-glsl-spirv.mlir +++ b/mlir/test/Conversion/MathToSPIRV/math-to-glsl-spirv.mlir @@ -1,6 +1,6 @@ // RUN: mlir-opt -split-input-file -convert-math-to-spirv -verify-diagnostics %s -o - | FileCheck %s -module attributes { spv.target_env = #spv.target_env<#spv.vce, {}> } { +module attributes { spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @float32_unary_scalar func.func @float32_unary_scalar(%arg0: f32) { diff --git a/mlir/test/Conversion/MathToSPIRV/math-to-opencl-spirv.mlir b/mlir/test/Conversion/MathToSPIRV/math-to-opencl-spirv.mlir --- a/mlir/test/Conversion/MathToSPIRV/math-to-opencl-spirv.mlir +++ b/mlir/test/Conversion/MathToSPIRV/math-to-opencl-spirv.mlir @@ -1,6 +1,6 @@ // RUN: mlir-opt -split-input-file -convert-math-to-spirv -verify-diagnostics %s -o - | FileCheck %s -module attributes { spv.target_env = #spv.target_env<#spv.vce, {}> } { +module attributes { spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @float32_unary_scalar func.func @float32_unary_scalar(%arg0: f32) { diff --git a/mlir/test/Conversion/MemRefToSPIRV/alloc.mlir b/mlir/test/Conversion/MemRefToSPIRV/alloc.mlir --- a/mlir/test/Conversion/MemRefToSPIRV/alloc.mlir +++ b/mlir/test/Conversion/MemRefToSPIRV/alloc.mlir @@ -2,7 +2,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { func.func @alloc_dealloc_workgroup_mem(%arg0 : index, %arg1 : index) { @@ -27,7 +27,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { func.func @alloc_dealloc_workgroup_mem(%arg0 : index, %arg1 : index) { @@ -56,7 +56,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { func.func @two_allocs() { @@ -76,7 +76,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { func.func @two_allocs_vector() { @@ -97,7 +97,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: func @alloc_dynamic_size @@ -113,7 +113,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: func @alloc_unsupported_memory_space @@ -130,7 +130,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: func @dealloc_dynamic_size @@ -145,7 +145,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: func @dealloc_unsupported_memory_space diff --git a/mlir/test/Conversion/MemRefToSPIRV/alloca.mlir b/mlir/test/Conversion/MemRefToSPIRV/alloca.mlir --- a/mlir/test/Conversion/MemRefToSPIRV/alloca.mlir +++ b/mlir/test/Conversion/MemRefToSPIRV/alloca.mlir @@ -1,6 +1,6 @@ // RUN: mlir-opt -split-input-file -convert-memref-to-spirv -canonicalize -verify-diagnostics %s -o - | FileCheck %s -module attributes {spv.target_env = #spv.target_env<#spv.vce, {}>} { +module attributes {spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>>} { func.func @alloc_function_variable(%arg0 : index, %arg1 : index) { %0 = memref.alloca() : memref<4x5xf32, 6> %1 = memref.load %0[%arg0, %arg1] : memref<4x5xf32, 6> @@ -19,7 +19,7 @@ // ----- -module attributes {spv.target_env = #spv.target_env<#spv.vce, {}>} { +module attributes {spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>>} { func.func @two_allocs() { %0 = memref.alloca() : memref<4x5xf32, 6> %1 = memref.alloca() : memref<2x3xi32, 6> @@ -33,7 +33,7 @@ // ----- -module attributes {spv.target_env = #spv.target_env<#spv.vce, {}>} { +module attributes {spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>>} { func.func @two_allocs_vector() { %0 = memref.alloca() : memref<4xvector<4xf32>, 6> %1 = memref.alloca() : memref<2xvector<2xi32>, 6> @@ -48,7 +48,7 @@ // ----- -module attributes {spv.target_env = #spv.target_env<#spv.vce, {}>} { +module attributes {spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>>} { // CHECK-LABEL: func @alloc_dynamic_size func.func @alloc_dynamic_size(%arg0 : index) -> f32 { // CHECK: memref.alloca @@ -60,7 +60,7 @@ // ----- -module attributes {spv.target_env = #spv.target_env<#spv.vce, {}>} { +module attributes {spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>>} { // CHECK-LABEL: func @alloc_unsupported_memory_space func.func @alloc_unsupported_memory_space(%arg0: index) -> f32 { // CHECK: memref.alloca diff --git a/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir b/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir --- a/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir +++ b/mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir @@ -11,7 +11,7 @@ StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8 ], - [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class]>, {}> + [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>> } { // CHECK-LABEL: @load_store_zero_rank_float @@ -114,7 +114,7 @@ // TODO: Test i64 types. module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @load_i1 @@ -314,7 +314,7 @@ module attributes { spv.target_env = #spv.target_env< #spv.vce, {}> + [SPV_KHR_storage_buffer_storage_class, SPV_KHR_16bit_storage]>, #spv.resource_limits<>> } { // CHECK-LABEL: @load_i8 diff --git a/mlir/test/Conversion/SCFToSPIRV/for.mlir b/mlir/test/Conversion/SCFToSPIRV/for.mlir --- a/mlir/test/Conversion/SCFToSPIRV/for.mlir +++ b/mlir/test/Conversion/SCFToSPIRV/for.mlir @@ -2,7 +2,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { func.func @loop_kernel(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) { diff --git a/mlir/test/Conversion/SCFToSPIRV/if.mlir b/mlir/test/Conversion/SCFToSPIRV/if.mlir --- a/mlir/test/Conversion/SCFToSPIRV/if.mlir +++ b/mlir/test/Conversion/SCFToSPIRV/if.mlir @@ -2,7 +2,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @kernel_simple_selection diff --git a/mlir/test/Conversion/SCFToSPIRV/while.mlir b/mlir/test/Conversion/SCFToSPIRV/while.mlir --- a/mlir/test/Conversion/SCFToSPIRV/while.mlir +++ b/mlir/test/Conversion/SCFToSPIRV/while.mlir @@ -2,7 +2,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @while_loop1 diff --git a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir b/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir --- a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir +++ b/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir @@ -1,6 +1,6 @@ // RUN: mlir-opt --lower-host-to-llvm %s | FileCheck %s -module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} { +module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits>} { // CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() : !llvm.struct<(array<6 x i32>)> // CHECK: llvm.func @__spv__foo_bar() @@ -32,7 +32,7 @@ } gpu.module @foo { - gpu.func @bar(%arg0: memref<6xi32>) kernel attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}} { + gpu.func @bar(%arg0: memref<6xi32>) kernel attributes {spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>>} { gpu.return } } diff --git a/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir b/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir --- a/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir +++ b/mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir @@ -1,6 +1,6 @@ // RUN: mlir-opt -split-input-file -convert-vector-to-spirv -verify-diagnostics %s -o - | FileCheck %s -module attributes { spv.target_env = #spv.target_env<#spv.vce, {}> } { +module attributes { spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: @bitcast // CHECK-SAME: %[[ARG0:.+]]: vector<2xf32>, %[[ARG1:.+]]: vector<2xf16> diff --git a/mlir/test/Dialect/SPIRV/IR/target-and-abi.mlir b/mlir/test/Dialect/SPIRV/IR/target-and-abi.mlir --- a/mlir/test/Dialect/SPIRV/IR/target-and-abi.mlir +++ b/mlir/test/Dialect/SPIRV/IR/target-and-abi.mlir @@ -26,23 +26,24 @@ // spv.entry_point_abi //===----------------------------------------------------------------------===// -// expected-error @+1 {{'spv.entry_point_abi' attribute must be a dictionary attribute containing one 32-bit integer elements attribute: 'local_size'}} +// expected-error @+1 {{'spv.entry_point_abi' attribute must be an entry point ABI attribute}} func.func @spv_entry_point() attributes { spv.entry_point_abi = 64 } { return } // ----- -// expected-error @+1 {{'spv.entry_point_abi' attribute must be a dictionary attribute containing one 32-bit integer elements attribute: 'local_size'}} func.func @spv_entry_point() attributes { - spv.entry_point_abi = {local_size = 64} + // expected-error @+2 {{failed to parse SPV_EntryPointABIAttr parameter 'local_size' which is to be a `DenseIntElementsAttr`}} + // expected-error @+1 {{invalid kind of attribute specified}} + spv.entry_point_abi = #spv.entry_point_abi } { return } // ----- func.func @spv_entry_point() attributes { - // CHECK: {spv.entry_point_abi = {local_size = dense<[64, 1, 1]> : vector<3xi32>}} - spv.entry_point_abi = {local_size = dense<[64, 1, 1]>: vector<3xi32>} + // CHECK: {spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>>} + spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>> } { return } // ----- @@ -104,25 +105,15 @@ // spv.target_env //===----------------------------------------------------------------------===// -func.func @target_env_wrong_limits() attributes { - spv.target_env = #spv.target_env< - #spv.vce, - // expected-error @+1 {{limits must be a dictionary attribute containing two 32-bit integer attributes 'max_compute_workgroup_invocations' and 'max_compute_workgroup_size'}} - {max_compute_workgroup_invocations = 128 : i64, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}> -} { return } - -// ----- - func.func @target_env() attributes { // CHECK: spv.target_env = #spv.target_env< // CHECK-SAME: #spv.vce, - // CHECK-SAME: {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}> + // CHECK-SAME: #spv.resource_limits> spv.target_env = #spv.target_env< #spv.vce, - { - max_compute_workgroup_invocations = 128 : i32, - max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32> - }> + #spv.resource_limits< + max_compute_workgroup_size = [128, 64, 64] + >> } { return } // ----- @@ -131,8 +122,8 @@ // CHECK: spv.target_env = #spv.target_env< // CHECK-SAME: #spv.vce, // CHECK-SAME: NVIDIA, - // CHECK-SAME: {}> - spv.target_env = #spv.target_env<#spv.vce, NVIDIA, {}> + // CHECK-SAME: #spv.resource_limits<>> + spv.target_env = #spv.target_env<#spv.vce, NVIDIA, #spv.resource_limits<>> } { return } // ----- @@ -141,8 +132,8 @@ // CHECK: spv.target_env = #spv.target_env< // CHECK-SAME: #spv.vce, // CHECK-SAME: AMD:DiscreteGPU, - // CHECK-SAME: {}> - spv.target_env = #spv.target_env<#spv.vce, AMD:DiscreteGPU, {}> + // CHECK-SAME: #spv.resource_limits<>> + spv.target_env = #spv.target_env<#spv.vce, AMD:DiscreteGPU, #spv.resource_limits<>> } { return } // ----- @@ -151,20 +142,17 @@ // CHECK: spv.target_env = #spv.target_env< // CHECK-SAME: #spv.vce, // CHECK-SAME: Qualcomm:IntegratedGPU:100925441, - // CHECK-SAME: {}> - spv.target_env = #spv.target_env<#spv.vce, Qualcomm:IntegratedGPU:0x6040001, {}> + // CHECK-SAME: #spv.resource_limits<>> + spv.target_env = #spv.target_env<#spv.vce, Qualcomm:IntegratedGPU:0x6040001, #spv.resource_limits<>> } { return } // ----- func.func @target_env_extra_fields() attributes { - // expected-error @+6 {{expected '>'}} + // expected-error @+3 {{expected '>'}} spv.target_env = #spv.target_env< #spv.vce, - { - max_compute_workgroup_invocations = 128 : i32, - max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32> - }, + #spv.resource_limits<>, more_stuff > } { return } @@ -174,37 +162,38 @@ func.func @target_env_cooperative_matrix() attributes{ // CHECK: spv.target_env = #spv.target_env< // CHECK-SAME: SPV_NV_cooperative_matrix - // CHECK-SAME: cooperative_matrix_properties_nv = [ - // CHECK-SAME: {a_type = i8, b_type = i8, c_type = i32, - // CHECK-SAME: k_size = 32 : i32, m_size = 8 : i32, n_size = 8 : i32 - // CHECK-SAME: result_type = i32, scope = 3 : i32} - // CHECK-SAME: {a_type = f16, b_type = f16, c_type = f16, - // CHECK-SAME: k_size = 16 : i32, m_size = 8 : i32, n_size = 8 : i32 - // CHECK-SAME: result_type = f16, scope = 3 : i32} + // CHECK-SAME: #spv.coop_matrix_props< + // CHECK-SAME: m_size = 8, n_size = 8, k_size = 32, + // CHECK-SAME: a_type = i8, b_type = i8, c_type = i32, + // CHECK-SAME: result_type = i32, scope = 3 : i32> + // CHECK-SAME: #spv.coop_matrix_props< + // CHECK-SAME: m_size = 8, n_size = 8, k_size = 16, + // CHECK-SAME: a_type = f16, b_type = f16, c_type = f16, + // CHECK-SAME: result_type = f16, scope = 3 : i32> spv.target_env = #spv.target_env< #spv.vce, - { - cooperative_matrix_properties_nv = [{ - m_size = 8: i32, - n_size = 8: i32, - k_size = 32: i32, + #spv.resource_limits< + cooperative_matrix_properties_nv = [#spv.coop_matrix_props< + m_size = 8, + n_size = 8, + k_size = 32, a_type = i8, b_type = i8, c_type = i32, result_type = i32, - scope = 3: i32 - }, { - m_size = 8: i32, - n_size = 8: i32, - k_size = 16: i32, + scope = 3 : i32 + >, #spv.coop_matrix_props< + m_size = 8, + n_size = 8, + k_size = 16, a_type = f16, b_type = f16, c_type = f16, result_type = f16, - scope = 3: i32 - }] - }> + scope = 3 : i32 + >] + >> } { return } // ----- diff --git a/mlir/test/Dialect/SPIRV/IR/target-env.mlir b/mlir/test/Dialect/SPIRV/IR/target-env.mlir --- a/mlir/test/Dialect/SPIRV/IR/target-env.mlir +++ b/mlir/test/Dialect/SPIRV/IR/target-env.mlir @@ -35,7 +35,7 @@ // CHECK-LABEL: @cmp_exchange_weak_suitable_version_capabilities func.func @cmp_exchange_weak_suitable_version_capabilities(%ptr: !spv.ptr, %value: i32, %comparator: i32) -> i32 attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: spv.AtomicCompareExchangeWeak "Workgroup" "AcquireRelease|AtomicCounterMemory" "Acquire" %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr, i32, i32) -> (i32) @@ -44,7 +44,7 @@ // CHECK-LABEL: @cmp_exchange_weak_unsupported_version func.func @cmp_exchange_weak_unsupported_version(%ptr: !spv.ptr, %value: i32, %comparator: i32) -> i32 attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: test.convert_to_atomic_compare_exchange_weak_op %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr, i32, i32) -> (i32) @@ -57,7 +57,7 @@ // CHECK-LABEL: @group_non_uniform_ballot_suitable_version func.func @group_non_uniform_ballot_suitable_version(%predicate: i1) -> vector<4xi32> attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: spv.GroupNonUniformBallot Workgroup %0 = "test.convert_to_group_non_uniform_ballot_op"(%predicate): (i1) -> (vector<4xi32>) @@ -66,7 +66,7 @@ // CHECK-LABEL: @group_non_uniform_ballot_unsupported_version func.func @group_non_uniform_ballot_unsupported_version(%predicate: i1) -> vector<4xi32> attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: test.convert_to_group_non_uniform_ballot_op %0 = "test.convert_to_group_non_uniform_ballot_op"(%predicate): (i1) -> (vector<4xi32>) @@ -79,7 +79,7 @@ // CHECK-LABEL: @cmp_exchange_weak_missing_capability_kernel func.func @cmp_exchange_weak_missing_capability_kernel(%ptr: !spv.ptr, %value: i32, %comparator: i32) -> i32 attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: test.convert_to_atomic_compare_exchange_weak_op %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr, i32, i32) -> (i32) @@ -88,7 +88,7 @@ // CHECK-LABEL: @cmp_exchange_weak_missing_capability_atomic_storage func.func @cmp_exchange_weak_missing_capability_atomic_storage(%ptr: !spv.ptr, %value: i32, %comparator: i32) -> i32 attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: test.convert_to_atomic_compare_exchange_weak_op %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr, i32, i32) -> (i32) @@ -97,7 +97,7 @@ // CHECK-LABEL: @subgroup_ballot_missing_capability func.func @subgroup_ballot_missing_capability(%predicate: i1) -> vector<4xi32> attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: test.convert_to_subgroup_ballot_op %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>) @@ -106,7 +106,7 @@ // CHECK-LABEL: @bit_reverse_directly_implied_capability func.func @bit_reverse_directly_implied_capability(%operand: i32) -> i32 attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: spv.BitReverse %0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32) @@ -115,7 +115,7 @@ // CHECK-LABEL: @bit_reverse_recursively_implied_capability func.func @bit_reverse_recursively_implied_capability(%operand: i32) -> i32 attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: spv.BitReverse %0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32) @@ -128,7 +128,7 @@ // CHECK-LABEL: @subgroup_ballot_suitable_extension func.func @subgroup_ballot_suitable_extension(%predicate: i1) -> vector<4xi32> attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: spv.SubgroupBallotKHR %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>) @@ -137,7 +137,7 @@ // CHECK-LABEL: @subgroup_ballot_missing_extension func.func @subgroup_ballot_missing_extension(%predicate: i1) -> vector<4xi32> attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: test.convert_to_subgroup_ballot_op %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>) @@ -146,7 +146,7 @@ // CHECK-LABEL: @module_suitable_extension1 func.func @module_suitable_extension1() attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: spv.module PhysicalStorageBuffer64 Vulkan "test.convert_to_module_op"() : () ->() @@ -155,7 +155,7 @@ // CHECK-LABEL: @module_suitable_extension2 func.func @module_suitable_extension2() attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: spv.module PhysicalStorageBuffer64 Vulkan "test.convert_to_module_op"() : () -> () @@ -164,7 +164,7 @@ // CHECK-LABEL: @module_missing_extension_mm func.func @module_missing_extension_mm() attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: test.convert_to_module_op "test.convert_to_module_op"() : () -> () @@ -173,7 +173,7 @@ // CHECK-LABEL: @module_missing_extension_am func.func @module_missing_extension_am() attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: test.convert_to_module_op "test.convert_to_module_op"() : () -> () @@ -183,7 +183,7 @@ // CHECK-LABEL: @module_implied_extension func.func @module_implied_extension() attributes { // Version 1.5 implies SPV_KHR_vulkan_memory_model and SPV_KHR_physical_storage_buffer. - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { // CHECK: spv.module PhysicalStorageBuffer64 Vulkan "test.convert_to_module_op"() : () -> () diff --git a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir --- a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir +++ b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir @@ -259,14 +259,14 @@ spv.func @kernel( %arg0: f32, %arg1: !spv.ptr)>, CrossWorkgroup>) "None" - attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>>} { spv.Return } spv.func @kernel_different_attr( %arg0: f32, %arg1: !spv.ptr)>, CrossWorkgroup>) "None" - attributes {spv.entry_point_abi = {local_size = dense<[64, 1, 1]> : vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>>} { spv.Return } } diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir @@ -1,7 +1,7 @@ // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s module attributes { - spv.target_env = #spv.target_env<#spv.vce, {}> + spv.target_env = #spv.target_env<#spv.vce, #spv.resource_limits<>> } { spv.module Physical64 OpenCL { // CHECK-LABEL: spv.module @@ -11,7 +11,7 @@ spv.func @kernel( %arg0: f32, %arg1: !spv.ptr)>, CrossWorkgroup>) "None" - attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>>} { spv.Return } } diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir @@ -2,7 +2,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.module @@ -15,7 +15,7 @@ {spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>}, %arg1: !spv.ptr)>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}) "None" - attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>>} { // CHECK: [[ARG1:%.*]] = spv.mlir.addressof [[VAR1]] // CHECK: [[ADDRESSARG0:%.*]] = spv.mlir.addressof [[VAR0]] // CHECK: [[CONST0:%.*]] = spv.Constant 0 : i32 diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir @@ -2,7 +2,7 @@ module attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { // CHECK-LABEL: spv.module @@ -38,7 +38,7 @@ {spv.interface_var_abi = #spv.interface_var_abi<(0, 5), StorageBuffer>}, %arg6: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 6), StorageBuffer>}) "None" - attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} { + attributes {spv.entry_point_abi = #spv.entry_point_abi : vector<3xi32>>} { // CHECK: [[ADDRESSARG6:%.*]] = spv.mlir.addressof [[VAR6]] // CHECK: [[CONST6:%.*]] = spv.Constant 0 : i32 // CHECK: [[ARG6PTR:%.*]] = spv.AccessChain [[ADDRESSARG6]]{{\[}}[[CONST6]] diff --git a/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir b/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir @@ -10,7 +10,7 @@ // CHECK: requires #spv.vce spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { spv.func @iadd(%val : i32) -> i32 "None" { %0 = spv.IAdd %val, %val: i32 @@ -24,7 +24,7 @@ // CHECK: requires #spv.vce spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { spv.func @group_non_uniform_ballot(%predicate : i1) -> vector<4xi32> "None" { %0 = spv.GroupNonUniformBallot Workgroup %predicate : vector<4xi32> @@ -41,7 +41,7 @@ // CHECK: requires #spv.vce spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { spv.func @iadd(%val : i32) -> i32 "None" { %0 = spv.IAdd %val, %val: i32 @@ -55,7 +55,7 @@ // CHECK: requires #spv.vce spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { spv.func @iadd(%val : i32) -> i32 "None" { %0 = spv.IAdd %val, %val: i32 @@ -76,7 +76,7 @@ // CHECK: requires #spv.vce spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" { %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32 @@ -87,7 +87,7 @@ // CHECK: requires #spv.vce spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" { %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32 @@ -101,7 +101,7 @@ // CHECK: requires #spv.vce spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { spv.func @iadd_function(%val : i8) -> i8 "None" { %0 = spv.IAdd %val, %val : i8 @@ -113,7 +113,7 @@ // CHECK: requires #spv.vce spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { spv.func @fadd_function(%val : f16) -> f16 "None" { %0 = spv.FAdd %val, %val : f16 @@ -125,7 +125,7 @@ // CHECK: requires #spv.vce spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { spv.func @iadd_v16_function(%val : vector<16xi32>) -> vector<16xi32> "None" { %0 = spv.IAdd %val, %val : vector<16xi32> @@ -144,7 +144,7 @@ spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, {}> + [SPV_KHR_shader_ballot, SPV_KHR_shader_clock, SPV_KHR_variable_pointers]>, #spv.resource_limits<>> } { spv.func @subgroup_ballot(%predicate : i1) -> vector<4xi32> "None" { %0 = spv.SubgroupBallotKHR %predicate: vector<4xi32> @@ -159,7 +159,7 @@ // CHECK: requires #spv.vce spv.module Logical Vulkan attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { spv.func @iadd(%val : i32) -> i32 "None" { %0 = spv.IAdd %val, %val: i32 @@ -174,7 +174,7 @@ // CHECK: requires #spv.vce spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { spv.func @iadd_storage_buffer(%ptr : !spv.ptr) -> i16 "None" { %0 = spv.Load "StorageBuffer" %ptr : i16 @@ -190,7 +190,7 @@ spv.module Logical GLSL450 attributes { spv.target_env = #spv.target_env< #spv.vce, - {}> + #spv.resource_limits<>> } { spv.GlobalVariable @data : !spv.ptr, Uniform> spv.GlobalVariable @img : !spv.ptr, UniformConstant> diff --git a/mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp b/mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp --- a/mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp +++ b/mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp @@ -31,6 +31,9 @@ "within the " "module, intended for testing only"; } + void getDependentDialects(DialectRegistry ®istry) const override { + registry.insert(); + } TestSpirvEntryPointABIPass() = default; TestSpirvEntryPointABIPass(const TestSpirvEntryPointABIPass &) {} void runOnOperation() override; diff --git a/mlir/test/mlir-spirv-cpu-runner/double.mlir b/mlir/test/mlir-spirv-cpu-runner/double.mlir --- a/mlir/test/mlir-spirv-cpu-runner/double.mlir +++ b/mlir/test/mlir-spirv-cpu-runner/double.mlir @@ -5,12 +5,13 @@ gpu.container_module, spv.target_env = #spv.target_env< #spv.vce, - {max_compute_workgroup_invocations = 128 : i32, - max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> + #spv.resource_limits< + max_compute_workgroup_invocations = 128, + max_compute_workgroup_size = [128, 128, 64]>> } { gpu.module @kernels { gpu.func @double(%arg0 : memref<6xi32>, %arg1 : memref<6xi32>) - kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32>}} { + kernel attributes { spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { %factor = arith.constant 2 : i32 %i0 = arith.constant 0 : index diff --git a/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir b/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir --- a/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir +++ b/mlir/test/mlir-spirv-cpu-runner/simple_add.mlir @@ -5,12 +5,13 @@ gpu.container_module, spv.target_env = #spv.target_env< #spv.vce, - {max_compute_workgroup_invocations = 128 : i32, - max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> + #spv.resource_limits< + max_compute_workgroup_invocations = 128, + max_compute_workgroup_size = [128, 128, 64]>> } { gpu.module @kernels { gpu.func @sum(%arg0 : memref<3xf32>, %arg1 : memref<3x3xf32>, %arg2 : memref<3x3x3xf32>) - kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32>}} { + kernel attributes { spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { %i0 = arith.constant 0 : index %i1 = arith.constant 1 : index %i2 = arith.constant 2 : index diff --git a/mlir/test/mlir-vulkan-runner/addf.mlir b/mlir/test/mlir-vulkan-runner/addf.mlir --- a/mlir/test/mlir-vulkan-runner/addf.mlir +++ b/mlir/test/mlir-vulkan-runner/addf.mlir @@ -4,11 +4,11 @@ module attributes { gpu.container_module, spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { gpu.module @kernels { gpu.func @kernel_add(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>) - kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32> }} { + kernel attributes { spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { %0 = gpu.block_id x %1 = memref.load %arg0[%0] : memref<8xf32> %2 = memref.load %arg1[%0] : memref<8xf32> diff --git a/mlir/test/mlir-vulkan-runner/addi.mlir b/mlir/test/mlir-vulkan-runner/addi.mlir --- a/mlir/test/mlir-vulkan-runner/addi.mlir +++ b/mlir/test/mlir-vulkan-runner/addi.mlir @@ -4,11 +4,11 @@ module attributes { gpu.container_module, spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { gpu.module @kernels { gpu.func @kernel_addi(%arg0 : memref<8xi32>, %arg1 : memref<8x8xi32>, %arg2 : memref<8x8x8xi32>) - kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32>}} { + kernel attributes { spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { %x = gpu.block_id x %y = gpu.block_id y %z = gpu.block_id z diff --git a/mlir/test/mlir-vulkan-runner/addi8.mlir b/mlir/test/mlir-vulkan-runner/addi8.mlir --- a/mlir/test/mlir-vulkan-runner/addi8.mlir +++ b/mlir/test/mlir-vulkan-runner/addi8.mlir @@ -4,11 +4,11 @@ module attributes { gpu.container_module, spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { gpu.module @kernels { gpu.func @kernel_addi(%arg0 : memref<8xi8>, %arg1 : memref<8x8xi8>, %arg2 : memref<8x8x8xi32>) - kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32>}} { + kernel attributes { spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { %x = gpu.block_id x %y = gpu.block_id y %z = gpu.block_id z diff --git a/mlir/test/mlir-vulkan-runner/mulf.mlir b/mlir/test/mlir-vulkan-runner/mulf.mlir --- a/mlir/test/mlir-vulkan-runner/mulf.mlir +++ b/mlir/test/mlir-vulkan-runner/mulf.mlir @@ -4,11 +4,11 @@ module attributes { gpu.container_module, spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits<>> } { gpu.module @kernels { gpu.func @kernel_mul(%arg0 : memref<4x4xf32>, %arg1 : memref<4x4xf32>, %arg2 : memref<4x4xf32>) - kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32> }} { + kernel attributes { spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { %x = gpu.block_id x %y = gpu.block_id y %1 = memref.load %arg0[%x, %y] : memref<4x4xf32> diff --git a/mlir/test/mlir-vulkan-runner/subf.mlir b/mlir/test/mlir-vulkan-runner/subf.mlir --- a/mlir/test/mlir-vulkan-runner/subf.mlir +++ b/mlir/test/mlir-vulkan-runner/subf.mlir @@ -4,11 +4,12 @@ module attributes { gpu.container_module, spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, + #spv.resource_limits<>> } { gpu.module @kernels { gpu.func @kernel_sub(%arg0 : memref<8x4x4xf32>, %arg1 : memref<4x4xf32>, %arg2 : memref<8x4x4xf32>) - kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32> }} { + kernel attributes { spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { %x = gpu.block_id x %y = gpu.block_id y %z = gpu.block_id z diff --git a/mlir/test/mlir-vulkan-runner/time.mlir b/mlir/test/mlir-vulkan-runner/time.mlir --- a/mlir/test/mlir-vulkan-runner/time.mlir +++ b/mlir/test/mlir-vulkan-runner/time.mlir @@ -7,11 +7,11 @@ module attributes { gpu.container_module, spv.target_env = #spv.target_env< - #spv.vce, {}> + #spv.vce, #spv.resource_limits> } { gpu.module @kernels { gpu.func @kernel_add(%arg0 : memref<16384xf32>, %arg1 : memref<16384xf32>, %arg2 : memref<16384xf32>) - kernel attributes { spv.entry_point_abi = {local_size = dense<[128, 1, 1]>: vector<3xi32> }} { + kernel attributes { spv.entry_point_abi = #spv.entry_point_abi: vector<3xi32>>} { %bid = gpu.block_id x %tid = gpu.thread_id x %cst = arith.constant 128 : index diff --git a/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp b/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp --- a/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp +++ b/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp @@ -250,9 +250,11 @@ FmtContext ctx; ctx.addSubst("_parser", "odsParser"); ctx.addSubst("_ctx", "odsParser.getContext()"); + ctx.withBuilder("odsBuilder"); if (isa(def)) ctx.addSubst("_type", "odsType"); os.indent(); + os << "::mlir::Builder odsBuilder(odsParser.getContext());\n"; // Declare variables to store all of the parameters. Allocated parameters // such as `ArrayRef` and `StringRef` must provide a `storageType`. Store @@ -668,7 +670,9 @@ FmtContext ctx; ctx.addSubst("_printer", "odsPrinter"); ctx.addSubst("_ctx", "getContext()"); + ctx.withBuilder("odsBuilder"); os.indent(); + os << "::mlir::Builder odsBuilder(getContext());\n"; // Generate printers. shouldEmitSpace = true; diff --git a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel --- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel +++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel @@ -4416,6 +4416,24 @@ deps = [":SPIRVOpsTdFiles"], ) +gentbl_cc_library( + name = "SPIRVAttributesIncGen", + strip_include_prefix = "include", + tbl_outs = [ + ( + ["-gen-attrdef-decls"], + "include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.h.inc", + ), + ( + ["-gen-attrdef-defs"], + "include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.cpp.inc", + ), + ], + tblgen = ":mlir-tblgen", + td_file = "include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.td", + deps = [":SPIRVOpsTdFiles"], +) + gentbl_cc_library( name = "SPIRVCanonicalizationIncGen", strip_include_prefix = "lib/Dialect/SPIRV/IR", @@ -4452,23 +4470,6 @@ deps = [":SPIRVOpsTdFiles"], ) -gentbl_cc_library( - name = "SPIRVTargetAndABIStructGen", - tbl_outs = [ - ( - ["-gen-struct-attr-decls"], - "include/mlir/Dialect/SPIRV/IR/TargetAndABI.h.inc", - ), - ( - ["-gen-struct-attr-defs"], - "include/mlir/Dialect/SPIRV/IR/TargetAndABI.cpp.inc", - ), - ], - tblgen = ":mlir-tblgen", - td_file = "include/mlir/Dialect/SPIRV/IR/TargetAndABI.td", - deps = [":SPIRVOpsTdFiles"], -) - gentbl_cc_library( name = "SPIRVAttrUtilsGen", strip_include_prefix = "include", @@ -4514,11 +4515,11 @@ ":InferTypeOpInterface", ":Parser", ":SPIRVAttrUtilsGen", + ":SPIRVAttributesIncGen", ":SPIRVAvailabilityIncGen", ":SPIRVCanonicalizationIncGen", ":SPIRVOpsIncGen", ":SPIRVSerializationGen", - ":SPIRVTargetAndABIStructGen", ":SideEffectInterfaces", ":Support", ":Transforms",