[mlir][spirv] Replace StructAttrs with AttrDefs

[Mogball]

Summary

Depends on D127370

[antiagainst]

Nice, thanks for migrating these!

[Mogball]

rebase

mlir/include/mlir/Dialect/SPIRV/IR/CMakeLists.txt

 add_public_tablegen_target(MLIRSPIRVSerializationGen)
 add_dependencies(mlir-headers MLIRSPIRVSerializationGen)
 
 set(LLVM_TARGET_DEFINITIONS SPIRVBase.td)
 mlir_tablegen(SPIRVAttrUtils.inc -gen-spirv-attr-utils)
 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)

mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.h

 #ifndef MLIR_DIALECT_SPIRV_IR_SPIRVATTRIBUTES_H
 #define MLIR_DIALECT_SPIRV_IR_SPIRVATTRIBUTES_H
 
 #include "mlir/Dialect/SPIRV/IR/SPIRVTypes.h"
 #include "mlir/IR/BuiltinAttributes.h"
 #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 {
 enum class Capability : uint32_t;
 enum class DeviceType : uint32_t;
 enum class Extension : uint32_t;
 enum class Vendor : uint32_t;
 enum class Version : uint32_t;
public:
   /// ID for unknown devices.
   static constexpr uint32_t kUnknownDeviceID = 0x7FFFFFFF;
 
   using Base::Base;
 
   /// 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();
 
   /// Returns the (version, capabilities, extensions) triple attribute.
   VerCapExtAttr getTripleAttr() const;
 
   /// Returns the target version.
public:
   /// Returns the device type.
   DeviceType getDeviceType() const;
 
   /// Returns the device ID.
   uint32_t getDeviceID() const;
 
   /// Returns the target resource limits.
   ResourceLimitsAttr getResourceLimits() const;
-
-  static LogicalResult verify(function_ref<InFlightDiagnostic()> emitError,
-                              VerCapExtAttr triple, Vendor vendorID,
-                              DeviceType deviceType, uint32_t deviceID,
-                              DictionaryAttr limits);
 };
 } // namespace spirv
 } // namespace mlir
 
 #endif // MLIR_DIALECT_SPIRV_IR_SPIRVATTRIBUTES_H

mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.td

 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI
 #define MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI
 
 include "mlir/Dialect/SPIRV/IR/SPIRVBase.td"
 
+class SPV_Attr<string attrName, string attrMnemonic, list<Trait> traits = []>
+    : AttrDef<SPIRV_Dialect, attrName> {
+  let mnemonic = attrMnemonic;
+}
+
 // 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<I32ElementsAttr>>
-]>;
+def SPV_EntryPointABIAttr : SPV_Attr<"EntryPointABI", "entry_point_abi"> {
+  let parameters = (ins OptionalParameter<"DenseIntElementsAttr">:$local_size);
+  let assemblyFormat = "`<` struct(params) `>`";
+}
 
 def SPV_ExtensionArrayAttr : TypedArrayAttrBase<
     SPV_ExtensionAttr, "SPIR-V extension array attribute">;
 
 def SPV_CapabilityArrayAttr : TypedArrayAttrBase<
     SPV_CapabilityAttr, "SPIR-V capability array attribute">;
 
 // Description of cooperative matrix operations supported on the
 // 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>
-]>;
+    SPV_Attr<"CooperativeMatrixPropertiesNV", "coop_matrix_props"> {
+  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 assemblyFormat = "`<` struct(params) `>`";
+}
 
 def SPV_CooperativeMatrixPropertiesNVArrayAttr :
     TypedArrayAttrBase<SPV_CooperativeMatrixPropertiesNVAttr,
                        "CooperativeMatrixPropertiesNV array attribute">;
 
 // This attribute specifies the limits for various resources on the target
 // architecture.
 //
 // See https://www.khronos.org/registry/vulkan/specs/1.2-extensions/html/vkspec.html#limits
 // for the complete list of limits and their explanation for the Vulkan API.
 // The following ones are those affecting SPIR-V CodeGen. Their default value
 // are the from Vulkan limit requirements:
 // https://www.khronos.org/registry/vulkan/specs/1.2-extensions/html/vkspec.html#limits-minmax
-def SPV_ResourceLimitsAttr : StructAttr<"ResourceLimitsAttr", SPIRV_Dialect, [
+def SPV_ResourceLimitsAttr : SPV_Attr<"ResourceLimits", "resource_limits"> {
+  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<I32Attr, "16384">>,
+    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<I32Attr, "128">>,
+    DefaultValuedParameter<"int", "128">:$max_compute_workgroup_invocations,
     // The maximum size of a local compute workgroup, per dimension.
-    StructFieldAttr<"max_compute_workgroup_size",
-                    DefaultValuedAttr<I32ElementsAttr, "{128, 128, 64}">>,
+    DefaultValuedParameter<
+      "ArrayAttr",
+      "$_builder.getI32ArrayAttr({128, 128, 64})"
+    >:$max_compute_workgroup_size,
 
     // The default number of invocations in each subgroup.
-    StructFieldAttr<"subgroup_size", DefaultValuedAttr<I32Attr, "32">>,
+    DefaultValuedParameter<"int", "32">:$subgroup_size,
 
     // The configurations of cooperative matrix operations
     // supported. Default is an empty list.
-    StructFieldAttr<
-        "cooperative_matrix_properties_nv",
-        DefaultValuedAttr<SPV_CooperativeMatrixPropertiesNVArrayAttr, "{}">>
-]>;
+    DefaultValuedParameter<
+      "ArrayAttr", 
+      "nullptr"
+    >:$cooperative_matrix_properties_nv
+  );
+  let assemblyFormat = "`<` struct(params) `>`";
+}
 
 #endif // MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI

mlir/include/mlir/Dialect/SPIRV/IR/SPIRVBase.td

let description = [{
     benefits transformations. The dialect also aims to maintain straightforward
     serialization into and deserialization from the SPIR-V binary format.
     See https://mlir.llvm.org/docs/Dialects/SPIR-V/ for more details regarding
     high-level designs and implementation structures of the SPIR-V dialect.
   }];
 
   let cppNamespace = "::mlir::spirv";
   let useDefaultTypePrinterParser = 1;
-  let useDefaultAttributePrinterParser = 1;
   let hasConstantMaterializer = 1;
   let hasOperationAttrVerify = 1;
   let hasRegionArgAttrVerify = 1;
   let hasRegionResultAttrVerify = 1;
 
   let extraClassDeclaration = [{
     void registerAttributes();
     void registerTypes();
 
     //===------------------------------------------------------------------===//
     // Attribute
     //===------------------------------------------------------------------===//
 
     /// 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;
   }];
 }
 
 //===----------------------------------------------------------------------===//
 // Utility definitions
 //===----------------------------------------------------------------------===//
 
 // A predicate that checks whether `$_self` is a known enum case for the

mlir/lib/Dialect/SPIRV/IR/CMakeLists.txt

add_mlir_dialect_library(MLIRSPIRV
   SPIRVOps.cpp
   SPIRVTypes.cpp
   TargetAndABI.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/SPIRV
 
   DEPENDS
+  MLIRSPIRVAttributeIncGen
   MLIRSPIRVAttrUtilsGen
   MLIRSPIRVAvailabilityIncGen
   MLIRSPIRVCanonicalizationIncGen
   MLIRSPIRVEnumAvailabilityIncGen
   MLIRSPIRVEnumsIncGen
   MLIRSPIRVOpsIncGen
-  MLIRSPIRVTargetAndABIIncGen
 
   LINK_LIBS PUBLIC
   MLIRControlFlowInterfaces
   MLIRIR
   MLIRParser
   MLIRSideEffectInterfaces
   MLIRSupport
   MLIRTransforms
 )

mlir/lib/Dialect/SPIRV/IR/SPIRVAttributes.cpp

 //===- SPIRVAttributes.cpp - SPIR-V attribute definitions -----------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h"
 #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
 //===----------------------------------------------------------------------===//
 
 namespace mlir {
 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
 //===----------------------------------------------------------------------===//
 
 namespace spirv {
 namespace detail {
 

 //===----------------------------------------------------------------------===//
 // TargetEnvAttr
 //===----------------------------------------------------------------------===//
 
 spirv::TargetEnvAttr spirv::TargetEnvAttr::get(spirv::VerCapExtAttr triple,
                                                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);
 }
 
 StringRef spirv::TargetEnvAttr::getKindName() { return "target_env"; }
 
 spirv::VerCapExtAttr spirv::TargetEnvAttr::getTripleAttr() const {

 uint32_t spirv::TargetEnvAttr::getDeviceID() const {
   return getImpl()->deviceID;
 }
 
 spirv::ResourceLimitsAttr spirv::TargetEnvAttr::getResourceLimits() const {
   return getImpl()->limits.cast<spirv::ResourceLimitsAttr>();
 }
 
-LogicalResult
-spirv::TargetEnvAttr::verify(function_ref<InFlightDiagnostic()> emitError,
-                             spirv::VerCapExtAttr /*triple*/,
-                             spirv::Vendor /*vendorID*/,
-                             spirv::DeviceType /*deviceType*/,
-                             uint32_t /*deviceID*/, DictionaryAttr limits) {
-  if (!limits.isa<spirv::ResourceLimitsAttr>())
-    return emitError() << "expected spirv::ResourceLimitsAttr for limits";
+//===----------------------------------------------------------------------===//
+// ODS Generated Attributes
+//===----------------------------------------------------------------------===//
+
+#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<LogicalResult(SMLoc, StringRef)> 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<uint32_t>(*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<uint32_t>(*versionSymbol));
+    } else {
+      parser.emitError(loc, "unknown version: ") << version;
+      return {};
+    }
+  }
+
+  ArrayAttr capabilitiesAttr;
+  {
+    SmallVector<Attribute, 4> capabilities;
+    SMLoc errorloc;
+    StringRef errorKeyword;
+
+    auto processCapability = [&](SMLoc loc, StringRef capability) {
+      if (auto capSymbol = spirv::symbolizeCapability(capability)) {
+        capabilities.push_back(
+            builder.getI32IntegerAttr(static_cast<uint32_t>(*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<Attribute, 1> 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<StringAttr>().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<TargetEnvAttr>())
+    print(targetEnv, printer);
+  else if (auto vceAttr = attr.dyn_cast<VerCapExtAttr>())
+    print(vceAttr, printer);
+  else if (auto interfaceVarABIAttr = attr.dyn_cast<InterfaceVarABIAttr>())
+    print(interfaceVarABIAttr, printer);
+  else
+    llvm_unreachable("unhandled SPIR-V attribute kind");
+}
 
 //===----------------------------------------------------------------------===//
 // SPIR-V Dialect
 //===----------------------------------------------------------------------===//
 
 void spirv::SPIRVDialect::registerAttributes() {
   addAttributes<InterfaceVarABIAttr, TargetEnvAttr, VerCapExtAttr>();
+  addAttributes<
+#define GET_ATTRDEF_LIST
+#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.cpp.inc"
+      >();
 }

mlir/lib/Dialect/SPIRV/IR/SPIRVDialect.cpp

void SPIRVDialect::printType(Type type, DialectAsmPrinter &os) const {
   TypeSwitch<Type>(type)
       .Case<ArrayType, CooperativeMatrixNVType, PointerType, RuntimeArrayType,
             ImageType, SampledImageType, StructType, MatrixType>(
           [&](auto type) { print(type, os); })
       .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<LogicalResult(SMLoc, StringRef)> 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<uint32_t>(*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<uint32_t>(*versionSymbol));
-    } else {
-      parser.emitError(loc, "unknown version: ") << version;
-      return {};
-    }
-  }
-
-  ArrayAttr capabilitiesAttr;
-  {
-    SmallVector<Attribute, 4> capabilities;
-    SMLoc errorloc;
-    StringRef errorKeyword;
-
-    auto processCapability = [&](SMLoc loc, StringRef capability) {
-      if (auto capSymbol = spirv::symbolizeCapability(capability)) {
-        capabilities.push_back(
-            builder.getI32IntegerAttr(static_cast<uint32_t>(*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<Attribute, 1> 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<spirv::ResourceLimitsAttr>()) {
-      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<StringAttr>().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<TargetEnvAttr>())
-    print(targetEnv, printer);
-  else if (auto vceAttr = attr.dyn_cast<VerCapExtAttr>())
-    print(vceAttr, printer);
-  else if (auto interfaceVarABIAttr = attr.dyn_cast<InterfaceVarABIAttr>())
-    print(interfaceVarABIAttr, printer);
-  else
-    llvm_unreachable("unhandled SPIR-V attribute kind");
-}
-
-//===----------------------------------------------------------------------===//
 // Constant
 //===----------------------------------------------------------------------===//
 
 Operation *SPIRVDialect::materializeConstant(OpBuilder &builder,
                                              Attribute value, Type type,
                                              Location loc) {
   if (!spirv::ConstantOp::isBuildableWith(type))
     return nullptr;
 
   return builder.create<spirv::ConstantOp>(loc, type, value);
 }
 
 //===----------------------------------------------------------------------===//
 // Shader Interface ABI
 //===----------------------------------------------------------------------===//
 
 LogicalResult SPIRVDialect::verifyOperationAttribute(Operation *op,
                                                      NamedAttribute attribute) {
   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<spirv::EntryPointABIAttr>())
+    if (!attr.isa<spirv::EntryPointABIAttr>()) {
       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<spirv::TargetEnvAttr>())
       return op->emitError("'") << symbol << "' must be a spirv::TargetEnvAttr";
   } else {
     return op->emitError("found unsupported '")
            << symbol << "' attribute on operation";
   }
 

mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp

bool spirv::needsInterfaceVarABIAttrs(spirv::TargetEnvAttr targetAttr) {
   return false;
 }
 
 StringRef spirv::getEntryPointABIAttrName() { return "spv.entry_point_abi"; }
 
 spirv::EntryPointABIAttr
 spirv::getEntryPointABIAttr(ArrayRef<int32_t> 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<int32_t>(
+      context, DenseElementsAttr::get<int32_t>(
                    VectorType::get(3, IntegerType::get(context, 32)), localSize)
-          .cast<DenseIntElementsAttr>(),
-      context);
+                   .cast<DenseIntElementsAttr>());
 }
 
 spirv::EntryPointABIAttr spirv::lookupEntryPointABI(Operation *op) {
   while (op && !isa<FunctionOpInterface>(op))
     op = op->getParentOp();
   if (!op)
     return {};
 
   if (auto attr = op->getAttrOfType<spirv::EntryPointABIAttr>(
           spirv::getEntryPointABIAttrName()))
     return attr;
 
   return {};
 }
 
 DenseIntElementsAttr spirv::lookupLocalWorkGroupSize(Operation *op) {
   if (auto entryPoint = spirv::lookupEntryPointABI(op))
-    return entryPoint.local_size();
+    return entryPoint.getLocal_size();
 
   return {};
 }
 
 spirv::ResourceLimitsAttr
 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"; }
 
 spirv::TargetEnvAttr spirv::getDefaultTargetEnv(MLIRContext *context) {
   auto triple = spirv::VerCapExtAttr::get(spirv::Version::V_1_0,
                                           {spirv::Capability::Shader},
                                           ArrayRef<Extension>(), context);

mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp

static LogicalResult lowerEntryPointABIAttr(spirv::FuncOp funcOp,
   if (failed(executionModel))
     return funcOp.emitRemark("lower entry point failure: could not select "
                              "execution model based on 'spv.target_env'");
 
   builder.create<spirv::EntryPointOp>(
       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<int32_t>();
     SmallVector<int32_t, 3> localSize(values);
     builder.create<spirv::ExecutionModeOp>(
         funcOp.getLoc(), funcOp, spirv::ExecutionMode::LocalSize, localSize);
     funcOp->removeAttr(entryPointAttrName);
   }
   return success();

mlir/test/Conversion/ArithmeticToSPIRV/arithmetic-to-spirv.mlir

 // RUN: mlir-opt -split-input-file -convert-arith-to-spirv -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // arithmetic ops
 //===----------------------------------------------------------------------===//
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64, Shader], []>, {}>
+    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64, Shader], []>, #spv.resource_limits<>>
 } {
 
 // Check integer operation conversions.
 // CHECK-LABEL: @int32_scalar
 func.func @int32_scalar(%lhs: i32, %rhs: i32) {
   // CHECK: spv.IAdd %{{.*}}, %{{.*}}: i32
   %0 = arith.addi %lhs, %rhs: i32
   // CHECK: spv.ISub %{{.*}}, %{{.*}}: i32

 }
 
 } // end module
 
 // -----
 
 // Check that types are converted to 32-bit when no special capabilities.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @int_vector23
 func.func @int_vector23(%arg0: vector<2xi8>, %arg1: vector<3xi16>) {
   // CHECK: spv.SDiv %{{.*}}, %{{.*}}: vector<2xi32>
   %0 = arith.divsi %arg0, %arg0: vector<2xi8>
   // CHECK: spv.SDiv %{{.*}}, %{{.*}}: vector<3xi32>
   %1 = arith.divsi %arg1, %arg1: vector<3xi16>

 
 } // end module
 
 // -----
 
 // 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<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #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>'}}
 func.func @int_vector4_invalid(%arg0: vector<4xi64>) {
   // expected-error @+2 {{bitwidth emulation is not implemented yet on unsigned op}}
   // expected-note @+1 {{see existing live user here}}
   %0 = arith.divui %arg0, %arg0: vector<4xi64>
   return
 }
 
 } // end module
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // std bit ops
 //===----------------------------------------------------------------------===//
 
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @bitwise_scalar
 func.func @bitwise_scalar(%arg0 : i32, %arg1 : i32) {
   // CHECK: spv.BitwiseAnd
   %0 = arith.andi %arg0, %arg1 : i32
   // CHECK: spv.BitwiseOr
   %1 = arith.ori %arg0, %arg1 : i32

 
 // -----
 
 //===----------------------------------------------------------------------===//
 // arith.cmpf
 //===----------------------------------------------------------------------===//
 
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @cmpf
 func.func @cmpf(%arg0 : f32, %arg1 : f32) {
   // CHECK: spv.FOrdEqual
   %1 = arith.cmpf oeq, %arg0, %arg1 : f32
   // CHECK: spv.FOrdGreaterThan
   %2 = arith.cmpf ogt, %arg0, %arg1 : f32

 }
 
 } // end module
 
 // -----
 
 // With Kernel capability, we can convert NaN check to spv.Ordered/spv.Unordered.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @cmpf
 func.func @cmpf(%arg0 : f32, %arg1 : f32) {
   // CHECK: spv.Ordered
   %0 = arith.cmpf ord, %arg0, %arg1 : f32
   // CHECK: spv.Unordered
   %1 = arith.cmpf uno, %arg0, %arg1 : f32
   return
 }
 
 } // end module
 
 // -----
 
 // Without Kernel capability, we need to convert NaN check to spv.IsNan.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @cmpf
 // CHECK-SAME: %[[LHS:.+]]: f32, %[[RHS:.+]]: f32
 func.func @cmpf(%arg0 : f32, %arg1 : f32) {
   // CHECK:      %[[LHS_NAN:.+]] = spv.IsNan %[[LHS]] : f32
   // CHECK-NEXT: %[[RHS_NAN:.+]] = spv.IsNan %[[RHS]] : f32
   // CHECK-NEXT: %[[OR:.+]] = spv.LogicalOr %[[LHS_NAN]], %[[RHS_NAN]] : i1

 
 // -----
 
 //===----------------------------------------------------------------------===//
 // arith.cmpi
 //===----------------------------------------------------------------------===//
 
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @cmpi
 func.func @cmpi(%arg0 : i32, %arg1 : i32) {
   // CHECK: spv.IEqual
   %0 = arith.cmpi eq, %arg0, %arg1 : i32
   // CHECK: spv.INotEqual
   %1 = arith.cmpi ne, %arg0, %arg1 : i32

 // -----
 
 //===----------------------------------------------------------------------===//
 // arith.constant
 //===----------------------------------------------------------------------===//
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
+    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @constant
 func.func @constant() {
   // CHECK: spv.Constant true
   %0 = arith.constant true
   // CHECK: spv.Constant 42 : i32
   %1 = arith.constant 42 : i32

 }
 
 } // end module
 
 // -----
 
 // Check that constants are converted to 32-bit when no special capability.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @constant_16bit
 func.func @constant_16bit() {
   // CHECK: spv.Constant 4 : i32
   %0 = arith.constant 4 : i16
   // CHECK: spv.Constant 5.000000e+00 : f32
   %1 = arith.constant 5.0 : f16

 // -----
 
 //===----------------------------------------------------------------------===//
 // std cast ops
 //===----------------------------------------------------------------------===//
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
+    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: index_cast1
 func.func @index_cast1(%arg0: i16) {
   // CHECK: spv.SConvert %{{.+}} : i16 to i32
   %0 = arith.index_cast %arg0 : i16 to index
   return
 }

 
 } // end module
 
 // -----
 
 // 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<v1.0, [Float64], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float64], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @fpext1
 // CHECK-SAME: %[[ARG:.*]]: f32
 func.func @fpext1(%arg0: f16) -> f64 {
   // CHECK-NEXT: spv.FConvert %[[ARG]] : f32 to f64
   %0 = arith.extf %arg0 : f16 to f64
   return %0: f64

 
 } // end module
 
 // -----
 
 // 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<v1.0, [Float16], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @fptrunc1
 // CHECK-SAME: %[[ARG:.*]]: f32
 func.func @fptrunc1(%arg0 : f64) -> f16 {
   // CHECK-NEXT: spv.FConvert %[[ARG]] : f32 to f16
   %0 = arith.truncf %arg0 : f64 to f16
   return %0: f16

 
 } // end module
 
 // -----
 
 // Check OpenCL lowering of arith.remsi
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Int16, Kernel], []>, {}>
+    #spv.vce<v1.0, [Int16, Kernel], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @scalar_srem
 // CHECK-SAME: (%[[LHS:.+]]: i32, %[[RHS:.+]]: i32)
 func.func @scalar_srem(%lhs: i32, %rhs: i32) {
   // CHECK: %[[LABS:.+]] = spv.OCL.s_abs %[[LHS]] : i32
   // CHECK: %[[RABS:.+]] = spv.OCL.s_abs %[[RHS]] : i32
   // CHECK:  %[[ABS:.+]] = spv.UMod %[[LABS]], %[[RABS]] : i32

 
 } // end module
 
 // -----
 
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader, Int8, Int16, Int64, Float16, Float64],
-             [SPV_KHR_storage_buffer_storage_class]>, {}>
+             [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @select
 func.func @select(%arg0 : i32, %arg1 : i32) {
   %0 = arith.cmpi sle, %arg0, %arg1 : i32
   // CHECK: spv.Select
   %1 = arith.select %0, %arg0, %arg1 : i32
   return
 }
 
 } // end module
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // arith.select
 //===----------------------------------------------------------------------===//
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64, Shader], []>, {}>
+    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64, Shader], []>, #spv.resource_limits<>>
 } {
 
 // Check integer operation conversions.
 // CHECK-LABEL: @int32_scalar
 func.func @int32_scalar(%lhs: i32, %rhs: i32) {
   // CHECK: spv.IAdd %{{.*}}, %{{.*}}: i32
   %0 = arith.addi %lhs, %rhs: i32
   // CHECK: spv.ISub %{{.*}}, %{{.*}}: i32

 }
 
 } // end module
 
 // -----
 
 // Check that types are converted to 32-bit when no special capabilities.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @int_vector23
 func.func @int_vector23(%arg0: vector<2xi8>, %arg1: vector<3xi16>) {
   // CHECK: spv.SDiv %{{.*}}, %{{.*}}: vector<2xi32>
   %0 = arith.divsi %arg0, %arg0: vector<2xi8>
   // CHECK: spv.SDiv %{{.*}}, %{{.*}}: vector<3xi32>
   %1 = arith.divsi %arg1, %arg1: vector<3xi16>

 
 } // end module
 
 // -----
 
 // 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<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // expected-error@below {{failed to materialize conversion for block argument #0 that remained live after conversion}}
 func.func @int_vector4_invalid(%arg0: vector<4xi64>) {
   // expected-error@below {{bitwidth emulation is not implemented yet on unsigned op}}
   // expected-note@below {{see existing live user here}}
   %0 = arith.divui %arg0, %arg0: vector<4xi64>
   return
 }
 
 } // end module
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // std bit ops
 //===----------------------------------------------------------------------===//
 
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @bitwise_scalar
 func.func @bitwise_scalar(%arg0 : i32, %arg1 : i32) {
   // CHECK: spv.BitwiseAnd
   %0 = arith.andi %arg0, %arg1 : i32
   // CHECK: spv.BitwiseOr
   %1 = arith.ori %arg0, %arg1 : i32

 
 // -----
 
 //===----------------------------------------------------------------------===//
 // arith.cmpf
 //===----------------------------------------------------------------------===//
 
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @cmpf
 func.func @cmpf(%arg0 : f32, %arg1 : f32) {
   // CHECK: spv.FOrdEqual
   %1 = arith.cmpf oeq, %arg0, %arg1 : f32
   // CHECK: spv.FOrdGreaterThan
   %2 = arith.cmpf ogt, %arg0, %arg1 : f32

 }
 
 } // end module
 
 // -----
 
 // With Kernel capability, we can convert NaN check to spv.Ordered/spv.Unordered.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @cmpf
 func.func @cmpf(%arg0 : f32, %arg1 : f32) {
   // CHECK: spv.Ordered
   %0 = arith.cmpf ord, %arg0, %arg1 : f32
   // CHECK: spv.Unordered
   %1 = arith.cmpf uno, %arg0, %arg1 : f32
   return
 }
 
 } // end module
 
 // -----
 
 // Without Kernel capability, we need to convert NaN check to spv.IsNan.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @cmpf
 // CHECK-SAME: %[[LHS:.+]]: f32, %[[RHS:.+]]: f32
 func.func @cmpf(%arg0 : f32, %arg1 : f32) {
   // CHECK:      %[[LHS_NAN:.+]] = spv.IsNan %[[LHS]] : f32
   // CHECK-NEXT: %[[RHS_NAN:.+]] = spv.IsNan %[[RHS]] : f32
   // CHECK-NEXT: %[[OR:.+]] = spv.LogicalOr %[[LHS_NAN]], %[[RHS_NAN]] : i1

 
 // -----
 
 //===----------------------------------------------------------------------===//
 // arith.cmpi
 //===----------------------------------------------------------------------===//
 
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @cmpi
 func.func @cmpi(%arg0 : i32, %arg1 : i32) {
   // CHECK: spv.IEqual
   %0 = arith.cmpi eq, %arg0, %arg1 : i32
   // CHECK: spv.INotEqual
   %1 = arith.cmpi ne, %arg0, %arg1 : i32

 // -----
 
 //===----------------------------------------------------------------------===//
 // arith.constant
 //===----------------------------------------------------------------------===//
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
+    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @constant
 func.func @constant() {
   // CHECK: spv.Constant true
   %0 = arith.constant true
   // CHECK: spv.Constant 42 : i32
   %1 = arith.constant 42 : i32

 }
 
 } // end module
 
 // -----
 
 // Check that constants are converted to 32-bit when no special capability.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @constant_16bit
 func.func @constant_16bit() {
   // CHECK: spv.Constant 4 : i32
   %0 = arith.constant 4 : i16
   // CHECK: spv.Constant 5.000000e+00 : f32
   %1 = arith.constant 5.0 : f16

 // -----
 
 //===----------------------------------------------------------------------===//
 // std cast ops
 //===----------------------------------------------------------------------===//
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
+    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: index_cast1
 func.func @index_cast1(%arg0: i16) {
   // CHECK: spv.SConvert %{{.+}} : i16 to i32
   %0 = arith.index_cast %arg0 : i16 to index
   return
 }

 
 } // end module
 
 // -----
 
 // 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<v1.0, [Float64], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float64], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @fpext1
 // CHECK-SAME: %[[ARG:.*]]: f32
 func.func @fpext1(%arg0: f16) -> f64 {
   // CHECK-NEXT: spv.FConvert %[[ARG]] : f32 to f64
   %0 = arith.extf %arg0 : f16 to f64
   return %0: f64

 
 } // end module
 
 // -----
 
 // 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<v1.0, [Float16], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @fptrunc1
 // CHECK-SAME: %[[ARG:.*]]: f32
 func.func @fptrunc1(%arg0 : f64) -> f16 {
   // CHECK-NEXT: spv.FConvert %[[ARG]] : f32 to f16
   %0 = arith.truncf %arg0 : f64 to f16
   return %0: f16

mlir/test/Conversion/ControlFlowToSPIRV/cf-ops-to-spirv.mlir

 // RUN: mlir-opt -split-input-file -convert-cf-to-spirv -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // cf.br, cf.cond_br
 //===----------------------------------------------------------------------===//
 
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: func @simple_loop
 func.func @simple_loop(%begin: i32, %end: i32, %step: i32) {
 // CHECK-NEXT:  spv.Branch ^bb1
   cf.br ^bb1
 
 // CHECK-NEXT: ^bb1:    // pred: ^bb0

mlir/test/Conversion/FuncToSPIRV/func-ops-to-spirv.mlir

 // RUN: mlir-opt -split-input-file -convert-func-to-spirv -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // func.return
 //===----------------------------------------------------------------------===//
 
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @return_none_val
 func.func @return_none_val() {
   // CHECK: spv.Return
   return
 }
 

mlir/test/Conversion/FuncToSPIRV/types-to-spirv.mlir

 // RUN: mlir-opt -split-input-file -convert-func-to-spirv %s -o - | FileCheck %s
 // RUN: mlir-opt -split-input-file -convert-func-to-spirv="emulate-non-32-bit-scalar-types=false" %s -o - | FileCheck %s --check-prefix=NOEMU
 
 //===----------------------------------------------------------------------===//
 // Integer types
 //===----------------------------------------------------------------------===//
 
 // 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<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @integer8
 // CHECK-SAME: i32
 // CHECK-SAME: si32
 // CHECK-SAME: ui32
 // NOEMU-LABEL: func @integer8
 // NOEMU-SAME: i8

 
 } // end module
 
 // -----
 
 // 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<v1.0, [Int8, Int16, Int64], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Int8, Int16, Int64], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @integer8
 // CHECK-SAME: i8
 // CHECK-SAME: si8
 // CHECK-SAME: ui8
 // NOEMU-LABEL: spv.func @integer8
 // NOEMU-SAME: i8

 func.func @integer64(%arg0: i64, %arg1: si64, %arg2: ui64) { return }
 
 } // end module
 
 // -----
 
 // Check that weird bitwidths are not supported.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-NOT: spv.func @integer4
 func.func @integer4(%arg0: i4) { return }
 
 // CHECK-NOT: spv.func @integer128
 func.func @integer128(%arg0: i128) { return }
 
 // CHECK-NOT: spv.func @integer42
 func.func @integer42(%arg0: i42) { return }
 
 } // end module
 // -----
 
 //===----------------------------------------------------------------------===//
 // Index type
 //===----------------------------------------------------------------------===//
 
 // The index type is always converted into i32.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @index_type
 // CHECK-SAME: %{{.*}}: i32
 func.func @index_type(%arg0: index) { return }
 
 } // end module
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // Float types
 //===----------------------------------------------------------------------===//
 
 // 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<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @float16
 // CHECK-SAME: f32
 // NOEMU-LABEL: func @float16
 // NOEMU-SAME: f16
 func.func @float16(%arg0: f16) { return }
 
 // CHECK-LABEL: spv.func @float64
 // CHECK-SAME: f32
 // NOEMU-LABEL: func @float64
 // NOEMU-SAME: f64
 func.func @float64(%arg0: f64) { return }
 
 } // end module
 
 // -----
 
 // 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<v1.0, [Float16, Float64], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16, Float64], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @float16
 // CHECK-SAME: f16
 // NOEMU-LABEL: spv.func @float16
 // NOEMU-SAME: f16
 func.func @float16(%arg0: f16) { return }
 
 // CHECK-LABEL: spv.func @float64
 // CHECK-SAME: f64
 // NOEMU-LABEL: spv.func @float64
 // NOEMU-SAME: f64
 func.func @float64(%arg0: f64) { return }
 
 } // end module
 
 // -----
 
 // Check that bf16 is not supported.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-NOT: spv.func @bf16_type
 func.func @bf16_type(%arg0: bf16) { return }
 
 } // end module
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // Vector types
 //===----------------------------------------------------------------------===//
 
 // 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<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @int_vector
 // CHECK-SAME: vector<2xi32>
 // CHECK-SAME: vector<3xsi32>
 // CHECK-SAME: vector<4xui32>
 func.func @int_vector(
   %arg0: vector<2xi8>,

 } // end module
 
 // -----
 
 // Check that capabilities for scalar types affects vector types too: having
 // special capabilities means keep vector types untouched.
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
+    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @int_vector
 // CHECK-SAME: vector<2xi8>
 // CHECK-SAME: vector<3xsi16>
 // CHECK-SAME: vector<4xui64>
 func.func @int_vector(
   %arg0: vector<2xi8>,

 func.func @one_element_vector(%arg0: vector<1xi32>) { return }
 
 } // end module
 
 // -----
 
 // Check that > 4-element vectors are not supported.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-NOT: spv.func @large_vector
 func.func @large_vector(%arg0: vector<1024xi32>) { return }
 
 } // end module
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // MemRef types
 //===----------------------------------------------------------------------===//
 
 // Check memory spaces.
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: func @memref_mem_space
 // CHECK-SAME: StorageBuffer
 // CHECK-SAME: Uniform
 // CHECK-SAME: Workgroup
 // CHECK-SAME: PushConstant
 // CHECK-SAME: Private

 } // end module
 
 // -----
 
 // Check that using non-32-bit scalar types in interface storage classes
 // requires special capability and extension: convert them to 32-bit if not
 // satisfied.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // An i1 is store in 8-bit, so 5xi1 has 40 bits, which is stored in 2xi32.
 // CHECK-LABEL: spv.func @memref_1bit_type
 // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.array<2 x i32, stride=4> [0])>, StorageBuffer>
 // NOEMU-LABEL: func @memref_1bit_type
 // NOEMU-SAME: memref<5xi1>
 func.func @memref_1bit_type(%arg0: memref<5xi1>) { return }

 // -----
 
 // Check that using non-32-bit scalar types in interface storage classes
 // requires special capability and extension: keep as-is when the capability
 // and extension is available.
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [StoragePushConstant8, StoragePushConstant16, Int64, Float64],
-             [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, {}>
+             [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @memref_8bit_PushConstant
 // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.array<16 x i8, stride=1> [0])>, PushConstant>
 // NOEMU-LABEL: spv.func @memref_8bit_PushConstant
 // NOEMU-SAME: !spv.ptr<!spv.struct<(!spv.array<16 x i8, stride=1> [0])>, PushConstant>
 func.func @memref_8bit_PushConstant(%arg0: memref<16xi8, 7>) { return }
 

 // -----
 
 // Check that using non-32-bit scalar types in interface storage classes
 // requires special capability and extension: keep as-is when the capability
 // and extension is available.
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [StorageBuffer8BitAccess, StorageBuffer16BitAccess, Int64, Float64],
-             [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, {}>
+             [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @memref_8bit_StorageBuffer
 // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.array<16 x i8, stride=1> [0])>, StorageBuffer>
 // NOEMU-LABEL: spv.func @memref_8bit_StorageBuffer
 // NOEMU-SAME: !spv.ptr<!spv.struct<(!spv.array<16 x i8, stride=1> [0])>, StorageBuffer>
 func.func @memref_8bit_StorageBuffer(%arg0: memref<16xi8, 0>) { return }
 

 // -----
 
 // Check that using non-32-bit scalar types in interface storage classes
 // requires special capability and extension: keep as-is when the capability
 // and extension is available.
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [UniformAndStorageBuffer8BitAccess, StorageUniform16, Int64, Float64],
-             [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, {}>
+             [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @memref_8bit_Uniform
 // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.array<16 x i8, stride=1> [0])>, Uniform>
 // NOEMU-LABEL: spv.func @memref_8bit_Uniform
 // NOEMU-SAME: !spv.ptr<!spv.struct<(!spv.array<16 x i8, stride=1> [0])>, Uniform>
 func.func @memref_8bit_Uniform(%arg0: memref<16xi8, 4>) { return }
 

 
 // -----
 
 // Check that using non-32-bit scalar types in interface storage classes
 // requires special capability and extension: keep as-is when the capability
 // and extension is available.
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [StorageInputOutput16, Int64, Float64], [SPV_KHR_16bit_storage]>, {}>
+    #spv.vce<v1.0, [StorageInputOutput16, Int64, Float64], [SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @memref_16bit_Input
 // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.array<16 x f16>)>, Input>
 // NOEMU-LABEL: spv.func @memref_16bit_Input
 // NOEMU-SAME: !spv.ptr<!spv.struct<(!spv.array<16 x f16>)>, Input>
 func.func @memref_16bit_Input(%arg3: memref<16xf16, 9>) { return }
 

 
 } // end module
 
 // -----
 
 // Check that memref offset and strides affect the array size.
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [StorageBuffer16BitAccess], [SPV_KHR_16bit_storage]>, {}>
+    #spv.vce<v1.0, [StorageBuffer16BitAccess], [SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @memref_offset_strides
 func.func @memref_offset_strides(
 // CHECK-SAME: !spv.array<64 x f32, stride=4> [0])>, StorageBuffer>
 // CHECK-SAME: !spv.array<72 x f32, stride=4> [0])>, StorageBuffer>
 // CHECK-SAME: !spv.array<256 x f32, stride=4> [0])>, StorageBuffer>
 // CHECK-SAME: !spv.array<64 x f32, stride=4> [0])>, StorageBuffer>

 ) { return }
 
 } // end module
 
 // -----
 
 // Dynamic shapes
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // Check that unranked shapes are not supported.
 // CHECK-LABEL: func @unranked_memref
 // CHECK-SAME: memref<*xi32>
 func.func @unranked_memref(%arg0: memref<*xi32>) { return }
 
 // CHECK-LABEL: func @memref_1bit_type

 func.func @memref_16bit_Output(%arg4: memref<?xf16, 10>) { return }
 
 } // end module
 
 // -----
 
 // Vector types
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: func @memref_vector
 // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.array<4 x vector<2xf32>, stride=8> [0])>, StorageBuffer>
 // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.array<4 x vector<4xf32>, stride=16> [0])>, Uniform>
 func.func @memref_vector(
     %arg0: memref<4xvector<2xf32>, 0>,
     %arg1: memref<4xvector<4xf32>, 4>)
 { return }
 
 // CHECK-LABEL: func @dynamic_dim_memref_vector
 // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.rtarray<vector<4xi32>, stride=16> [0])>, StorageBuffer>
 // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.rtarray<vector<2xf32>, stride=8> [0])>, StorageBuffer>
 func.func @dynamic_dim_memref_vector(%arg0: memref<8x?xvector<4xi32>>,
                          %arg1: memref<?x?xvector<2xf32>>)
 { return }
 
 } // end module
 
 // -----
 
 // Vector types, check that sizes not available in SPIR-V are not transformed.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: func @memref_vector_wrong_size
 // CHECK-SAME: memref<4xvector<5xf32>>
 func.func @memref_vector_wrong_size(
     %arg0: memref<4xvector<5xf32>, 0>)
 { return }
 
 } // end module
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // Tensor types
 //===----------------------------------------------------------------------===//
 
 // Check that tensor element types are kept untouched with proper capabilities.
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, {}>
+    #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @int_tensor_types
 // CHECK-SAME: !spv.array<32 x i64>
 // CHECK-SAME: !spv.array<32 x i32>
 // CHECK-SAME: !spv.array<32 x i16>
 // CHECK-SAME: !spv.array<32 x i8>
 func.func @int_tensor_types(

 ) { return }
 
 } // end module
 
 // -----
 
 // Check that tensor element types are changed to 32-bit without capabilities.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.func @int_tensor_types
 // CHECK-SAME: !spv.array<32 x i32>
 // CHECK-SAME: !spv.array<32 x i32>
 // CHECK-SAME: !spv.array<32 x i32>
 // CHECK-SAME: !spv.array<32 x i32>
 func.func @int_tensor_types(

 ) { return }
 
 } // end module
 
 // -----
 
 // Check that dynamic shapes are not supported.
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: func @unranked_tensor
 // CHECK-SAME: tensor<*xi32>
 func.func @unranked_tensor(%arg0: tensor<*xi32>) { return }
 
 // CHECK-LABEL: func @dynamic_dim_tensor
 // CHECK-SAME: tensor<8x?xi32>
 func.func @dynamic_dim_tensor(%arg0: tensor<8x?xi32>) { return }
 
 } // end module

mlir/test/Conversion/GPUToSPIRV/builtins.mlir

 // RUN: mlir-opt -split-input-file -convert-gpu-to-spirv %s -o - | FileCheck %s
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.block_id x
       gpu.return
     }
   }
 }
gpu.launch_func @kernels::@builtin_workgroup_id_y
         dynamic_shared_memory_size %c256
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
       %0 = gpu.block_id y
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_z
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
       %0 = gpu.block_id z
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // 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<local_size = dense<[32, 1, 1]>: 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
       // point ABI we want here.
       // CHECK: spv.Constant 32 : i32
       %0 = gpu.block_dim x
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_y
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // 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<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
       // The constant value is obtained from the spv.entry_point_abi.
       // CHECK: spv.Constant 4 : i32
       %0 = gpu.block_dim y
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_z
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // 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<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
       // The constant value is obtained from the spv.entry_point_abi.
       // CHECK: spv.Constant 1 : i32
       %0 = gpu.block_dim z
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_local_id_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[LOCALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.thread_id x
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_num_workgroups_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[NUMWORKGROUPS]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.grid_dim x
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[SUBGROUPID]]
       // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
       %0 = gpu.subgroup_id : index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[NUMSUBGROUPS]]
       // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
       %0 = gpu.num_subgroups : index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}}
   // 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{{\]}}
       %0 = gpu.block_dim x
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_y
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}}
   // 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{{\]}}
       %0 = gpu.block_dim y
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_size_z
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}}
   // 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{{\]}}
       %0 = gpu.block_dim z
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_global_id_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[GLOBALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = gpu.global_id x
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_global_id_y
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[GLOBALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
       %0 = gpu.global_id y
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func.func @builtin() {
     %c0 = arith.constant 1 : index
     gpu.launch_func @kernels::@builtin_global_id_z
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[GLOBALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
       %0 = gpu.global_id z
       gpu.return
     }
   }
 }
 
 
 // -----
 
 module attributes {gpu.container_module} {
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[SUBGROUPSIZE]]
       // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
       %0 = gpu.subgroup_size : index
       gpu.return
     }
   }
 }

mlir/test/Conversion/GPUToSPIRV/entry-point.mlir

 // RUN: mlir-opt -test-spirv-entry-point-abi %s | FileCheck %s -check-prefix=DEFAULT
 // 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<local_size = dense<1> : 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<local_size = dense<[32, 1, 1]> : vector<3xi32>>
 
 gpu.module @kernels {
   gpu.func @foo() kernel {
     gpu.return
   }
 }

mlir/test/Conversion/GPUToSPIRV/load-store.mlir

 // RUN: mlir-opt -convert-gpu-to-spirv %s -o - | FileCheck %s
 
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
 } {
   func.func @load_store(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>) {
     %c0 = arith.constant 0 : index
     %c12 = arith.constant 12 : index
     %0 = arith.subi %c12, %c0 : index
     %c1 = arith.constant 1 : index
     %c0_0 = arith.constant 0 : index
     %c4 = arith.constant 4 : index
gpu.module @kernels {
     // CHECK-SAME: %[[ARG0:.*]]: !spv.ptr<!spv.struct<(!spv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 0)>}
     // CHECK-SAME: %[[ARG1:.*]]: !spv.ptr<!spv.struct<(!spv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
     // CHECK-SAME: %[[ARG2:.*]]: !spv.ptr<!spv.struct<(!spv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 2)>}
     // CHECK-SAME: %[[ARG3:.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 3), StorageBuffer>}
     // CHECK-SAME: %[[ARG4:.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 4), StorageBuffer>}
     // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>} {
       // CHECK: %[[ADDRESSWORKGROUPID:.*]] = spv.mlir.addressof @[[$WORKGROUPIDVAR]]
       // CHECK: %[[WORKGROUPID:.*]] = spv.Load "Input" %[[ADDRESSWORKGROUPID]]
       // CHECK: %[[WORKGROUPIDX:.*]] = spv.CompositeExtract %[[WORKGROUPID]]{{\[}}0 : i32{{\]}}
       // CHECK: %[[ADDRESSLOCALINVOCATIONID:.*]] = spv.mlir.addressof @[[$LOCALINVOCATIONIDVAR]]
       // CHECK: %[[LOCALINVOCATIONID:.*]] = spv.Load "Input" %[[ADDRESSLOCALINVOCATIONID]]
       // CHECK: %[[LOCALINVOCATIONIDX:.*]] = spv.CompositeExtract %[[LOCALINVOCATIONID]]{{\[}}0 : i32{{\]}}
       %0 = gpu.block_id x
       %1 = gpu.block_id y

mlir/test/Conversion/GPUToSPIRV/module-structure-opencl.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   gpu.container_module,
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel, Addresses], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel, Addresses], []>, #spv.resource_limits<>>
 } {
   gpu.module @kernels {
     // CHECK-LABEL: spv.module @{{.*}} Physical64 OpenCL
     //       CHECK:   spv.func
     //  CHECK-SAME:     {{%.*}}: f32
     //   CHECK-NOT:     spv.interface_var_abi
     //  CHECK-SAME:     {{%.*}}: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, 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<local_size = dense<[32, 4, 1]> : 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<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
       gpu.return
     }
   }
 
   func.func @main() {
     %0 = "op"() : () -> (f32)
     %1 = "op"() : () -> (memref<12xf32, 11>)
     %cst = arith.constant 1 : index
     gpu.launch_func @kernels::@basic_module_structure
         blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
         args(%0 : f32, %1 : memref<12xf32, 11>)
     return
   }
 }

mlir/test/Conversion/GPUToSPIRV/simple.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -split-input-file -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {gpu.container_module} {
   gpu.module @kernels {
     // CHECK:       spv.module @{{.*}} Logical GLSL450 {
     // 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<!spv.struct<(!spv.array<12 x f32, stride=4> [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<local_size = dense<[32, 4, 1]> : 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<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
       // CHECK: spv.Return
       gpu.return
     }
   }
 
   func.func @main() {
     %0 = "op"() : () -> (f32)
     %1 = "op"() : () -> (memref<12xf32>)

 module attributes {gpu.container_module} {
   gpu.module @kernels {
     // CHECK:       spv.module @{{.*}} Logical GLSL450 {
     // CHECK-LABEL: spv.func @basic_module_structure_preset_ABI
     // CHECK-SAME: {{%[a-zA-Z0-9_]*}}: f32
     // CHECK-SAME: spv.interface_var_abi = #spv.interface_var_abi<(1, 2), StorageBuffer>
     // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [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<local_size = dense<[32, 4, 1]> : 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<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
       // CHECK: spv.Return
       gpu.return
     }
   }
 }
 
 // -----
 
module attributes {gpu.container_module} {
   gpu.module @kernels {
     // expected-error @below {{failed to legalize operation 'gpu.func'}}
     // expected-remark @below {{match failure: missing 'spv.interface_var_abi' attribute at argument 1}}
     gpu.func @missing_entry_point_abi(
       %arg0 : f32
         {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<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   gpu.module @kernels {
     // expected-error @below {{failed to legalize operation 'gpu.func'}}
     // expected-remark @below {{match failure: missing 'spv.interface_var_abi' attribute at argument 0}}
     gpu.func @missing_entry_point_abi(
       %arg0 : f32,
       %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<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   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<local_size = dense<[32, 4, 1]>: vector<3xi32>>} {
       // CHECK: spv.ControlBarrier Workgroup, Workgroup, "AcquireRelease|WorkgroupMemory"
       gpu.barrier
       gpu.return
     }
   }
 
   func.func @main() {
     %0 = "op"() : () -> (f32)
     %1 = "op"() : () -> (memref<12xf32>)
     %cst = arith.constant 1 : index
     gpu.launch_func @kernels::@barrier
         blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
         args(%0 : f32, %1 : memref<12xf32>)
     return
   }
 }

mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir

 // RUN: mlir-opt -split-input-file -convert-linalg-to-spirv -canonicalize -verify-diagnostics %s -o - | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // Single workgroup reduction
 //===----------------------------------------------------------------------===//
 
 #single_workgroup_reduction_trait = {
   iterator_types = ["reduction"],
   indexing_maps = [
     affine_map<(i) -> (i)>,
     affine_map<(i) -> (0)>
   ]
 }
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
+    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #spv.resource_limits<>>
 } {
 
 // CHECK:      spv.GlobalVariable
 // CHECK-SAME: built_in("LocalInvocationId")
 
 // CHECK:      @single_workgroup_reduction
 // CHECK-SAME: (%[[INPUT:.+]]: !spv.ptr{{.+}}, %[[OUTPUT:.+]]: !spv.ptr{{.+}})
 

 // CHECK:          spv.AtomicIAdd "Device" "AcquireRelease" %[[OUTPTR]], %[[ADD]]
 // CHECK:          spv.Branch ^bb2
 // CHECK:        ^bb2:
 // CHECK:          spv.mlir.merge
 // CHECK:        }
 // 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<local_size = dense<[16, 1, 1]>: vector<3xi32>>
 } {
   linalg.generic #single_workgroup_reduction_trait
       ins(%input : memref<16xi32>)
      outs(%output : memref<1xi32>) {
     ^bb(%in: i32, %out: i32):
       %sum = arith.addi %in, %out : i32
       linalg.yield %sum : i32
   }
#single_workgroup_reduction_trait = {
   indexing_maps = [
     affine_map<(i) -> (i)>,
     affine_map<(i) -> (0)>
   ]
 }
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
+    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #spv.resource_limits<>>
 } {
 func.func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>) {
   // expected-error @+1 {{failed to legalize operation 'linalg.generic'}}
   linalg.generic #single_workgroup_reduction_trait
       ins(%input : memref<16xi32>)
      outs(%output : memref<1xi32>) {
     ^bb(%in: i32, %out: i32):
       %sum = arith.addi %in, %out : i32
#single_workgroup_reduction_trait = {
   indexing_maps = [
     affine_map<(i) -> (i)>,
     affine_map<(i) -> (0)>
   ]
 }
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
+    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #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<local_size = dense<[32, 1, 1]>: vector<3xi32>>
 } {
   // expected-error @+1 {{failed to legalize operation 'linalg.generic'}}
   linalg.generic #single_workgroup_reduction_trait
       ins(%input : memref<16xi32>)
      outs(%output : memref<1xi32>) {
     ^bb(%in: i32, %out: i32):
       %sum = arith.addi %in, %out : i32
       linalg.yield %sum : i32
#single_workgroup_reduction_trait = {
   indexing_maps = [
     affine_map<(i, j) -> (i, j)>,
     affine_map<(i, j) -> (i)>
   ]
 }
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
+    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #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<local_size = dense<[16, 8, 1]>: vector<3xi32>>
 } {
   // expected-error @+1 {{failed to legalize operation 'linalg.generic'}}
   linalg.generic #single_workgroup_reduction_trait
       ins(%input : memref<16x8xi32>)
      outs(%output : memref<16xi32>) {
     ^bb(%in: i32, %out: i32):
       %sum = arith.addi %in, %out : i32
       linalg.yield %sum : i32
   }
   spv.Return
 }
 }

mlir/test/Conversion/MathToSPIRV/math-to-core-spirv.mlir

 //       CHECK:   %[[VAND:.+]] = spv.BitwiseAnd %[[VCAST]], %[[VMASK]] : i32
 //       CHECK:   %[[SAND:.+]] = spv.BitwiseAnd %[[SCAST]], %[[SMASK]] : i32
 //       CHECK:   %[[OR:.+]] = spv.BitwiseOr %[[VAND]], %[[SAND]] : i32
 //       CHECK:   %[[RESULT:.+]] = spv.Bitcast %[[OR]] : i32 to f32
 //       CHECK:   return %[[RESULT]]
 
 // -----
 
-module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16, Int16], []>, {}> } {
+module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16, Int16], []>, #spv.resource_limits<>> } {
 
 func.func @copy_sign_vector(%value: vector<3xf16>, %sign: vector<3xf16>) -> vector<3xf16> {
   %0 = math.copysign %value, %sign : vector<3xf16>
   return %0: vector<3xf16>
 }
 
 }
 

mlir/test/Conversion/MathToSPIRV/math-to-glsl-spirv.mlir

 // 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<v1.0, [Shader], []>, {}> } {
+module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], []>, #spv.resource_limits<>> } {
 
 // CHECK-LABEL: @float32_unary_scalar
 func.func @float32_unary_scalar(%arg0: f32) {
   // CHECK: spv.GLSL.Cos %{{.*}}: f32
   %0 = math.cos %arg0 : f32
   // CHECK: spv.GLSL.Exp %{{.*}}: f32
   %1 = math.exp %arg0 : f32
   // CHECK: %[[EXP:.+]] = spv.GLSL.Exp %arg0

mlir/test/Conversion/MathToSPIRV/math-to-opencl-spirv.mlir

 // 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<v1.0, [Kernel], []>, {}> } {
+module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel], []>, #spv.resource_limits<>> } {
 
 // CHECK-LABEL: @float32_unary_scalar
 func.func @float32_unary_scalar(%arg0: f32) {
   // CHECK: spv.OCL.cos %{{.*}}: f32
   %0 = math.cos %arg0 : f32
   // CHECK: spv.OCL.exp %{{.*}}: f32
   %1 = math.exp %arg0 : f32
   // CHECK: %[[EXP:.+]] = spv.OCL.exp %arg0

mlir/test/Conversion/MemRefToSPIRV/alloc.mlir

 // 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<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
   }
 {
   func.func @alloc_dealloc_workgroup_mem(%arg0 : index, %arg1 : index) {
     %0 = memref.alloc() : memref<4x5xf32, 3>
     %1 = memref.load %0[%arg0, %arg1] : memref<4x5xf32, 3>
     memref.store %1, %0[%arg0, %arg1] : memref<4x5xf32, 3>
     memref.dealloc %0 : memref<4x5xf32, 3>
     return
   }
 }
 //     CHECK: spv.GlobalVariable @[[VAR:.+]] : !spv.ptr<!spv.struct<(!spv.array<20 x f32>)>, Workgroup>
 //     CHECK: func @alloc_dealloc_workgroup_mem
 // CHECK-NOT:   memref.alloc
 //     CHECK:   %[[PTR:.+]] = spv.mlir.addressof @[[VAR]]
 //     CHECK:   %[[LOADPTR:.+]] = spv.AccessChain %[[PTR]]
 //     CHECK:   %[[VAL:.+]] = spv.Load "Workgroup" %[[LOADPTR]] : f32
 //     CHECK:   %[[STOREPTR:.+]] = spv.AccessChain %[[PTR]]
 //     CHECK:   spv.Store "Workgroup" %[[STOREPTR]], %[[VAL]] : f32
 // CHECK-NOT:   memref.dealloc
 
 // -----
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
   }
 {
   func.func @alloc_dealloc_workgroup_mem(%arg0 : index, %arg1 : index) {
     %0 = memref.alloc() : memref<4x5xi16, 3>
     %1 = memref.load %0[%arg0, %arg1] : memref<4x5xi16, 3>
     memref.store %1, %0[%arg0, %arg1] : memref<4x5xi16, 3>
     memref.dealloc %0 : memref<4x5xi16, 3>
     return

 //       CHECK:   %{{.+}} = spv.AtomicAnd "Workgroup" "AcquireRelease" %[[PTR]], %{{.+}} : !spv.ptr<i32, Workgroup>
 //       CHECK:   %{{.+}} = spv.AtomicOr "Workgroup" "AcquireRelease" %[[PTR]], %{{.+}} : !spv.ptr<i32, Workgroup>
 
 
 // -----
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
   }
 {
   func.func @two_allocs() {
     %0 = memref.alloc() : memref<4x5xf32, 3>
     %1 = memref.alloc() : memref<2x3xi32, 3>
     return
   }
 }
 
 //  CHECK-DAG: spv.GlobalVariable @__workgroup_mem__{{[0-9]+}}
 // CHECK-SAME:   !spv.ptr<!spv.struct<(!spv.array<6 x i32>)>, Workgroup>
 //  CHECK-DAG: spv.GlobalVariable @__workgroup_mem__{{[0-9]+}}
 // CHECK-SAME:   !spv.ptr<!spv.struct<(!spv.array<20 x f32>)>, Workgroup>
 //      CHECK: func @two_allocs()
 
 // -----
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
   }
 {
   func.func @two_allocs_vector() {
     %0 = memref.alloc() : memref<4xvector<4xf32>, 3>
     %1 = memref.alloc() : memref<2xvector<2xi32>, 3>
     return
   }
 }
 
 //  CHECK-DAG: spv.GlobalVariable @__workgroup_mem__{{[0-9]+}}
 // CHECK-SAME:   !spv.ptr<!spv.struct<(!spv.array<2 x vector<2xi32>>)>, Workgroup>
 //  CHECK-DAG: spv.GlobalVariable @__workgroup_mem__{{[0-9]+}}
 // CHECK-SAME:   !spv.ptr<!spv.struct<(!spv.array<4 x vector<4xf32>>)>, Workgroup>
 //      CHECK: func @two_allocs_vector()
 
 
 // -----
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
   }
 {
   // CHECK-LABEL: func @alloc_dynamic_size
   func.func @alloc_dynamic_size(%arg0 : index) -> f32 {
     // CHECK: memref.alloc
     %0 = memref.alloc(%arg0) : memref<4x?xf32, 3>
     %1 = memref.load %0[%arg0, %arg0] : memref<4x?xf32, 3>
     return %1: f32
   }
 }
 
 // -----
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
   }
 {
   // CHECK-LABEL: func @alloc_unsupported_memory_space
   func.func @alloc_unsupported_memory_space(%arg0: index) -> f32 {
     // CHECK: memref.alloc
     %0 = memref.alloc() : memref<4x5xf32>
     %1 = memref.load %0[%arg0, %arg0] : memref<4x5xf32>
     return %1: f32
   }
 }
 
 
 // -----
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
   }
 {
   // CHECK-LABEL: func @dealloc_dynamic_size
   func.func @dealloc_dynamic_size(%arg0 : memref<4x?xf32, 3>) {
     // CHECK: memref.dealloc
     memref.dealloc %arg0 : memref<4x?xf32, 3>
     return
   }
 }
 
 // -----
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
   }
 {
   // CHECK-LABEL: func @dealloc_unsupported_memory_space
   func.func @dealloc_unsupported_memory_space(%arg0 : memref<4x5xf32>) {
     // CHECK: memref.dealloc
     memref.dealloc %arg0 : memref<4x5xf32>
     return
   }
 }

mlir/test/Conversion/MemRefToSPIRV/alloca.mlir

 // 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<v1.3, [Shader], []>, {}>} {
+module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, #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>
     memref.store %1, %0[%arg0, %arg1] : memref<4x5xf32, 6>
     return
   }
 }
 
 // CHECK-LABEL: func @alloc_function_variable
 //       CHECK:   %[[VAR:.+]] = spv.Variable : !spv.ptr<!spv.struct<(!spv.array<20 x f32>)>, Function>
 //       CHECK:   %[[LOADPTR:.+]] = spv.AccessChain %[[VAR]]
 //       CHECK:   %[[VAL:.+]] = spv.Load "Function" %[[LOADPTR]] : f32
 //       CHECK:   %[[STOREPTR:.+]] = spv.AccessChain %[[VAR]]
 //       CHECK:   spv.Store "Function" %[[STOREPTR]], %[[VAL]] : f32
 
 
 // -----
 
-module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
+module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, #spv.resource_limits<>>} {
   func.func @two_allocs() {
     %0 = memref.alloca() : memref<4x5xf32, 6>
     %1 = memref.alloca() : memref<2x3xi32, 6>
     return
   }
 }
 
 // CHECK-LABEL: func @two_allocs
 //   CHECK-DAG: spv.Variable : !spv.ptr<!spv.struct<(!spv.array<6 x i32>)>, Function>
 //   CHECK-DAG: spv.Variable : !spv.ptr<!spv.struct<(!spv.array<20 x f32>)>, Function>
 
 // -----
 
-module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
+module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, #spv.resource_limits<>>} {
   func.func @two_allocs_vector() {
     %0 = memref.alloca() : memref<4xvector<4xf32>, 6>
     %1 = memref.alloca() : memref<2xvector<2xi32>, 6>
     return
   }
 }
 
 // CHECK-LABEL: func @two_allocs_vector
 //   CHECK-DAG: spv.Variable : !spv.ptr<!spv.struct<(!spv.array<2 x vector<2xi32>>)>, Function>
 //   CHECK-DAG: spv.Variable : !spv.ptr<!spv.struct<(!spv.array<4 x vector<4xf32>>)>, Function>
 
 
 // -----
 
-module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
+module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, #spv.resource_limits<>>} {
   // CHECK-LABEL: func @alloc_dynamic_size
   func.func @alloc_dynamic_size(%arg0 : index) -> f32 {
     // CHECK: memref.alloca
     %0 = memref.alloca(%arg0) : memref<4x?xf32, 6>
     %1 = memref.load %0[%arg0, %arg0] : memref<4x?xf32, 6>
     return %1: f32
   }
 }
 
 // -----
 
-module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
+module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, #spv.resource_limits<>>} {
   // CHECK-LABEL: func @alloc_unsupported_memory_space
   func.func @alloc_unsupported_memory_space(%arg0: index) -> f32 {
     // CHECK: memref.alloca
     %0 = memref.alloca() : memref<4x5xf32>
     %1 = memref.load %0[%arg0, %arg0] : memref<4x5xf32>
     return %1: f32
   }
 }

mlir/test/Conversion/MemRefToSPIRV/memref-to-spirv.mlir

 // RUN: mlir-opt -split-input-file -convert-memref-to-spirv="bool-num-bits=8" %s -o - | FileCheck %s
 
 // Check that with proper compute and storage extensions, we don't need to
 // perform special tricks.
 
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0,
       [
         Shader, Int8, Int16, Int64, Float16, Float64,
         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
 func.func @load_store_zero_rank_float(%arg0: memref<f32>, %arg1: memref<f32>) {
   //      CHECK: [[ARG0:%.*]] = builtin.unrealized_conversion_cast {{.+}} : memref<f32> to !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>
   //      CHECK: [[ARG1:%.*]] = builtin.unrealized_conversion_cast {{.+}} : memref<f32> to !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>
   //      CHECK: [[ZERO1:%.*]] = spv.Constant 0 : i32
   //      CHECK: spv.AccessChain [[ARG0]][

 
 // -----
 
 // Check that access chain indices are properly adjusted if non-32-bit types are
 // emulated via 32-bit types.
 // TODO: Test i64 types.
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @load_i1
 func.func @load_i1(%arg0: memref<i1>) -> i1 {
   //     CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
   //     CHECK: %[[FOUR1:.+]] = spv.Constant 4 : i32
   //     CHECK: %[[QUOTIENT:.+]] = spv.SDiv %[[ZERO]], %[[FOUR1]] : i32
   //     CHECK: %[[PTR:.+]] = spv.AccessChain %{{.+}}[%[[ZERO]], %[[QUOTIENT]]]

 
 // -----
 
 // Check that access chain indices are properly adjusted if non-16/32-bit types
 // are emulated via 32-bit types.
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Int16, StorageBuffer16BitAccess, Shader],
-    [SPV_KHR_storage_buffer_storage_class, SPV_KHR_16bit_storage]>, {}>
+    [SPV_KHR_storage_buffer_storage_class, SPV_KHR_16bit_storage]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @load_i8
 func.func @load_i8(%arg0: memref<i8>) {
   //     CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
   //     CHECK: %[[FOUR1:.+]] = spv.Constant 4 : i32
   //     CHECK: %[[QUOTIENT:.+]] = spv.SDiv %[[ZERO]], %[[FOUR1]] : i32
   //     CHECK: %[[PTR:.+]] = spv.AccessChain %{{.+}}[%[[ZERO]], %[[QUOTIENT]]]

mlir/test/Conversion/SCFToSPIRV/for.mlir

 // RUN: mlir-opt -convert-scf-to-spirv %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
 } {
 
 func.func @loop_kernel(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) {
   // CHECK: %[[LB:.*]] = spv.Constant 4 : i32
   %lb = arith.constant 4 : index
   // CHECK: %[[UB:.*]] = spv.Constant 42 : i32
   %ub = arith.constant 42 : index
   // CHECK: %[[STEP:.*]] = spv.Constant 2 : i32

mlir/test/Conversion/SCFToSPIRV/if.mlir

 // RUN: mlir-opt -convert-scf-to-spirv %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @kernel_simple_selection
 func.func @kernel_simple_selection(%arg2 : memref<10xf32>, %arg3 : i1) {
   %value = arith.constant 0.0 : f32
   %i = arith.constant 0 : index
 
   // CHECK:       spv.mlir.selection {

mlir/test/Conversion/SCFToSPIRV/while.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -convert-scf-to-spirv %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader, Int64], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader, Int64], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: @while_loop1
 func.func @while_loop1(%arg0: i32, %arg1: i32) -> i32 {
   // CHECK-SAME: (%[[ARG1:.*]]: i32, %[[ARG2:.*]]: i32)
   // CHECK: %[[INITVAR:.*]] = spv.Constant 2 : i32
   // CHECK: %[[VAR1:.*]] = spv.Variable : !spv.ptr<i32, Function>
   // CHECK: spv.mlir.loop {

mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir

 // RUN: mlir-opt --lower-host-to-llvm %s | FileCheck %s
 
-module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, {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<v1.0, [Shader], [SPV_KHR_variable_pointers]>, #spv.resource_limits<max_compute_workgroup_invocations = 128, max_compute_workgroup_size = [128, 128, 64]>>} {
 
   //       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()
 
   //       CHECK: spv.module @__spv__foo
   //       CHECK:   spv.GlobalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
   //       CHECK:   spv.func @__spv__foo_bar
 
spv.func @bar() "None" attributes {workgroup_attributions = 0 : i64} {
       %0 = spv.mlir.addressof @bar_arg_0 : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
       spv.Return
     }
     spv.EntryPoint "GLCompute" @bar
     spv.ExecutionMode @bar "LocalSize", 1, 1, 1
   }
 
   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<local_size = dense<1> : vector<3xi32>>} {
       gpu.return
     }
   }
 
   func.func @main() {
     %buffer = memref.alloc() : memref<6xi32>
     %one = arith.constant 1 : index
     gpu.launch_func @foo::@bar blocks in (%one, %one, %one)
         threads in (%one, %one, %one) args(%buffer : memref<6xi32>)
     return
   }
 }

mlir/test/Conversion/VectorToSPIRV/vector-to-spirv.mlir

 // 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<v1.0, [Float16], []>, {}> } {
+module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.0, [Float16], []>, #spv.resource_limits<>> } {
 
 // CHECK-LABEL: @bitcast
 //  CHECK-SAME: %[[ARG0:.+]]: vector<2xf32>, %[[ARG1:.+]]: vector<2xf16>
 //       CHECK:   spv.Bitcast %[[ARG0]] : vector<2xf32> to vector<4xf16>
 //       CHECK:   spv.Bitcast %[[ARG1]] : vector<2xf16> to f32
 func.func @bitcast(%arg0 : vector<2xf32>, %arg1: vector<2xf16>) -> (vector<4xf16>, vector<1xf32>) {
   %0 = vector.bitcast %arg0 : vector<2xf32> to vector<4xf16>
   %1 = vector.bitcast %arg1 : vector<2xf16> to vector<1xf32>

mlir/test/Dialect/SPIRV/IR/target-and-abi.mlir

 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // 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<local_size = 64>
 } { 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<local_size = dense<[64, 1, 1]> : vector<3xi32>>}
+  spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[64, 1, 1]>: vector<3xi32>>
 } { return }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.interface_var_abi
 //===----------------------------------------------------------------------===//
 

 ) { return }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.target_env
 //===----------------------------------------------------------------------===//
 
-func.func @target_env_wrong_limits() attributes {
-  spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
-    // 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<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
-  // CHECK-SAME:   {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
+  // CHECK-SAME:   #spv.resource_limits<max_compute_workgroup_size = [128, 64, 64]>>
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
-    {
-      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 }
 
 // -----
 
 func.func @target_env_vendor_id() attributes {
   // CHECK:      spv.target_env = #spv.target_env<
   // CHECK-SAME:   #spv.vce<v1.0, [], []>,
   // CHECK-SAME:   NVIDIA,
-  // CHECK-SAME:   {}>
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, NVIDIA, {}>
+  // CHECK-SAME:   #spv.resource_limits<>>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, NVIDIA, #spv.resource_limits<>>
 } { return }
 
 // -----
 
 func.func @target_env_vendor_id_device_type() attributes {
   // CHECK:      spv.target_env = #spv.target_env<
   // CHECK-SAME:   #spv.vce<v1.0, [], []>,
   // CHECK-SAME:   AMD:DiscreteGPU,
-  // CHECK-SAME:   {}>
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, AMD:DiscreteGPU, {}>
+  // CHECK-SAME:   #spv.resource_limits<>>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, AMD:DiscreteGPU, #spv.resource_limits<>>
 } { return }
 
 // -----
 
 func.func @target_env_vendor_id_device_type_device_id() attributes {
   // CHECK:      spv.target_env = #spv.target_env<
   // CHECK-SAME:   #spv.vce<v1.0, [], []>,
   // CHECK-SAME:   Qualcomm:IntegratedGPU:100925441,
-  // CHECK-SAME:   {}>
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, Qualcomm:IntegratedGPU:0x6040001, {}>
+  // CHECK-SAME:   #spv.resource_limits<>>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, 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<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
-    {
-      max_compute_workgroup_invocations = 128 : i32,
-      max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>
-    },
+    #spv.resource_limits<>,
     more_stuff
   >
 } { return }
 
 // -----
 
 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<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class,
                             SPV_NV_cooperative_matrix]>,
-  {
-    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 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.vce
 //===----------------------------------------------------------------------===//
 

mlir/test/Dialect/SPIRV/IR/target-env.mlir

 // which extensions are incorporated into SPIR-V 1.5.
 
 //===----------------------------------------------------------------------===//
 // MaxVersion
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @cmp_exchange_weak_suitable_version_capabilities
 func.func @cmp_exchange_weak_suitable_version_capabilities(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.1, [Kernel, AtomicStorage], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.1, [Kernel, AtomicStorage], []>, #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, Workgroup>, i32, i32) -> (i32)
   return %0: i32
 }
 
 // CHECK-LABEL: @cmp_exchange_weak_unsupported_version
 func.func @cmp_exchange_weak_unsupported_version(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.4, [Kernel, AtomicStorage], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.4, [Kernel, AtomicStorage], []>, #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, Workgroup>, i32, i32) -> (i32)
   return %0: i32
 }
 
 //===----------------------------------------------------------------------===//
 // MinVersion
 //===----------------------------------------------------------------------===//
 
 // 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<v1.4, [GroupNonUniformBallot], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.4, [GroupNonUniformBallot], []>, #spv.resource_limits<>>
 } {
   // CHECK: spv.GroupNonUniformBallot Workgroup
   %0 = "test.convert_to_group_non_uniform_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
   return %0: vector<4xi32>
 }
 
 // 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<v1.1, [GroupNonUniformBallot], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.1, [GroupNonUniformBallot], []>, #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>)
   return %0: vector<4xi32>
 }
 
 //===----------------------------------------------------------------------===//
 // Capability
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @cmp_exchange_weak_missing_capability_kernel
 func.func @cmp_exchange_weak_missing_capability_kernel(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.3, [AtomicStorage], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.3, [AtomicStorage], []>, #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, Workgroup>, i32, i32) -> (i32)
   return %0: i32
 }
 
 // CHECK-LABEL: @cmp_exchange_weak_missing_capability_atomic_storage
 func.func @cmp_exchange_weak_missing_capability_atomic_storage(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.3, [Kernel], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.3, [Kernel], []>, #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, Workgroup>, i32, i32) -> (i32)
   return %0: i32
 }
 
 // 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<v1.4, [], [SPV_KHR_shader_ballot]>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.4, [], [SPV_KHR_shader_ballot]>, #spv.resource_limits<>>
 } {
   // CHECK: test.convert_to_subgroup_ballot_op
   %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
   return %0: vector<4xi32>
 }
 
 // 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<v1.0, [Geometry], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Geometry], []>, #spv.resource_limits<>>
 } {
   // CHECK: spv.BitReverse
   %0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32)
   return %0: i32
 }
 
 // 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<v1.0, [GeometryPointSize], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [GeometryPointSize], []>, #spv.resource_limits<>>
 } {
   // CHECK: spv.BitReverse
   %0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32)
   return %0: i32
 }
 
 //===----------------------------------------------------------------------===//
 // Extension
 //===----------------------------------------------------------------------===//
 
 // 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<v1.4, [SubgroupBallotKHR], [SPV_KHR_shader_ballot]>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], [SPV_KHR_shader_ballot]>, #spv.resource_limits<>>
 } {
   // CHECK: spv.SubgroupBallotKHR
   %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
   return %0: vector<4xi32>
 }
 
 // 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<v1.4, [SubgroupBallotKHR], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], []>, #spv.resource_limits<>>
 } {
   // CHECK: test.convert_to_subgroup_ballot_op
   %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
   return %0: vector<4xi32>
 }
 
 // CHECK-LABEL: @module_suitable_extension1
 func.func @module_suitable_extension1() attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_EXT_physical_storage_buffer]>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_EXT_physical_storage_buffer]>, #spv.resource_limits<>>
 } {
   // CHECK: spv.module PhysicalStorageBuffer64 Vulkan
   "test.convert_to_module_op"() : () ->()
   return
 }
 
 // CHECK-LABEL: @module_suitable_extension2
 func.func @module_suitable_extension2() attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_KHR_physical_storage_buffer]>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_KHR_physical_storage_buffer]>, #spv.resource_limits<>>
 } {
   // CHECK: spv.module PhysicalStorageBuffer64 Vulkan
   "test.convert_to_module_op"() : () -> ()
   return
 }
 
 // CHECK-LABEL: @module_missing_extension_mm
 func.func @module_missing_extension_mm() attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_physical_storage_buffer]>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_physical_storage_buffer]>, #spv.resource_limits<>>
 } {
   // CHECK: test.convert_to_module_op
   "test.convert_to_module_op"() : () -> ()
   return
 }
 
 // CHECK-LABEL: @module_missing_extension_am
 func.func @module_missing_extension_am() attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model]>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model]>, #spv.resource_limits<>>
 } {
   // CHECK: test.convert_to_module_op
   "test.convert_to_module_op"() : () -> ()
   return
 }
 
 // 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<v1.5, [VulkanMemoryModel, PhysicalStorageBufferAddresses], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.5, [VulkanMemoryModel, PhysicalStorageBufferAddresses], []>, #spv.resource_limits<>>
 } {
   // CHECK: spv.module PhysicalStorageBuffer64 Vulkan
   "test.convert_to_module_op"() : () -> ()
   return
 }

mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir

spv.module Logical GLSL450 {
 
   spv.func @baz_no_return_another_control(%arg0: i32) -> () "Inline|Pure" {
     spv.Return
   }
 
   spv.func @kernel(
     %arg0: f32,
     %arg1: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, CrossWorkgroup>) "None"
-  attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
+  attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
     spv.Return
   }
 
   spv.func @kernel_different_attr(
     %arg0: f32,
     %arg1: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, CrossWorkgroup>) "None"
-  attributes {spv.entry_point_abi = {local_size = dense<[64, 1, 1]> : vector<3xi32>}} {
+  attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[64, 1, 1]> : vector<3xi32>>} {
     spv.Return
   }
 }
 }

mlir/test/Dialect/SPIRV/Transforms/abi-interface-opencl.mlir

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
-  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel, Addresses], []>, {}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Kernel, Addresses], []>, #spv.resource_limits<>>
 } {
   spv.module Physical64 OpenCL {
     // CHECK-LABEL: spv.module
     //       CHECK:   spv.func [[FN:@.*]]({{%.*}}: f32, {{%.*}}: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, CrossWorkgroup>
     //       CHECK:   spv.EntryPoint "Kernel" [[FN]]
     //       CHECK:   spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
     spv.func @kernel(
       %arg0: f32,
       %arg1: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, CrossWorkgroup>) "None"
-    attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
+    attributes {spv.entry_point_abi = #spv.entry_point_abi<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
       spv.Return
     }
   }
 }

mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.module
 spv.module Logical GLSL450 {
   // CHECK-DAG:    spv.GlobalVariable [[VAR0:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<(f32 [0])>, StorageBuffer>
   // CHECK-DAG:    spv.GlobalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
   // CHECK:    spv.func [[FN:@.*]]()
   spv.func @kernel(
     %arg0: f32
            {spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>},
     %arg1: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, 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<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
     // CHECK: [[ARG1:%.*]] = spv.mlir.addressof [[VAR1]]
     // CHECK: [[ADDRESSARG0:%.*]] = spv.mlir.addressof [[VAR0]]
     // CHECK: [[CONST0:%.*]] = spv.Constant 0 : i32
     // CHECK: [[ARG0PTR:%.*]] = spv.AccessChain [[ADDRESSARG0]]{{\[}}[[CONST0]]
     // CHECK: [[ARG0:%.*]] = spv.Load "StorageBuffer" [[ARG0PTR]]
     // CHECK: spv.Return
     spv.Return
   }
   // CHECK: spv.EntryPoint "GLCompute" [[FN]]
   // CHECK: spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
 } // end spv.module
 
 } // end module

mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spv.resource_limits<>>
 } {
 
 // CHECK-LABEL: spv.module
 spv.module Logical GLSL450 {
   // CHECK-DAG: spv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
   spv.GlobalVariable @__builtin_var_WorkgroupSize__ built_in("WorkgroupSize") : !spv.ptr<vector<3xi32>, Input>
   // CHECK-DAG: spv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
   spv.GlobalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
spv.func @load_store_kernel(
     %arg3: i32
     {spv.interface_var_abi = #spv.interface_var_abi<(0, 3), StorageBuffer>},
     %arg4: i32
     {spv.interface_var_abi = #spv.interface_var_abi<(0, 4), StorageBuffer>},
     %arg5: i32
     {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<local_size = dense<[32, 1, 1]> : vector<3xi32>>} {
     // CHECK: [[ADDRESSARG6:%.*]] = spv.mlir.addressof [[VAR6]]
     // CHECK: [[CONST6:%.*]] = spv.Constant 0 : i32
     // CHECK: [[ARG6PTR:%.*]] = spv.AccessChain [[ADDRESSARG6]]{{\[}}[[CONST6]]
     // CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG6PTR]]
     // CHECK: [[ADDRESSARG5:%.*]] = spv.mlir.addressof [[VAR5]]
     // CHECK: [[CONST5:%.*]] = spv.Constant 0 : i32
     // CHECK: [[ARG5PTR:%.*]] = spv.AccessChain [[ADDRESSARG5]]{{\[}}[[CONST5]]
     // CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG5PTR]]

mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir

 // RUN: mlir-opt -spirv-update-vce %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // Version
 //===----------------------------------------------------------------------===//
 
 // Test deducing minimal version.
 // spv.IAdd is available from v1.0.
 
 // CHECK: requires #spv.vce<v1.0, [Shader], []>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.5, [Shader], []>, {}>
+    #spv.vce<v1.5, [Shader], []>, #spv.resource_limits<>>
 } {
   spv.func @iadd(%val : i32) -> i32 "None" {
     %0 = spv.IAdd %val, %val: i32
     spv.ReturnValue %0: i32
   }
 }
 
 // Test deducing minimal version.
 // spv.GroupNonUniformBallot is available since v1.3.
 
 // CHECK: requires #spv.vce<v1.3, [GroupNonUniformBallot, Shader], []>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.5, [Shader, GroupNonUniformBallot], []>, {}>
+    #spv.vce<v1.5, [Shader, GroupNonUniformBallot], []>, #spv.resource_limits<>>
 } {
   spv.func @group_non_uniform_ballot(%predicate : i1) -> vector<4xi32> "None" {
     %0 = spv.GroupNonUniformBallot Workgroup %predicate : vector<4xi32>
     spv.ReturnValue %0: vector<4xi32>
   }
 }
 
 //===----------------------------------------------------------------------===//
 // Capability
 //===----------------------------------------------------------------------===//
 
 // Test minimal capabilities.
 
 // CHECK: requires #spv.vce<v1.0, [Shader], []>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader, Float16, Float64, Int16, Int64, VariablePointers], []>, {}>
+    #spv.vce<v1.0, [Shader, Float16, Float64, Int16, Int64, VariablePointers], []>, #spv.resource_limits<>>
 } {
   spv.func @iadd(%val : i32) -> i32 "None" {
     %0 = spv.IAdd %val, %val: i32
     spv.ReturnValue %0: i32
   }
 }
 
 // Test Physical Storage Buffers are deduced correctly.
 
 // CHECK: spv.module PhysicalStorageBuffer64 GLSL450 requires #spv.vce<v1.0, [PhysicalStorageBufferAddresses, Shader], [SPV_EXT_physical_storage_buffer]>
 spv.module PhysicalStorageBuffer64 GLSL450 attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader, PhysicalStorageBufferAddresses], [SPV_EXT_physical_storage_buffer]>, {}>
+    #spv.vce<v1.0, [Shader, PhysicalStorageBufferAddresses], [SPV_EXT_physical_storage_buffer]>, #spv.resource_limits<>>
 } {
   spv.func @physical_ptr(%val : !spv.ptr<f32, PhysicalStorageBuffer>) "None" {
     spv.Return
   }
 }
 
 // Test deducing implied capability.
 // AtomicStorage implies Shader.
 
 // CHECK: requires #spv.vce<v1.0, [Shader], []>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [AtomicStorage], []>, {}>
+    #spv.vce<v1.0, [AtomicStorage], []>, #spv.resource_limits<>>
 } {
   spv.func @iadd(%val : i32) -> i32 "None" {
     %0 = spv.IAdd %val, %val: i32
     spv.ReturnValue %0: i32
   }
 }
 
 // Test selecting the capability available in the target environment.
 // spv.GroupNonUniform op itself can be enabled via any of
 // * GroupNonUniformArithmetic
 // * GroupNonUniformClustered
 // * GroupNonUniformPartitionedNV
 // Its 'Reduce' group operation can be enabled via any of
 // * Kernel
 // * GroupNonUniformArithmetic
 // * GroupNonUniformBallot
 
 // CHECK: requires #spv.vce<v1.3, [GroupNonUniformArithmetic, Shader], []>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
+    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, #spv.resource_limits<>>
 } {
   spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" {
     %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32
     spv.ReturnValue %0: i32
   }
 }
 
 // CHECK: requires #spv.vce<v1.3, [GroupNonUniformClustered, GroupNonUniformBallot, Shader], []>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.3, [Shader, GroupNonUniformClustered, GroupNonUniformBallot], []>, {}>
+    #spv.vce<v1.3, [Shader, GroupNonUniformClustered, GroupNonUniformBallot], []>, #spv.resource_limits<>>
 } {
   spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" {
     %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32
     spv.ReturnValue %0: i32
   }
 }
 
 // Test type required capabilities
 
 // Using 8-bit integers in non-interface storage class requires Int8.
 // CHECK: requires #spv.vce<v1.0, [Int8, Shader], []>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.3, [Shader, Int8], []>, {}>
+    #spv.vce<v1.3, [Shader, Int8], []>, #spv.resource_limits<>>
 } {
   spv.func @iadd_function(%val : i8) -> i8 "None" {
     %0 = spv.IAdd %val, %val : i8
     spv.ReturnValue %0: i8
   }
 }
 
 // Using 16-bit floats in non-interface storage class requires Float16.
 // CHECK: requires #spv.vce<v1.0, [Float16, Shader], []>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.3, [Shader, Float16], []>, {}>
+    #spv.vce<v1.3, [Shader, Float16], []>, #spv.resource_limits<>>
 } {
   spv.func @fadd_function(%val : f16) -> f16 "None" {
     %0 = spv.FAdd %val, %val : f16
     spv.ReturnValue %0: f16
   }
 }
 
 // Using 16-element vectors requires Vector16.
 // CHECK: requires #spv.vce<v1.0, [Vector16, Shader], []>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.3, [Shader, Vector16], []>, {}>
+    #spv.vce<v1.3, [Shader, Vector16], []>, #spv.resource_limits<>>
 } {
   spv.func @iadd_v16_function(%val : vector<16xi32>) -> vector<16xi32> "None" {
     %0 = spv.IAdd %val, %val : vector<16xi32>
     spv.ReturnValue %0: vector<16xi32>
   }
 }
 
 //===----------------------------------------------------------------------===//
 // Extension
 //===----------------------------------------------------------------------===//
 
 // Test deducing minimal extensions.
 // spv.SubgroupBallotKHR requires the SPV_KHR_shader_ballot extension.
 
 // CHECK: requires #spv.vce<v1.0, [SubgroupBallotKHR, Shader], [SPV_KHR_shader_ballot]>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader, SubgroupBallotKHR],
-             [SPV_KHR_shader_ballot, SPV_KHR_shader_clock, SPV_KHR_variable_pointers]>, {}>
+             [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>
     spv.ReturnValue %0: vector<4xi32>
   }
 }
 
 // Test deducing implied extension.
 // Vulkan memory model requires SPV_KHR_vulkan_memory_model, which is enabled
 // implicitly by v1.5.
 
 // CHECK: requires #spv.vce<v1.0, [VulkanMemoryModel], [SPV_KHR_vulkan_memory_model]>
 spv.module Logical Vulkan attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.5, [Shader, VulkanMemoryModel], []>, {}>
+    #spv.vce<v1.5, [Shader, VulkanMemoryModel], []>, #spv.resource_limits<>>
 } {
   spv.func @iadd(%val : i32) -> i32 "None" {
     %0 = spv.IAdd %val, %val: i32
     spv.ReturnValue %0: i32
   }
 }
 
 // Test type required extensions
 
 // Using 8-bit integers in interface storage class requires additional
 // extensions and capabilities.
 // CHECK: requires #spv.vce<v1.0, [StorageBuffer16BitAccess, Shader, Int16], [SPV_KHR_16bit_storage, SPV_KHR_storage_buffer_storage_class]>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.3, [Shader, StorageBuffer16BitAccess, Int16], []>, {}>
+    #spv.vce<v1.3, [Shader, StorageBuffer16BitAccess, Int16], []>, #spv.resource_limits<>>
 } {
   spv.func @iadd_storage_buffer(%ptr : !spv.ptr<i16, StorageBuffer>) -> i16 "None" {
     %0 = spv.Load "StorageBuffer" %ptr : i16
     %1 = spv.IAdd %0, %0 : i16
     spv.ReturnValue %1: i16
   }
 }
 
 // Complicated nested types
 // * Buffer requires ImageBuffer or SampledBuffer.
 // * Rg32f requires StorageImageExtendedFormats.
 // CHECK: requires #spv.vce<v1.0, [UniformAndStorageBuffer8BitAccess, StorageUniform16, Int64, Shader, ImageBuffer, StorageImageExtendedFormats], [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.5, [Shader, UniformAndStorageBuffer8BitAccess, StorageBuffer16BitAccess, StorageUniform16, Int16, Int64, ImageBuffer, StorageImageExtendedFormats], []>,
-    {}>
+    #spv.resource_limits<>>
 } {
   spv.GlobalVariable @data : !spv.ptr<!spv.struct<(i8 [0], f16 [2], i64 [4])>, Uniform>
   spv.GlobalVariable @img  : !spv.ptr<!spv.image<f32, Buffer, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Rg32f>, UniformConstant>
 }

mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp

struct TestSpirvEntryPointABIPass
   MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestSpirvEntryPointABIPass)
 
   StringRef getArgument() const final { return "test-spirv-entry-point-abi"; }
   StringRef getDescription() const final {
     return "Set the spv.entry_point_abi attribute on GPU kernel function "
            "within the "
            "module, intended for testing only";
   }
+  void getDependentDialects(DialectRegistry &registry) const override {
+    registry.insert<spirv::SPIRVDialect>();
+  }
   TestSpirvEntryPointABIPass() = default;
   TestSpirvEntryPointABIPass(const TestSpirvEntryPointABIPass &) {}
   void runOnOperation() override;
 
 private:
   Pass::ListOption<int32_t> workgroupSize{
       *this, "workgroup-size",
       llvm::cl::desc(

mlir/test/mlir-spirv-cpu-runner/double.mlir

 // RUN: mlir-spirv-cpu-runner %s -e main --entry-point-result=void --shared-libs=%mlir_runner_utils_dir/libmlir_runner_utils%shlibext,%spirv_wrapper_library_dir/libmlir_test_spirv_cpu_runner_c_wrappers%shlibext
 
 // CHECK: [8,  8,  8,  8,  8,  8]
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>,
-    {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<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
       %factor = arith.constant 2 : i32
 
       %i0 = arith.constant 0 : index
       %i1 = arith.constant 1 : index
       %i2 = arith.constant 2 : index
       %i3 = arith.constant 3 : index
       %i4 = arith.constant 4 : index
       %i5 = arith.constant 5 : index

mlir/test/mlir-spirv-cpu-runner/simple_add.mlir

 // RUN: mlir-spirv-cpu-runner %s -e main --entry-point-result=void --shared-libs=%mlir_runner_utils_dir/libmlir_runner_utils%shlibext,%spirv_wrapper_library_dir/libmlir_test_spirv_cpu_runner_c_wrappers%shlibext
 
 // CHECK: [[[7.7,    0,    0], [7.7,    0,    0], [7.7,    0,    0]], [[0,    7.7,    0], [0,    7.7,    0], [0,    7.7,    0]], [[0,    0,    7.7], [0,    0,    7.7], [0,    0,    7.7]]]
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_8bit_storage]>,
-    {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<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
       %i0 = arith.constant 0 : index
       %i1 = arith.constant 1 : index
       %i2 = arith.constant 2 : index
 
       %x = memref.load %arg0[%i0] : memref<3xf32>
       %y = memref.load %arg1[%i0, %i0] : memref<3x3xf32>
       %sum = arith.addf %x, %y : f32
 

mlir/test/mlir-vulkan-runner/addf.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%vulkan_wrapper_library_dir/libvulkan-runtime-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK: [3.3,  3.3,  3.3,  3.3,  3.3,  3.3,  3.3,  3.3]
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #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<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
       %0 = gpu.block_id x
       %1 = memref.load %arg0[%0] : memref<8xf32>
       %2 = memref.load %arg1[%0] : memref<8xf32>
       %3 = arith.addf %1, %2 : f32
       memref.store %3, %arg2[%0] : memref<8xf32>
       gpu.return
     }
   }

mlir/test/mlir-vulkan-runner/addi.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%vulkan_wrapper_library_dir/libvulkan-runtime-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK-COUNT-64: [3, 3, 3, 3, 3, 3, 3, 3]
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #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<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
       %x = gpu.block_id x
       %y = gpu.block_id y
       %z = gpu.block_id z
       %0 = memref.load %arg0[%x] : memref<8xi32>
       %1 = memref.load %arg1[%y, %x] : memref<8x8xi32>
       %2 = arith.addi %0, %1 : i32
       memref.store %2, %arg2[%z, %y, %x] : memref<8x8x8xi32>
       gpu.return

mlir/test/mlir-vulkan-runner/addi8.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%vulkan_wrapper_library_dir/libvulkan-runtime-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK-COUNT-64: [3, 3, 3, 3, 3, 3, 3, 3]
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_8bit_storage]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_8bit_storage]>, #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<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
       %x = gpu.block_id x
       %y = gpu.block_id y
       %z = gpu.block_id z
       %0 = memref.load %arg0[%x] : memref<8xi8>
       %1 = memref.load %arg1[%y, %x] : memref<8x8xi8>
       %2 = arith.addi %0, %1 : i8
       %3 = arith.extui %2 : i8 to i32
       memref.store %3, %arg2[%z, %y, %x] : memref<8x8x8xi32>

mlir/test/mlir-vulkan-runner/mulf.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%vulkan_wrapper_library_dir/libvulkan-runtime-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK-COUNT-4: [6, 6, 6, 6]
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #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<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
       %x = gpu.block_id x
       %y = gpu.block_id y
       %1 = memref.load %arg0[%x, %y] : memref<4x4xf32>
       %2 = memref.load %arg1[%x, %y] : memref<4x4xf32>
       %3 = arith.mulf %1, %2 : f32
       memref.store %3, %arg2[%x, %y] : memref<4x4xf32>
       gpu.return
     }

mlir/test/mlir-vulkan-runner/subf.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%vulkan_wrapper_library_dir/libvulkan-runtime-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK-COUNT-32: [2.2, 2.2, 2.2, 2.2]
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, 
+    #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<local_size = dense<[1, 1, 1]>: vector<3xi32>>} {
       %x = gpu.block_id x
       %y = gpu.block_id y
       %z = gpu.block_id z
       %1 = memref.load %arg0[%x, %y, %z] : memref<8x4x4xf32>
       %2 = memref.load %arg1[%y, %z] : memref<4x4xf32>
       %3 = arith.subf %1, %2 : f32
       memref.store %3, %arg2[%x, %y, %z] : memref<8x4x4xf32>
       gpu.return

mlir/test/mlir-vulkan-runner/time.mlir

 // RUN: mlir-vulkan-runner %s --shared-libs=%vulkan_wrapper_library_dir/libvulkan-runtime-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s
 
 // CHECK: Compute shader execution time
 // CHECK: Command buffer submit time
 // CHECK: Wait idle time
 
 module attributes {
   gpu.container_module,
   spv.target_env = #spv.target_env<
-    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #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<local_size = dense<[128, 1, 1]>: vector<3xi32>>} {
       %bid = gpu.block_id x
       %tid = gpu.thread_id x
       %cst = arith.constant 128 : index
       %b = arith.muli %bid, %cst : index
       %0 = arith.addi %b, %tid : index
       %1 = memref.load %arg0[%0] : memref<16384xf32>
       %2 = memref.load %arg1[%0] : memref<16384xf32>
       %3 = arith.addf %1, %2 : f32

mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp

 //===----------------------------------------------------------------------===//
 // ParserGen
 //===----------------------------------------------------------------------===//
 
 void DefFormat::genParser(MethodBody &os) {
   FmtContext ctx;
   ctx.addSubst("_parser", "odsParser");
   ctx.addSubst("_ctx", "odsParser.getContext()");
+  ctx.withBuilder("odsBuilder");
   if (isa<AttrDef>(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
   // FailureOr<T> to defer type construction for parameters that are parsed in
   // a loop (parsers return FailureOr anyways).
   ArrayRef<AttrOrTypeParameter> params = def.getParameters();
   for (const AttrOrTypeParameter &param : params) {
     if (isa<AttributeSelfTypeParameter>(param))

 //===----------------------------------------------------------------------===//
 // PrinterGen
 //===----------------------------------------------------------------------===//
 
 void DefFormat::genPrinter(MethodBody &os) {
   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;
   lastWasPunctuation = false;
   for (FormatElement *el : elements)
     genElementPrinter(el, ctx, os);
 }
 

utils/bazel/llvm-project-overlay/mlir/BUILD.bazel

tbl_outs = [
         ),
     ],
     tblgen = ":mlir-tblgen",
     td_file = "include/mlir/Dialect/SPIRV/IR/SPIRVOps.td",
     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",
     tbl_outs = [
         (
             ["-gen-rewriters"],
             "lib/Dialect/SPIRV/IR/SPIRVCanonicalization.inc",
         ),
     ],
tbl_outs = [
         ),
     ],
     tblgen = ":mlir-tblgen",
     td_file = "include/mlir/Dialect/SPIRV/IR/SPIRVOps.td",
     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",
     tbl_outs = [
         (
             ["-gen-spirv-attr-utils"],
             "include/mlir/Dialect/SPIRV/IR/SPIRVAttrUtils.inc",
         ),
     ],
cc_library(
     includes = ["include"],
     deps = [
         ":CommonFolders",
         ":ControlFlowInterfaces",
         ":IR",
         ":InferTypeOpInterface",
         ":Parser",
         ":SPIRVAttrUtilsGen",
+        ":SPIRVAttributesIncGen",
         ":SPIRVAvailabilityIncGen",
         ":SPIRVCanonicalizationIncGen",
         ":SPIRVOpsIncGen",
         ":SPIRVSerializationGen",
-        ":SPIRVTargetAndABIStructGen",
         ":SideEffectInterfaces",
         ":Support",
         ":Transforms",
         "//llvm:Support",
     ],
 )
 
 gentbl_cc_library(