[mlir][spirv] Use spv.vce in spv.module and wire up (de)serialization

[antiagainst]

Summary

This commits changes the definition of spv.module to use the #spv.vce
attribute for specifying (version, capabilities, extensions) triple
so that we can have better API and custom assembly form. Since now
we have proper modelling of the triple, (de)serialization is wired up
to use them.

With the new UpdateVCEPass, we don't need to manually specify the
required extensions and capabilities anymore when creating a spv.module.
One just need to call UpdateVCEPass before serialization to get the
needed version/extensions/capabilities.

Depends On D75871

[mravishankar]

Nice cleanup!

[mehdi_amini]

mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp

1

Doesn't seem intended?

[antiagainst]

Address comments

[antiagainst]

mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp

1

Oops! Thanks for the good catch!

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

 def SPV_ISUI_NoSampler      : I32EnumAttrCase<"NoSampler", 2>;
 
 def SPV_SamplerUseAttr:
     SPV_I32EnumAttr<
       "ImageSamplerUseInfo", "valid SPIR-V Sampler Use specification",
       [SPV_ISUI_SamplerUnknown, SPV_ISUI_NeedSampler, SPV_ISUI_NoSampler]>;
 
 //===----------------------------------------------------------------------===//
+// SPIR-V attribute definitions
+//===----------------------------------------------------------------------===//
+
+def SPV_VerCapExtAttr : Attr<
+    CPred<"$_self.isa<::mlir::spirv::VerCapExtAttr>()">,
+    "version-capability-extension attribute"> {
+  let storageType = "::mlir::spirv::VerCapExtAttr";
+  let returnType = "::mlir::spirv::VerCapExtAttr";
+  let convertFromStorage = "$_self";
+}
+
+//===----------------------------------------------------------------------===//
 // SPIR-V type definitions
 //===----------------------------------------------------------------------===//
 
 class IOrUI<int width>
     : Type<Or<[CPred<"$_self.isSignlessInteger(" # width # ")">,
                CPred<"$_self.isUnsignedInteger(" # width # ")">]>,
            width # "-bit signless/unsigned integer"> {
   int bitwidth = width;

mlir/include/mlir/Dialect/SPIRV/SPIRVBinaryUtils.h

 
 /// The serializer tool ID registered to the Khronos Group
 constexpr uint32_t kGeneratorNumber = 22;
 
 /// Auto-generated getOpcode<*Op>() specializations
 #define GET_SPIRV_SERIALIZATION_UTILS
 #include "mlir/Dialect/SPIRV/SPIRVSerialization.inc"
 
-/// Appends a SPRI-V module header to `header` with the given `idBound`.
-void appendModuleHeader(SmallVectorImpl<uint32_t> &header, uint32_t idBound);
+/// Appends a SPRI-V module header to `header` with the given `version` and
+/// `idBound`.
+void appendModuleHeader(SmallVectorImpl<uint32_t> &header,
+                        spirv::Version version, uint32_t idBound);
 
 /// Returns the word-count-prefixed opcode for an SPIR-V instruction.
 uint32_t getPrefixedOpcode(uint32_t wordCount, spirv::Opcode opcode);
 
 /// Encodes an SPIR-V `literal` string into the given `binary` vector.
 LogicalResult encodeStringLiteralInto(SmallVectorImpl<uint32_t> &binary,
                                       StringRef literal);
 } // end namespace spirv
 } // end namespace mlir
 
 #endif // MLIR_DIALECT_SPIRV_SPIRV_BINARY_UTILS_H_

mlir/include/mlir/Dialect/SPIRV/SPIRVOps.h

 #include "mlir/Interfaces/ControlFlowInterfaces.h"
 #include "mlir/Interfaces/SideEffects.h"
 #include "llvm/Support/PointerLikeTypeTraits.h"
 
 namespace mlir {
 class OpBuilder;
 
 namespace spirv {
+class VerCapExtAttr;
 
 // TableGen'erated operation interfaces for querying versions, extensions, and
 // capabilities.
 #include "mlir/Dialect/SPIRV/SPIRVAvailability.h.inc"
 
 // TablenGen'erated operation declarations.
 #define GET_OP_CLASSES
 #include "mlir/Dialect/SPIRV/SPIRVOps.h.inc"

mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td

let description = [{
     implicitly capture values from the enclosing environment.
 
     This op has only one region, which only contains one block. The block
     must be terminated via the `spv._module_end` op.
 
     ### Custom assembly form
 
     ```
-    addressing-model ::= `"Logical"` | `"Physical32"` | `"Physical64"`
-    memory-model ::= `"Simple"` | `"GLSL450"` | `"OpenCL"` | `"VulkanKHR"`
+    addressing-model ::= `Logical` | `Physical32` | `Physical64` | ...
+    memory-model ::= `Simple` | `GLSL450` | `OpenCL` | `Vulkan` | ...
     spv-module-op ::= `spv.module` addressing-model memory-model
-                      region
+                      (requires  spirv-vce-attribute)?
                       (`attributes` attribute-dict)?
+                      region
     ```
 
     For example:
 
     ```
-    spv.module "Logical" "VulkanKHR" { }
+    spv.module Logical GLSL450  {}
 
-    spv.module "Logical" "VulkanKHR" {
-      func @do_nothing() -> () {
+    spv.module Logical Vulkan
+        requires #spv.vce<v1.0, [Shader], [SPV_KHR_vulkan_memory_model]>
+        attributes { some_additional_attr = ... } {
+      spv.func @do_nothing() -> () {
         spv.Return
       }
-    } attributes {
-      capability = ["Shader"],
-      extension = ["SPV_KHR_16bit_storage"]
     }
     ```
   }];
 
   let arguments = (ins
     SPV_AddressingModelAttr:$addressing_model,
     SPV_MemoryModelAttr:$memory_model,
-    OptionalAttr<StrArrayAttr>:$capabilities,
-    OptionalAttr<StrArrayAttr>:$extensions,
-    OptionalAttr<StrArrayAttr>:$extended_instruction_sets
+    OptionalAttr<SPV_VerCapExtAttr>:$vce_triple
   );
 
   let results = (outs);
 
   let regions = (region SizedRegion<1>:$body);
 
-  let builders =
-    [OpBuilder<"Builder *, OperationState &state">,
-     OpBuilder<[{Builder *, OperationState &state,
-                 IntegerAttr addressing_model,
-                 IntegerAttr memory_model}]>,
-     OpBuilder<[{Builder *, OperationState &state,
-                 spirv::AddressingModel addressing_model,
-                 spirv::MemoryModel memory_model,
-                 /*optional*/ ArrayRef<spirv::Capability> capabilities = {},
-                 /*optional*/ ArrayRef<spirv::Extension> extensions = {},
-                 /*optional*/ ArrayAttr extended_instruction_sets = nullptr}]>];
+  let builders = [
+    OpBuilder<[{Builder *, OperationState &state}]>,
+                spirv::MemoryModel memory_model}]>
+  ];
 
   // We need to ensure the block inside the region is properly terminated;
   // the auto-generated builders do not guarantee that.
   let skipDefaultBuilders = 1;
 
   let hasOpcode = 0;
 
   let autogenSerialization = 0;
 
   let extraClassDeclaration = [{
+    static StringRef getVCETripleAttrName() { return "vce_triple"; }
+
     Block& getBlock() {
       return this->getOperation()->getRegion(0).front();
     }
   }];
 }
 
 // -----
 

mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp

 
 //===----------------------------------------------------------------------===//
 // ModuleOp with gpu.module.
 //===----------------------------------------------------------------------===//
 
 PatternMatchResult GPUModuleConversion::matchAndRewrite(
     gpu::GPUModuleOp moduleOp, ArrayRef<Value> operands,
     ConversionPatternRewriter &rewriter) const {
-  // TODO : Generalize this to account for different extensions,
-  // capabilities, extended_instruction_sets, other addressing models
-  // and memory models.
   auto spvModule = rewriter.create<spirv::ModuleOp>(
       moduleOp.getLoc(), spirv::AddressingModel::Logical,
-      spirv::MemoryModel::GLSL450, spirv::Capability::Shader,
-      spirv::Extension::SPV_KHR_storage_buffer_storage_class);
+      spirv::MemoryModel::GLSL450);
+
   // Move the region from the module op into the SPIR-V module.
   Region &spvModuleRegion = spvModule.body();
   rewriter.inlineRegionBefore(moduleOp.body(), spvModuleRegion,
                               spvModuleRegion.begin());
   // The spv.module build method adds a block with a terminator. Remove that
   // block. The terminator of the module op in the remaining block will be
   // legalized later.
   spvModuleRegion.back().erase();

mlir/lib/Dialect/SPIRV/SPIRVOps.cpp

 //===- SPIRVOps.cpp - MLIR SPIR-V operations ------------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines the operations in the SPIR-V dialect.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SPIRV/SPIRVOps.h"
 
+#include "mlir/Dialect/SPIRV/SPIRVAttributes.h"
 #include "mlir/Dialect/SPIRV/SPIRVDialect.h"
 #include "mlir/Dialect/SPIRV/SPIRVTypes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Function.h"
 #include "mlir/IR/FunctionImplementation.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/StandardTypes.h"
 #include "mlir/Interfaces/CallInterfaces.h"
getStrArrayAttrForEnumList(Builder &builder, ArrayRef<Ty> enumValues,
   SmallVector<StringRef, 1> enumValStrs;
   enumValStrs.reserve(enumValues.size());
   for (auto val : enumValues) {
     enumValStrs.emplace_back(stringifyFn(val));
   }
   return builder.getStrArrayAttr(enumValStrs);
 }
 
+/// Parses the next string attribute in `parser` as an enumerant of the given
+/// `EnumClass`.
 template <typename EnumClass>
 static ParseResult
-parseEnumAttribute(EnumClass &value, OpAsmParser &parser,
-                   StringRef attrName = spirv::attributeName<EnumClass>()) {
+parseEnumStrAttr(EnumClass &value, OpAsmParser &parser,
   Attribute attrVal;
   SmallVector<NamedAttribute, 1> attr;
   auto loc = parser.getCurrentLocation();
   if (parser.parseAttribute(attrVal, parser.getBuilder().getNoneType(),
                             attrName, attr)) {
     return failure();
   }
   if (!attrVal.isa<StringAttr>()) {
     return parser.emitError(loc, "expected ")
            << attrName << " attribute specified as string";
   }
   auto attrOptional =
       spirv::symbolizeEnum<EnumClass>()(attrVal.cast<StringAttr>().getValue());
   if (!attrOptional) {
     return parser.emitError(loc, "invalid ")
            << attrName << " attribute specification: " << attrVal;
   }
   value = attrOptional.getValue();
   return success();
 }
 
+/// Parses the next string attribute in `parser` as an enumerant of the given
+/// `EnumClass` and inserts the enumerant into `state` as an 32-bit integer
+/// attribute with the enum class's name as attribute name.
 template <typename EnumClass>
 static ParseResult
-parseEnumAttribute(EnumClass &value, OpAsmParser &parser, OperationState &state,
-                   StringRef attrName = spirv::attributeName<EnumClass>()) {
-  if (parseEnumAttribute(value, parser)) {
+parseEnumStrAttr(EnumClass &value, OpAsmParser &parser, OperationState &state,
+  if (parseEnumStrAttr(value, parser)) {
+    return failure();
+  }
+  state.addAttribute(attrName, parser.getBuilder().getI32IntegerAttr(
+                                   llvm::bit_cast<int32_t>(value)));
+  return success();
+}
+
+/// Parses the next keyword in `parser` as an enumerant of the given
+/// `EnumClass`.
+template <typename EnumClass>
+static ParseResult
+parseEnumKeywordAttr(EnumClass &value, OpAsmParser &parser,
+                     StringRef attrName = spirv::attributeName<EnumClass>()) {
+  StringRef keyword;
+  SmallVector<NamedAttribute, 1> attr;
+  auto loc = parser.getCurrentLocation();
+  if (parser.parseKeyword(&keyword))
+    return failure();
+  if (Optional<EnumClass> attr = spirv::symbolizeEnum<EnumClass>()(keyword)) {
+    value = attr.getValue();
+    return success();
+  }
+  return parser.emitError(loc, "invalid ")
+         << attrName << " attribute specification: " << keyword;
+}
+
+/// Parses the next keyword in `parser` as an enumerant of the given `EnumClass`
+/// and inserts the enumerant into `state` as an 32-bit integer attribute with
+/// the enum class's name as attribute name.
+template <typename EnumClass>
+static ParseResult
+parseEnumKeywordAttr(EnumClass &value, OpAsmParser &parser,
+                     OperationState &state,
+                     StringRef attrName = spirv::attributeName<EnumClass>()) {
+  if (parseEnumKeywordAttr(value, parser)) {
     return failure();
   }
   state.addAttribute(attrName, parser.getBuilder().getI32IntegerAttr(
                                    llvm::bit_cast<int32_t>(value)));
   return success();
 }
 
 static ParseResult parseMemoryAccessAttributes(OpAsmParser &parser,
                                                OperationState &state) {
   // Parse an optional list of attributes staring with '['
   if (parser.parseOptionalLSquare()) {
     // Nothing to do
     return success();
   }
 
   spirv::MemoryAccess memoryAccessAttr;
-  if (parseEnumAttribute(memoryAccessAttr, parser, state)) {
+  if (parseEnumStrAttr(memoryAccessAttr, parser, state)) {
     return failure();
   }
 
   if (spirv::bitEnumContains(memoryAccessAttr, spirv::MemoryAccess::Aligned)) {
     // Parse integer attribute for alignment.
     Attribute alignmentAttr;
     Type i32Type = parser.getBuilder().getIntegerType(32);
     if (parser.parseComma() ||

 static ParseResult parseAtomicUpdateOp(OpAsmParser &parser,
                                        OperationState &state, bool hasValue) {
   spirv::Scope scope;
   spirv::MemorySemantics memoryScope;
   SmallVector<OpAsmParser::OperandType, 2> operandInfo;
   OpAsmParser::OperandType ptrInfo, valueInfo;
   Type type;
   llvm::SMLoc loc;
-  if (parseEnumAttribute(scope, parser, state, kMemoryScopeAttrName) ||
-      parseEnumAttribute(memoryScope, parser, state, kSemanticsAttrName) ||
+  if (parseEnumStrAttr(scope, parser, state, kMemoryScopeAttrName) ||
+      parseEnumStrAttr(memoryScope, parser, state, kSemanticsAttrName) ||
       parser.parseOperandList(operandInfo, (hasValue ? 2 : 1)) ||
       parser.getCurrentLocation(&loc) || parser.parseColonType(type))
     return failure();
 
   auto ptrType = type.dyn_cast<spirv::PointerType>();
   if (!ptrType)
     return parser.emitError(loc, "expected pointer type");
 
static LogicalResult verifyAtomicUpdateOp(Operation *op) {
   return success();
 }
 
 static ParseResult parseGroupNonUniformArithmeticOp(OpAsmParser &parser,
                                                     OperationState &state) {
   spirv::Scope executionScope;
   spirv::GroupOperation groupOperation;
   OpAsmParser::OperandType valueInfo;
-  if (parseEnumAttribute(executionScope, parser, state,
-                         kExecutionScopeAttrName) ||
-      parseEnumAttribute(groupOperation, parser, state,
-                         kGroupOperationAttrName) ||
+  if (parseEnumStrAttr(executionScope, parser, state,
+      parseEnumStrAttr(groupOperation, parser, state,
       parser.parseOperand(valueInfo))
     return failure();
 
   Optional<OpAsmParser::OperandType> clusterSizeInfo;
   if (succeeded(parser.parseOptionalKeyword(kClusterSize))) {
     clusterSizeInfo = OpAsmParser::OperandType();
     if (parser.parseLParen() || parser.parseOperand(*clusterSizeInfo) ||
         parser.parseRParen())

 //===----------------------------------------------------------------------===//
 
 static ParseResult parseAtomicCompareExchangeWeakOp(OpAsmParser &parser,
                                                     OperationState &state) {
   spirv::Scope memoryScope;
   spirv::MemorySemantics equalSemantics, unequalSemantics;
   SmallVector<OpAsmParser::OperandType, 3> operandInfo;
   Type type;
-  if (parseEnumAttribute(memoryScope, parser, state, kMemoryScopeAttrName) ||
-      parseEnumAttribute(equalSemantics, parser, state,
-                         kEqualSemanticsAttrName) ||
-      parseEnumAttribute(unequalSemantics, parser, state,
-                         kUnequalSemanticsAttrName) ||
+  if (parseEnumStrAttr(memoryScope, parser, state, kMemoryScopeAttrName) ||
+      parseEnumStrAttr(equalSemantics, parser, state,
+      parseEnumStrAttr(unequalSemantics, parser, state,
       parser.parseOperandList(operandInfo, 3))
     return failure();
 
   auto loc = parser.getCurrentLocation();
   if (parser.parseColonType(type))
     return failure();
 
   auto ptrType = type.dyn_cast<spirv::PointerType>();

 static ParseResult parseEntryPointOp(OpAsmParser &parser,
                                      OperationState &state) {
   spirv::ExecutionModel execModel;
   SmallVector<OpAsmParser::OperandType, 0> identifiers;
   SmallVector<Type, 0> idTypes;
   SmallVector<Attribute, 4> interfaceVars;
 
   FlatSymbolRefAttr fn;
-  if (parseEnumAttribute(execModel, parser, state) ||
+  if (parseEnumStrAttr(execModel, parser, state) ||
       parser.parseAttribute(fn, Type(), kFnNameAttrName, state.attributes)) {
     return failure();
   }
 
   if (!parser.parseOptionalComma()) {
     // Parse the interface variables
     do {
       // The name of the interface variable attribute isnt important
build(builder, state, builder->getSymbolRefAttr(function),
         builder->getI32ArrayAttr(params));
 }
 
 static ParseResult parseExecutionModeOp(OpAsmParser &parser,
                                         OperationState &state) {
   spirv::ExecutionMode execMode;
   Attribute fn;
   if (parser.parseAttribute(fn, kFnNameAttrName, state.attributes) ||
-      parseEnumAttribute(execMode, parser, state)) {
+      parseEnumStrAttr(execMode, parser, state)) {
     return failure();
   }
 
   SmallVector<int32_t, 4> values;
   Type i32Type = parser.getBuilder().getIntegerType(32);
   while (!parser.parseOptionalComma()) {
     SmallVector<NamedAttribute, 1> attr;
     Attribute value;
if (impl::parseFunctionSignature(parser, /*allowVariadic=*/false, entryArgs,
                                    resultAttrs))
     return failure();
 
   auto fnType = builder.getFunctionType(argTypes, resultTypes);
   state.addAttribute(impl::getTypeAttrName(), TypeAttr::get(fnType));
 
   // Parse the optional function control keyword.
   spirv::FunctionControl fnControl;
-  if (parseEnumAttribute(fnControl, parser, state))
+  if (parseEnumStrAttr(fnControl, parser, state))
     return failure();
 
   // If additional attributes are present, parse them.
   if (parser.parseOptionalAttrDictWithKeyword(state.attributes))
     return failure();
 
   // Add the attributes to the function arguments.
   assert(argAttrs.size() == argTypes.size());
build(builder, state, ptrType.getPointeeType(), basePtr, memory_access,
         alignment);
 }
 
 static ParseResult parseLoadOp(OpAsmParser &parser, OperationState &state) {
   // Parse the storage class specification
   spirv::StorageClass storageClass;
   OpAsmParser::OperandType ptrInfo;
   Type elementType;
-  if (parseEnumAttribute(storageClass, parser) ||
-      parser.parseOperand(ptrInfo) ||
+  if (parseEnumStrAttr(storageClass, parser) || parser.parseOperand(ptrInfo) ||
       parseMemoryAccessAttributes(parser, state) ||
       parser.parseOptionalAttrDict(state.attributes) || parser.parseColon() ||
       parser.parseType(elementType)) {
     return failure();
   }
 
   auto ptrType = spirv::PointerType::get(elementType, storageClass);
   if (parser.resolveOperand(ptrInfo, ptrType, state.operands)) {

 //===----------------------------------------------------------------------===//
 // spv.module
 //===----------------------------------------------------------------------===//
 
 void spirv::ModuleOp::build(Builder *builder, OperationState &state) {
   ensureTerminator(*state.addRegion(), *builder, state.location);
 }
 
-// TODO(ravishankarm): This is only here for resolving some dependency outside
-// of mlir. Remove once it is done.
-void spirv::ModuleOp::build(Builder *builder, OperationState &state,
-                            IntegerAttr addressing_model,
-                            IntegerAttr memory_model) {
-  state.addAttribute("addressing_model", addressing_model);
-  state.addAttribute("memory_model", memory_model);
-  build(builder, state);
-}
-
 void spirv::ModuleOp::build(Builder *builder, OperationState &state,
                             spirv::AddressingModel addressing_model,
-                            spirv::MemoryModel memory_model,
-                            ArrayRef<spirv::Capability> capabilities,
-                            ArrayRef<spirv::Extension> extensions,
-                            ArrayAttr extended_instruction_sets) {
+                            spirv::MemoryModel memory_model) {
   state.addAttribute(
       "addressing_model",
       builder->getI32IntegerAttr(static_cast<int32_t>(addressing_model)));
   state.addAttribute("memory_model", builder->getI32IntegerAttr(
                                          static_cast<int32_t>(memory_model)));
-  if (!capabilities.empty())
-    state.addAttribute("capabilities",
-                       getStrArrayAttrForEnumList<spirv::Capability>(
-                           *builder, capabilities, spirv::stringifyCapability));
-  if (!extensions.empty())
-    state.addAttribute("extensions",
-                       getStrArrayAttrForEnumList<spirv::Extension>(
-                           *builder, extensions, spirv::stringifyExtension));
-  if (extended_instruction_sets)
-    state.addAttribute("extended_instruction_sets", extended_instruction_sets);
-  build(builder, state);
+  ensureTerminator(*state.addRegion(), *builder, state.location);
 }
 
 static ParseResult parseModuleOp(OpAsmParser &parser, OperationState &state) {
   Region *body = state.addRegion();
 
   // Parse attributes
   spirv::AddressingModel addrModel;
   spirv::MemoryModel memoryModel;
-  if (parseEnumAttribute(addrModel, parser, state) ||
-      parseEnumAttribute(memoryModel, parser, state)) {
+  if (parseEnumKeywordAttr(addrModel, parser, state) ||
+      parseEnumKeywordAttr(memoryModel, parser, state))
     return failure();
-  }
 
-  if (parser.parseRegion(*body, /*arguments=*/{}, /*argTypes=*/{}))
+  if (succeeded(parser.parseOptionalKeyword("requires"))) {
+    spirv::VerCapExtAttr vceTriple;
+    if (parser.parseAttribute(vceTriple,
+                              spirv::ModuleOp::getVCETripleAttrName(),
+                              state.attributes))
       return failure();
+  }
 
   if (parser.parseOptionalAttrDictWithKeyword(state.attributes))
     return failure();
 
+  if (parser.parseRegion(*body, /*arguments=*/{}, /*argTypes=*/{}))
+    return failure();
+
   spirv::ModuleOp::ensureTerminator(*body, parser.getBuilder(), state.location);
   return success();
 }
 
 static void print(spirv::ModuleOp moduleOp, OpAsmPrinter &printer) {
   printer << spirv::ModuleOp::getOperationName();
 
-  // Only print out addressing model and memory model in a nicer way if both
-  // presents. Otherwise, print them in the general form. This helps
-  // debugging ill-formed ModuleOp.
   SmallVector<StringRef, 2> elidedAttrs;
+
+  printer << " " << spirv::stringifyAddressingModel(moduleOp.addressing_model())
+          << " " << spirv::stringifyMemoryModel(moduleOp.memory_model());
   auto addressingModelAttrName = spirv::attributeName<spirv::AddressingModel>();
   auto memoryModelAttrName = spirv::attributeName<spirv::MemoryModel>();
-  if (moduleOp.getAttr(addressingModelAttrName) &&
-      moduleOp.getAttr(memoryModelAttrName)) {
-    printer << " \""
-            << spirv::stringifyAddressingModel(moduleOp.addressing_model())
-            << "\" \"" << spirv::stringifyMemoryModel(moduleOp.memory_model())
-            << '"';
-    elidedAttrs.assign({addressingModelAttrName, memoryModelAttrName});
+
+  if (Optional<spirv::VerCapExtAttr> triple = moduleOp.vce_triple()) {
+    printer << " requires " << *triple;
+    elidedAttrs.push_back(spirv::ModuleOp::getVCETripleAttrName());
   }
 
+  printer.printOptionalAttrDictWithKeyword(moduleOp.getAttrs(), elidedAttrs);
   printer.printRegion(moduleOp.body(), /*printEntryBlockArgs=*/false,
                       /*printBlockTerminators=*/false);
-  printer.printOptionalAttrDictWithKeyword(moduleOp.getAttrs(), elidedAttrs);
 }
 
 static LogicalResult verify(spirv::ModuleOp moduleOp) {
   auto &op = *moduleOp.getOperation();
   auto *dialect = op.getDialect();
-  auto &body = op.getRegion(0).front();
   DenseMap<std::pair<spirv::FuncOp, spirv::ExecutionModel>, spirv::EntryPointOp>
       entryPoints;
   SymbolTable table(moduleOp);
 
-  for (auto &op : body) {
+  for (auto &op : moduleOp.getBlock()) {
     if (op.getDialect() != dialect)
       return op.emitError("'spv.module' can only contain spv.* ops");
 
     // For EntryPoint op, check that the function and execution model is not
     // duplicated in EntryPointOps. Also verify that the interface specified
     // comes from globalVariables here to make this check cheaper.
     if (auto entryPointOp = dyn_cast<spirv::EntryPointOp>(op)) {
       auto funcOp = table.lookup<spirv::FuncOp>(entryPointOp.fn());
if (auto entryPointOp = dyn_cast<spirv::EntryPointOp>(op)) {
         for (auto &op : block) {
           if (op.getDialect() != dialect)
             return op.emitError(
                 "functions in 'spv.module' can only contain spv.* ops");
         }
     }
   }
 
-  // Verify capabilities. ODS already guarantees that we have an array of
-  // string attributes.
-  if (auto caps = moduleOp.getAttrOfType<ArrayAttr>("capabilities")) {
-    for (auto cap : caps.getValue()) {
-      auto capStr = cap.cast<StringAttr>().getValue();
-      if (!spirv::symbolizeCapability(capStr))
-        return moduleOp.emitOpError("uses unknown capability: ") << capStr;
-    }
-  }
-
-  // Verify extensions. ODS already guarantees that we have an array of
-  // string attributes.
-  if (auto exts = moduleOp.getAttrOfType<ArrayAttr>("extensions")) {
-    for (auto ext : exts.getValue()) {
-      auto extStr = ext.cast<StringAttr>().getValue();
-      if (!spirv::symbolizeExtension(extStr))
-        return moduleOp.emitOpError("uses unknown extension: ") << extStr;
-    }
-  }
-
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // spv._reference_of
 //===----------------------------------------------------------------------===//
 
 static LogicalResult verify(spirv::ReferenceOfOp referenceOfOp) {

 //===----------------------------------------------------------------------===//
 
 static ParseResult parseStoreOp(OpAsmParser &parser, OperationState &state) {
   // Parse the storage class specification
   spirv::StorageClass storageClass;
   SmallVector<OpAsmParser::OperandType, 2> operandInfo;
   auto loc = parser.getCurrentLocation();
   Type elementType;
-  if (parseEnumAttribute(storageClass, parser) ||
+  if (parseEnumStrAttr(storageClass, parser) ||
       parser.parseOperandList(operandInfo, 2) ||
       parseMemoryAccessAttributes(parser, state) || parser.parseColon() ||
       parser.parseType(elementType)) {
     return failure();
   }
 
   auto ptrType = spirv::PointerType::get(elementType, storageClass);
   if (parser.resolveOperands(operandInfo, {ptrType, elementType}, loc,

mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp

 //===- Deserializer.cpp - MLIR SPIR-V Deserialization ---------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines the SPIR-V binary to MLIR SPIR-V module deserialization.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SPIRV/Serialization.h"
 
+#include "mlir/Dialect/SPIRV/SPIRVAttributes.h"
 #include "mlir/Dialect/SPIRV/SPIRVBinaryUtils.h"
 #include "mlir/Dialect/SPIRV/SPIRVOps.h"
 #include "mlir/Dialect/SPIRV/SPIRVTypes.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Location.h"
 #include "mlir/Support/LogicalResult.h"
 #include "mlir/Support/StringExtras.h"
private:
 
   /// Processes SPIR-V module header in `binary`.
   LogicalResult processHeader();
 
   /// Processes the SPIR-V OpCapability with `operands` and updates bookkeeping
   /// in the deserializer.
   LogicalResult processCapability(ArrayRef<uint32_t> operands);
 
-  /// Attaches all collected capabilities to `module` as an attribute.
-  void attachCapabilities();
-
   /// Processes the SPIR-V OpExtension with `operands` and updates bookkeeping
   /// in the deserializer.
   LogicalResult processExtension(ArrayRef<uint32_t> words);
 
   /// Processes the SPIR-V OpExtInstImport with `operands` and updates
   /// bookkeeping in the deserializer.
   LogicalResult processExtInstImport(ArrayRef<uint32_t> words);
 
-  /// Attaches all collected extensions to `module` as an attribute.
-  void attachExtensions();
+  /// Attaches (version, capabilities, extensions) triple to `module` as an
+  /// attribute.
+  void attachVCETriple();
 
   /// Processes the SPIR-V OpMemoryModel with `operands` and updates `module`.
   LogicalResult processMemoryModel(ArrayRef<uint32_t> operands);
 
   /// Process SPIR-V OpName with `operands`.
   LogicalResult processName(ArrayRef<uint32_t> operands);
 
   /// Processes an OpDecorate instruction.
private:
   /// The current function under construction.
   Optional<spirv::FuncOp> curFunction;
 
   /// The current block under construction.
   Block *curBlock = nullptr;
 
   OpBuilder opBuilder;
 
+  spirv::Version version;
+
   /// The list of capabilities used by the module.
   llvm::SmallSetVector<spirv::Capability, 4> capabilities;
 
   /// The list of extensions used by the module.
-  llvm::SmallSetVector<StringRef, 2> extensions;
+  llvm::SmallSetVector<spirv::Extension, 2> extensions;
 
   // Result <id> to type mapping.
   DenseMap<uint32_t, Type> typeMap;
 
   // Result <id> to constant attribute and type mapping.
   ///
   /// In the SPIR-V binary format, all constants are placed in the module and
   /// shared by instructions at module level and in subsequent functions. But in
assert(curOffset == binarySize &&
          "deserializer should never index beyond the binary end");
 
   for (auto &deferred : deferredInstructions) {
     if (failed(processInstruction(deferred.first, deferred.second, false))) {
       return failure();
     }
   }
 
-  // Attaches the capabilities/extensions as an attribute to the module.
-  attachCapabilities();
-  attachExtensions();
+  attachVCETriple();
 
   LLVM_DEBUG(llvm::dbgs() << "+++ completed deserialization +++\n");
   return success();
 }
 
 Optional<spirv::ModuleOp> Deserializer::collect() { return module; }
 
 //===----------------------------------------------------------------------===//
 // Module structure
 //===----------------------------------------------------------------------===//
 
 spirv::ModuleOp Deserializer::createModuleOp() {
   Builder builder(context);
   OperationState state(unknownLoc, spirv::ModuleOp::getOperationName());
-  // TODO(antiagainst): use target environment to select the version
-  state.addAttribute("major_version", builder.getI32IntegerAttr(1));
-  state.addAttribute("minor_version", builder.getI32IntegerAttr(0));
   spirv::ModuleOp::build(&builder, state);
   return cast<spirv::ModuleOp>(Operation::create(state));
 }
 
 LogicalResult Deserializer::processHeader() {
   if (binary.size() < spirv::kHeaderWordCount)
     return emitError(unknownLoc,
                      "SPIR-V binary module must have a 5-word header");
 
   if (binary[0] != spirv::kMagicNumber)
     return emitError(unknownLoc, "incorrect magic number");
 
+  // Version number bytes: 0 | major number | minor number | 0
+  uint32_t majorVersion = (binary[1] << 8) >> 24;
+  uint32_t minorVersion = (binary[1] << 16) >> 24;
+  if (majorVersion == 1) {
+    switch (minorVersion) {
+#define MIN_VERSION_CASE(v)                                                    \
+  case v:                                                                      \
+    version = spirv::Version::V_1_##v;                                         \
+    break
+
+      MIN_VERSION_CASE(0);
+      MIN_VERSION_CASE(1);
+      MIN_VERSION_CASE(2);
+      MIN_VERSION_CASE(3);
+      MIN_VERSION_CASE(4);
+      MIN_VERSION_CASE(5);
+#undef MIN_VERSION_CASE
+    default:
+      return emitError(unknownLoc, "unspported SPIR-V minor version: ")
+             << minorVersion;
+    }
+  } else {
+    return emitError(unknownLoc, "unspported SPIR-V major version: ")
+           << majorVersion;
+  }
+
   // TODO(antiagainst): generator number, bound, schema
   curOffset = spirv::kHeaderWordCount;
   return success();
 }
 
 LogicalResult Deserializer::processCapability(ArrayRef<uint32_t> operands) {
   if (operands.size() != 1)
     return emitError(unknownLoc, "OpMemoryModel must have one parameter");
 
   auto cap = spirv::symbolizeCapability(operands[0]);
   if (!cap)
     return emitError(unknownLoc, "unknown capability: ") << operands[0];
 
   capabilities.insert(*cap);
   return success();
 }
 
-void Deserializer::attachCapabilities() {
-  if (capabilities.empty())
-    return;
-
-  SmallVector<StringRef, 2> caps;
-  caps.reserve(capabilities.size());
-
-  for (auto cap : capabilities) {
-    caps.push_back(spirv::stringifyCapability(cap));
-  }
-
-  module->setAttr("capabilities", opBuilder.getStrArrayAttr(caps));
-}
-
 LogicalResult Deserializer::processExtension(ArrayRef<uint32_t> words) {
   if (words.empty()) {
     return emitError(
         unknownLoc,
         "OpExtension must have a literal string for the extension name");
   }
 
   unsigned wordIndex = 0;
   StringRef extName = decodeStringLiteral(words, wordIndex);
-  if (wordIndex != words.size()) {
+  if (wordIndex != words.size())
     return emitError(unknownLoc,
                      "unexpected trailing words in OpExtension instruction");
-  }
+  auto ext = spirv::symbolizeExtension(extName);
+  if (!ext)
+    return emitError(unknownLoc, "unknown extension: ") << extName;
 
-  extensions.insert(extName);
+  extensions.insert(*ext);
   return success();
 }
 
 LogicalResult Deserializer::processExtInstImport(ArrayRef<uint32_t> words) {
   if (words.size() < 2) {
     return emitError(unknownLoc,
                      "OpExtInstImport must have a result <id> and a literal "
                      "string for the extended instruction set name");
   }
 
   unsigned wordIndex = 1;
   extendedInstSets[words[0]] = decodeStringLiteral(words, wordIndex);
   if (wordIndex != words.size()) {
     return emitError(unknownLoc,
                      "unexpected trailing words in OpExtInstImport");
   }
   return success();
 }
 
-void Deserializer::attachExtensions() {
-  if (extensions.empty())
-    return;
-
-  module->setAttr("extensions",
-                  opBuilder.getStrArrayAttr(extensions.getArrayRef()));
+void Deserializer::attachVCETriple() {
+  module->setAttr(spirv::ModuleOp::getVCETripleAttrName(),
+                  spirv::VerCapExtAttr::get(version, capabilities.getArrayRef(),
+                                            extensions.getArrayRef(), context));
 }
 
 LogicalResult Deserializer::processMemoryModel(ArrayRef<uint32_t> operands) {
   if (operands.size() != 2)
     return emitError(unknownLoc, "OpMemoryModel must have two operands");
 
   module->setAttr(
       "addressing_model",

mlir/lib/Dialect/SPIRV/Serialization/SPIRVBinaryUtils.cpp

 //===- SPIRVBinaryUtils.cpp - MLIR SPIR-V Binary Module Utilities ---------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines common utilities for SPIR-V binary module.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SPIRV/SPIRVBinaryUtils.h"
+#include "mlir/Dialect/SPIRV/SPIRVTypes.h"
 
 using namespace mlir;
 
 void spirv::appendModuleHeader(SmallVectorImpl<uint32_t> &header,
-                               uint32_t idBound) {
-  // The major and minor version number for the generated SPIR-V binary.
-  // TODO(antiagainst): use target environment to select the version
-  constexpr uint8_t kMajorVersion = 1;
-  constexpr uint8_t kMinorVersion = 0;
+                               spirv::Version version, uint32_t idBound) {
+  uint32_t majorVersion = 1;
+  uint32_t minorVersion = 0;
+  switch (version) {
+#define MIN_VERSION_CASE(v)                                                    \
+  case spirv::Version::V_1_##v:                                                \
+    minorVersion = v;                                                          \
+    break
+
+    MIN_VERSION_CASE(0);
+    MIN_VERSION_CASE(1);
+    MIN_VERSION_CASE(2);
+    MIN_VERSION_CASE(3);
+    MIN_VERSION_CASE(4);
+    MIN_VERSION_CASE(5);
+#undef MIN_VERSION_CASE
+  }
 
   // See "2.3. Physical Layout of a SPIR-V Module and Instruction" in the SPIR-V
   // spec for the definition of the binary module header.
   //
   // The first five words of a SPIR-V module must be:
   // +-------------------------------------------------------------------------+
   // | Magic number                                                            |
   // +-------------------------------------------------------------------------+
   // | Version number (bytes: 0 | major number | minor number | 0)             |
   // +-------------------------------------------------------------------------+
   // | Generator magic number                                                  |
   // +-------------------------------------------------------------------------+
   // | Bound (all result <id>s in the module guaranteed to be less than it)    |
   // +-------------------------------------------------------------------------+
   // | 0 (reserved for instruction schema)                                     |
   // +-------------------------------------------------------------------------+
   header.push_back(spirv::kMagicNumber);
-  header.push_back((kMajorVersion << 16) | (kMinorVersion << 8));
+  header.push_back((majorVersion << 16) | (minorVersion << 8));
   header.push_back(kGeneratorNumber);
   header.push_back(idBound); // <id> bound
   header.push_back(0);       // Schema (reserved word)
 }
 
 /// Returns the word-count-prefixed opcode for an SPIR-V instruction.
 uint32_t spirv::getPrefixedOpcode(uint32_t wordCount, spirv::Opcode opcode) {
   assert(((wordCount >> 16) == 0) && "word count out of range!");

mlir/lib/Dialect/SPIRV/Serialization/Serializer.cpp

 //===- Serializer.cpp - MLIR SPIR-V Serialization -------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines the MLIR SPIR-V module to SPIR-V binary serialization.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SPIRV/Serialization.h"
 
 #include "mlir/ADT/TypeSwitch.h"
+#include "mlir/Dialect/SPIRV/SPIRVAttributes.h"
 #include "mlir/Dialect/SPIRV/SPIRVBinaryUtils.h"
 #include "mlir/Dialect/SPIRV/SPIRVDialect.h"
 #include "mlir/Dialect/SPIRV/SPIRVOps.h"
 #include "mlir/Dialect/SPIRV/SPIRVTypes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/RegionGraphTraits.h"
 #include "mlir/Support/LogicalResult.h"
 #include "mlir/Support/StringExtras.h"
auto moduleSize = spirv::kHeaderWordCount + capabilities.size() +
                     extensions.size() + extendedSets.size() +
                     memoryModel.size() + entryPoints.size() +
                     executionModes.size() + decorations.size() +
                     typesGlobalValues.size() + functions.size();
 
   binary.clear();
   binary.reserve(moduleSize);
 
-  spirv::appendModuleHeader(binary, nextID);
+  spirv::appendModuleHeader(binary, module.vce_triple()->getVersion(), nextID);
   binary.append(capabilities.begin(), capabilities.end());
   binary.append(extensions.begin(), extensions.end());
   binary.append(extendedSets.begin(), extendedSets.end());
   binary.append(memoryModel.begin(), memoryModel.end());
   binary.append(entryPoints.begin(), entryPoints.end());
   binary.append(executionModes.begin(), executionModes.end());
   binary.append(names.begin(), names.end());
   binary.append(decorations.begin(), decorations.end());
uint32_t Serializer::getOrCreateFunctionID(StringRef fnName) {
   if (!funcID) {
     funcID = getNextID();
     funcIDMap[fnName] = funcID;
   }
   return funcID;
 }
 
 void Serializer::processCapability() {
-  auto caps = module.getAttrOfType<ArrayAttr>("capabilities");
-  if (!caps)
-    return;
-
-  for (auto cap : caps.getValue()) {
-    auto capStr = cap.cast<StringAttr>().getValue();
-    auto capVal = spirv::symbolizeCapability(capStr);
+  for (auto cap : module.vce_triple()->getCapabilities())
     encodeInstructionInto(capabilities, spirv::Opcode::OpCapability,
-                          {static_cast<uint32_t>(*capVal)});
-  }
+                          {static_cast<uint32_t>(cap)});
 }
 
 void Serializer::processExtension() {
-  auto exts = module.getAttrOfType<ArrayAttr>("extensions");
-  if (!exts)
-    return;
-
-  SmallVector<uint32_t, 16> extName;
-  for (auto ext : exts.getValue()) {
-    auto extStr = ext.cast<StringAttr>().getValue();
+  llvm::SmallVector<uint32_t, 16> extName;
+  for (spirv::Extension ext : module.vce_triple()->getExtensions()) {
     extName.clear();
-    spirv::encodeStringLiteralInto(extName, extStr);
+    spirv::encodeStringLiteralInto(extName, spirv::stringifyExtension(ext));
     encodeInstructionInto(extensions, spirv::Opcode::OpExtension, extName);
   }
 }
 
 void Serializer::processMemoryModel() {
   uint32_t mm = module.getAttrOfType<IntegerAttr>("memory_model").getInt();
   uint32_t am = module.getAttrOfType<IntegerAttr>("addressing_model").getInt();
 
LogicalResult Serializer::emitDecoration(uint32_t target,
   decorations.push_back(target);
   decorations.push_back(static_cast<uint32_t>(decoration));
   decorations.append(params.begin(), params.end());
   return success();
 }
 
 LogicalResult spirv::serialize(spirv::ModuleOp module,
                                SmallVectorImpl<uint32_t> &binary) {
+  if (!module.vce_triple().hasValue())
+    return module.emitError(
+        "module must have 'vce_triple' attribute to be serializeable");
+
   Serializer serializer(module);
 
   if (failed(serializer.serialize()))
     return failure();
 
   LLVM_DEBUG(serializer.printValueIDMap(llvm::dbgs()));
 
   serializer.collect(binary);
   return success();
 }

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

if (walkResult.wasInterrupted())
     return signalPassFailure();
 
   // TODO(antiagainst): verify that the deduced version is consistent with
   // SPIR-V ops' maximal version requirements.
 
   auto triple = spirv::VerCapExtAttr::get(
       deducedVersion, deducedCapabilities.getArrayRef(),
       deducedExtensions.getArrayRef(), &getContext());
-  module.setAttr("vce_triple", triple);
+  module.setAttr(spirv::ModuleOp::getVCETripleAttrName(), triple);
 }
 
 std::unique_ptr<OpPassBase<spirv::ModuleOp>>
 mlir::spirv::createUpdateVersionCapabilityExtensionPass() {
   return std::make_unique<UpdateVCEPass>();
 }
 
 static PassRegistration<UpdateVCEPass>
     pass("spirv-update-vce",
          "Deduce and attach minimal (version, capabilities, extensions) "
          "requirements to spv.module ops");

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 @builtin() {
     %c0 = constant 1 : index
     "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_id_x", kernel_module = @kernels} : (index, index, index, index, index, index) -> ()
     return
   }
 
-  // CHECK-LABEL:  spv.module "Logical" "GLSL450"
+  // CHECK-LABEL:  spv.module Logical GLSL450
   // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_x()
       attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv._address_of [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_id_y", kernel_module = @kernels} : (index, index, index, index, index, index) -> ()
     return
   }
 
-  // CHECK-LABEL:  spv.module "Logical" "GLSL450"
+  // CHECK-LABEL:  spv.module Logical GLSL450
   // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_y()
       attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv._address_of [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "y"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_id_z", kernel_module = @kernels} : (index, index, index, index, index, index) -> ()
     return
   }
 
-  // CHECK-LABEL:  spv.module "Logical" "GLSL450"
+  // CHECK-LABEL:  spv.module Logical GLSL450
   // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_z()
       attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv._address_of [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "z"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_size_x", kernel_module = @kernels} : (index, index, index, index, index, index) -> ()
     return
   }
 
-  // CHECK-LABEL:  spv.module "Logical" "GLSL450"
+  // CHECK-LABEL:  spv.module Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_x()
       attributes {gpu.kernel, 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"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_size_y", kernel_module = @kernels} : (index, index, index, index, index, index) -> ()
     return
   }
 
-  // CHECK-LABEL:  spv.module "Logical" "GLSL450"
+  // CHECK-LABEL:  spv.module Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_y()
       attributes {gpu.kernel, 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"() {dimension = "y"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_size_z", kernel_module = @kernels} : (index, index, index, index, index, index) -> ()
     return
   }
 
-  // CHECK-LABEL:  spv.module "Logical" "GLSL450"
+  // CHECK-LABEL:  spv.module Logical GLSL450
   gpu.module @kernels {
     gpu.func @builtin_workgroup_size_z()
       attributes {gpu.kernel, 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"() {dimension = "z"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_local_id_x", kernel_module = @kernels} : (index, index, index, index, index, index) -> ()
     return
   }
 
-  // CHECK-LABEL:  spv.module "Logical" "GLSL450"
+  // CHECK-LABEL:  spv.module Logical GLSL450
   // CHECK: spv.globalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
   gpu.module @kernels {
     gpu.func @builtin_local_id_x()
       attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv._address_of [[LOCALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.thread_id"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_num_workgroups_x", kernel_module = @kernels} : (index, index, index, index, index, index) -> ()
     return
   }
 
-  // CHECK-LABEL:  spv.module "Logical" "GLSL450"
+  // CHECK-LABEL:  spv.module Logical GLSL450
   // CHECK: spv.globalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
   gpu.module @kernels {
     gpu.func @builtin_num_workgroups_x()
       attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv._address_of [[NUMWORKGROUPS]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.grid_dim"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }

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

func @load_store(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>) {
     %c4 = constant 4 : index
     %1 = subi %c4, %c0_0 : index
     %c1_1 = constant 1 : index
     %c1_2 = constant 1 : index
     "gpu.launch_func"(%0, %c1_2, %c1_2, %1, %c1_2, %c1_2, %arg0, %arg1, %arg2, %c0, %c0_0, %c1, %c1_1) {kernel = "load_store_kernel", kernel_module = @kernels} : (index, index, index, index, index, index, memref<12x4xf32>, memref<12x4xf32>, memref<12x4xf32>, index, index, index, index) -> ()
     return
   }
 
-  // CHECK-LABEL: spv.module "Logical" "GLSL450"
+  // CHECK-LABEL: spv.module Logical GLSL450
   gpu.module @kernels {
     // CHECK-DAG: spv.globalVariable [[NUMWORKGROUPSVAR:@.*]] built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
     // CHECK-DAG: spv.globalVariable [[LOCALINVOCATIONIDVAR:@.*]] built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
     // CHECK-DAG: spv.globalVariable [[WORKGROUPIDVAR:@.*]] built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
     // CHECK-LABEL:    spv.func @load_store_kernel
     // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32 [4]> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32 [4]> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: [[ARG2:%.*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32 [4]> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 2 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}

mlir/test/Conversion/GPUToSPIRV/simple.mlir

 // RUN: mlir-opt -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:       spv.module Logical GLSL450 {
     // CHECK-LABEL: spv.func @basic_module_structure
     // CHECK-SAME: {{%.*}}: f32 {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: {{%.*}}: !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
     // CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}
     gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32>)
       attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
       // CHECK: spv.Return
       gpu.return
     }
-    // CHECK: attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}
   }
 
   func @main() {
     %0 = "op"() : () -> (f32)
     %1 = "op"() : () -> (memref<12xf32>)
     %cst = constant 1 : index
     "gpu.launch_func"(%cst, %cst, %cst, %cst, %cst, %cst, %0, %1) { kernel = "basic_module_structure", kernel_module = @kernels }
         : (index, index, index, index, index, index, f32, memref<12xf32>) -> ()

mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir

 // RUN: mlir-opt %s -convert-gpu-launch-to-vulkan-launch | FileCheck %s
 
 // CHECK: %[[resource:.*]] = alloc() : memref<12xf32>
 // CHECK: %[[index:.*]] = constant 1 : index
 // CHECK: call @vulkanLaunch(%[[index]], %[[index]], %[[index]], %[[index]], %[[index]], %[[index]], %[[resource]]) {spirv_blob = "{{.*}}", spirv_entry_point = "kernel"}
 
 module attributes {gpu.container_module} {
-  spv.module "Logical" "GLSL450" {
+  spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
     spv.globalVariable @kernel_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer>
     spv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
       %0 = spv._address_of @kernel_arg_0 : !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer>
       %2 = spv.constant 0 : i32
       %3 = spv._address_of @kernel_arg_0 : !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer>
       %4 = spv.AccessChain %0[%2, %2] : !spv.ptr<!spv.struct<!spv.array<12 x f32 [4]> [0]>, StorageBuffer>
       %5 = spv.Load "StorageBuffer" %4 : f32
       spv.Return
     }
     spv.EntryPoint "GLCompute" @kernel
     spv.ExecutionMode @kernel "LocalSize", 1, 1, 1
-  } attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}
+  }
   gpu.module @kernels {
     gpu.func @kernel(%arg0: memref<12xf32>) kernel {
       gpu.return
     }
   }
   func @foo() {
     %0 = alloc() : memref<12xf32>
     %c1 = constant 1 : index
     "gpu.launch_func"(%c1, %c1, %c1, %c1, %c1, %c1, %0) {kernel = "kernel", kernel_module = @kernels} : (index, index, index, index, index, index, memref<12xf32>) -> ()
     return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/arithmetic-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @fmul(%arg0 : f32, %arg1 : f32) "None" {
     // CHECK: {{%.*}}= spv.FMul {{%.*}}, {{%.*}} : f32
     %0 = spv.FMul %arg0, %arg1 : f32
     spv.Return
   }
   spv.func @fadd(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) "None" {
     // CHECK: {{%.*}} = spv.FAdd {{%.*}}, {{%.*}} : vector<4xf32>
     %0 = spv.FAdd %arg0, %arg1 : vector<4xf32>

mlir/test/Dialect/SPIRV/Serialization/array.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @array_stride(%arg0 : !spv.ptr<!spv.array<4x!spv.array<4xf32 [4]> [128]>, StorageBuffer>, %arg1 : i32, %arg2 : i32) "None" {
     // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.array<4 x !spv.array<4 x f32 [4]> [128]>, StorageBuffer>
     %2 = spv.AccessChain %arg0[%arg1, %arg2] : !spv.ptr<!spv.array<4x!spv.array<4xf32 [4]> [128]>, StorageBuffer>
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: spv.globalVariable {{@.*}} : !spv.ptr<!spv.rtarray<f32>, StorageBuffer>
   spv.globalVariable @var0 : !spv.ptr<!spv.rtarray<f32>, StorageBuffer>
   // CHECK: spv.globalVariable {{@.*}} : !spv.ptr<!spv.rtarray<vector<4xf16>>, Input>
   spv.globalVariable @var1 : !spv.ptr<!spv.rtarray<vector<4xf16>>, Input>
 }

mlir/test/Dialect/SPIRV/Serialization/atomic-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @atomic_compare_exchange_weak
   spv.func @atomic_compare_exchange_weak(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 "None" {
     // CHECK: spv.AtomicCompareExchangeWeak "Workgroup" "Release" "Acquire" %{{.*}}, %{{.*}}, %{{.*}} : !spv.ptr<i32, Workgroup>
     %0 = spv.AtomicCompareExchangeWeak "Workgroup" "Release" "Acquire" %ptr, %value, %comparator: !spv.ptr<i32, Workgroup>
     // CHECK: spv.AtomicAnd "Device" "None" %{{.*}}, %{{.*}} : !spv.ptr<i32, Workgroup>
     %1 = spv.AtomicAnd "Device" "None" %ptr, %value : !spv.ptr<i32, Workgroup>
     // CHECK: spv.AtomicIAdd "Workgroup" "Acquire" %{{.*}}, %{{.*}} : !spv.ptr<i32, Workgroup>
     %2 = spv.AtomicIAdd "Workgroup" "Acquire" %ptr, %value : !spv.ptr<i32, Workgroup>

mlir/test/Dialect/SPIRV/Serialization/barrier.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @memory_barrier_0() -> () "None" {
     // CHECK: spv.MemoryBarrier "Device", "Release|UniformMemory"
     spv.MemoryBarrier "Device", "Release|UniformMemory"
     spv.Return
   }
   spv.func @memory_barrier_1() -> () "None" {
     // CHECK: spv.MemoryBarrier "Subgroup", "AcquireRelease|SubgroupMemory"
     spv.MemoryBarrier "Subgroup", "AcquireRelease|SubgroupMemory"

mlir/test/Dialect/SPIRV/Serialization/bit-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @bitcount(%arg: i32) -> i32 "None" {
     // CHECK: spv.BitCount {{%.*}} : i32
     %0 = spv.BitCount %arg : i32
     spv.ReturnValue %0 : i32
   }
   spv.func @bit_field_insert(%base: vector<3xi32>, %insert: vector<3xi32>, %offset: i32, %count: i16) -> vector<3xi32> "None" {
     // CHECK: {{%.*}} = spv.BitFieldInsert {{%.*}}, {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi32>, i32, i16
     %0 = spv.BitFieldInsert %base, %insert, %offset, %count : vector<3xi32>, i32, i16

mlir/test/Dialect/SPIRV/Serialization/cast-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @bit_cast(%arg0 : f32) "None" {
     // CHECK: {{%.*}} = spv.Bitcast {{%.*}} : f32 to i32
     %0 = spv.Bitcast %arg0 : f32 to i32
     // CHECK: {{%.*}} = spv.Bitcast {{%.*}} : i32 to si32
     %1 = spv.Bitcast %0 : i32 to si32
     // CHECK: {{%.*}} = spv.Bitcast {{%.*}} : si32 to i32
     %2 = spv.Bitcast %1 : si32 to ui32
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @convert_f_to_s(%arg0 : f32) -> i32 "None" {
     // CHECK: {{%.*}} = spv.ConvertFToS {{%.*}} : f32 to i32
     %0 = spv.ConvertFToS %arg0 : f32 to i32
     spv.ReturnValue %0 : i32
   }
   spv.func @convert_f_to_u(%arg0 : f32) -> i32 "None" {
     // CHECK: {{%.*}} = spv.ConvertFToU {{%.*}} : f32 to i32
     %0 = spv.ConvertFToU %arg0 : f32 to i32

mlir/test/Dialect/SPIRV/Serialization/composite-op.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @composite_insert(%arg0 : !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>>, %arg1: !spv.array<4xf32>) -> !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>> "None" {
     // CHECK: spv.CompositeInsert {{%.*}}, {{%.*}}[1 : i32, 0 : i32] : !spv.array<4 x f32> into !spv.struct<f32, !spv.struct<!spv.array<4 x f32>, f32>>
     %0 = spv.CompositeInsert %arg1, %arg0[1 : i32, 0 : i32] : !spv.array<4xf32> into !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>>
     spv.ReturnValue %0: !spv.struct<f32, !spv.struct<!spv.array<4xf32>, f32>>
   }
   spv.func @composite_construct_vector(%arg0: f32, %arg1: f32, %arg2 : f32) -> vector<3xf32> "None" {
     // CHECK: spv.CompositeConstruct {{%.*}}, {{%.*}}, {{%.*}} : vector<3xf32>
     %0 = spv.CompositeConstruct %arg0, %arg1, %arg2 : vector<3xf32>
     spv.ReturnValue %0: vector<3xf32>
   }
 }

mlir/test/Dialect/SPIRV/Serialization/constant.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @bool_const
   spv.func @bool_const() -> () "None" {
     // CHECK: spv.constant true
     %0 = spv.constant true
     // CHECK: spv.constant false
     %1 = spv.constant false
 
     %2 = spv.Variable init(%0): !spv.ptr<i1, Function>

mlir/test/Dialect/SPIRV/Serialization/entry-point.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @noop() -> () "None" {
     spv.Return
   }
   // CHECK:      spv.EntryPoint "GLCompute" @noop
   // CHECK-NEXT: spv.ExecutionMode @noop "ContractionOff"
   spv.EntryPoint "GLCompute" @noop
   spv.ExecutionMode @noop "ContractionOff"
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK:       spv.globalVariable @var2 : !spv.ptr<f32, Input>
   // CHECK-NEXT:  spv.globalVariable @var3 : !spv.ptr<f32, Output>
   // CHECK-NEXT:  spv.func @noop({{%.*}}: !spv.ptr<f32, Input>, {{%.*}}: !spv.ptr<f32, Output>) "None"
   // CHECK:       spv.EntryPoint "GLCompute" @noop, @var2, @var3
   spv.globalVariable @var2 : !spv.ptr<f32, Input>
   spv.globalVariable @var3 : !spv.ptr<f32, Output>
   spv.func @noop(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @noop, @var2, @var3
   spv.ExecutionMode @noop "ContractionOff"
 }

mlir/test/Dialect/SPIRV/Serialization/execution-mode.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @foo
   // CHECK: spv.ExecutionMode @foo "LocalSizeHint", 3, 4, 5
   spv.ExecutionMode @foo "LocalSizeHint", 3, 4, 5
 }

mlir/test/Dialect/SPIRV/Serialization/function-call.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.globalVariable @var1 : !spv.ptr<!spv.array<4xf32>, Input>
   spv.func @fmain() -> i32 "None" {
     %0 = spv.constant 16 : i32
     %1 = spv._address_of @var1 : !spv.ptr<!spv.array<4xf32>, Input>
     // CHECK: {{%.*}} = spv.FunctionCall @f_0({{%.*}}) : (i32) -> i32
     %3 = spv.FunctionCall @f_0(%0) : (i32) -> i32
     // CHECK: spv.FunctionCall @f_1({{%.*}}, {{%.*}}) : (i32, !spv.ptr<!spv.array<4 x f32>, Input>) -> ()
     spv.FunctionCall @f_1(%3, %1) : (i32, !spv.ptr<!spv.array<4xf32>, Input>) ->  ()

mlir/test/Dialect/SPIRV/Serialization/global-variable.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 // CHECK:      spv.globalVariable @var0 bind(1, 0) : !spv.ptr<f32, Input>
 // CHECK-NEXT: spv.globalVariable @var1 bind(0, 1) : !spv.ptr<f32, Output>
 // CHECK-NEXT: spv.globalVariable @var2 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 // CHECK-NEXT: spv.globalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.globalVariable @var0 bind(1, 0) : !spv.ptr<f32, Input>
   spv.globalVariable @var1 bind(0, 1) : !spv.ptr<f32, Output>
   spv.globalVariable @var2 {built_in = "GlobalInvocationId"} : !spv.ptr<vector<3xi32>, Input>
   spv.globalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK:         spv.globalVariable @var1 : !spv.ptr<f32, Input>
   // CHECK-NEXT:    spv.globalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr<f32, Input>
   spv.globalVariable @var1 : !spv.ptr<f32, Input>
   spv.globalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr<f32, Input>
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
   spv.func @foo() "None" {
     // CHECK: %[[ADDR:.*]] = spv._address_of @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
     %0 = spv._address_of @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
     %1 = spv.constant 0: i32
     // CHECK: spv.AccessChain %[[ADDR]]
     %2 = spv.AccessChain %0[%1] : !spv.ptr<vector<3xi32>, Input>
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/glsl-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @fmul(%arg0 : f32, %arg1 : f32) "None" {
     // CHECK: {{%.*}} = spv.GLSL.Exp {{%.*}} : f32
     %0 = spv.GLSL.Exp %arg0 : f32
     // CHECK: {{%.*}} = spv.GLSL.FMax {{%.*}}, {{%.*}} : f32
     %1 = spv.GLSL.FMax %arg0, %arg1 : f32
     // CHECK: {{%.*}} = spv.GLSL.Sqrt {{%.*}} : f32
     %2 = spv.GLSL.Sqrt %arg0 : f32
     // CHECK: {{%.*}} = spv.GLSL.Cos {{%.*}} : f32
     %3 = spv.GLSL.Cos %arg0 : f32
     // CHECK: {{%.*}} = spv.GLSL.Sin {{%.*}} : f32
     %4 = spv.GLSL.Sin %arg0 : f32
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/group-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @subgroup_ballot
   spv.func @subgroup_ballot(%predicate: i1) -> vector<4xi32> "None" {
     // CHECK: %{{.*}} = spv.SubgroupBallotKHR %{{.*}}: vector<4xi32>
     %0 = spv.SubgroupBallotKHR %predicate: vector<4xi32>
     spv.ReturnValue %0: vector<4xi32>
   }
 }

mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @iequal_scalar(%arg0: i32, %arg1: i32)  "None" {
     // CHECK: {{.*}} = spv.IEqual {{.*}}, {{.*}} : i32
     %0 = spv.IEqual %arg0, %arg1 : i32
     spv.Return
   }
   spv.func @inotequal_vector(%arg0: vector<4xi32>, %arg1: vector<4xi32>) "None" {
     // CHECK: {{.*}} = spv.INotEqual {{.*}}, {{.*}} : vector<4xi32>
     %0 = spv.INotEqual %arg0, %arg1 : vector<4xi32>
spv.func @cmpf(%arg0 : f32, %arg1 : f32) "None" {
     // CHECK: spv.FUnordNotEqual
     %12 = spv.FUnordNotEqual %arg0, %arg1 : f32
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.specConstant @condition_scalar = true
   spv.func @select() -> () "None" {
     %0 = spv.constant 4.0 : f32
     %1 = spv.constant 5.0 : f32
     %2 = spv._reference_of @condition_scalar : i1
     // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : i1, f32
     %3 = spv.Select %2, %0, %1 : i1, f32
     %4 = spv.constant dense<[2.0, 3.0, 4.0, 5.0]> : vector<4xf32>

mlir/test/Dialect/SPIRV/Serialization/loop.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 // Single loop
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // for (int i = 0; i < count; ++i) {}
   spv.func @loop(%count : i32) -> () "None" {
     %zero = spv.constant 0: i32
     %one = spv.constant 1: i32
     %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
 // CHECK:        spv.Branch ^bb1
 // CHECK-NEXT: ^bb1:
// CHECK-NEXT:     spv._merge
     }
     spv.Return
   }
 
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
-} attributes {
-  capabilities = ["Shader"]
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.globalVariable @GV1 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
   spv.globalVariable @GV2 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
   spv.func @loop_kernel() "None" {
     %0 = spv._address_of @GV1 : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
     %1 = spv.constant 0 : i32
     %2 = spv.AccessChain %0[%1] : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
     %3 = spv._address_of @GV2 : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
     %5 = spv.AccessChain %3[%1] : !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
// CHECK-NEXT:   ^bb3:
     ^merge:
 // CHECK-NEXT:     spv._merge
       spv._merge
     }
     spv.Return
   }
   spv.EntryPoint "GLCompute" @loop_kernel
   spv.ExecutionMode @loop_kernel "LocalSize", 1, 1, 1
-} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}
+}
 
 // -----
 
 // Nested loop
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // for (int i = 0; i < count; ++i) {
   //   for (int j = 0; j < count; ++j) { }
   // }
   spv.func @loop(%count : i32) -> () "None" {
     %zero = spv.constant 0: i32
     %one = spv.constant 1: i32
     %ivar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
     %jvar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
// CHECK-NEXT:     spv._merge
     } // end outer loop
     spv.Return
   }
 
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
-} attributes {
-  capabilities = ["Shader"]
 }
 

mlir/test/Dialect/SPIRV/Serialization/memory-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 // CHECK:           spv.func {{@.*}}([[ARG1:%.*]]: !spv.ptr<f32, Input>, [[ARG2:%.*]]: !spv.ptr<f32, Output>) "None" {
 // CHECK-NEXT:        [[VALUE:%.*]] = spv.Load "Input" [[ARG1]] : f32
 // CHECK-NEXT:        spv.Store "Output" [[ARG2]], [[VALUE]] : f32
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @load_store(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) "None" {
     %1 = spv.Load "Input" %arg0 : f32
     spv.Store "Output" %arg1, %1 : f32
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @access_chain(%arg0 : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>, %arg1 : i32, %arg2 : i32) "None" {
     // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.array<4 x !spv.array<4 x f32>>, Function>
     // CHECK-NEXT: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.array<4 x !spv.array<4 x f32>>, Function>
     %1 = spv.AccessChain %arg0[%arg1] : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>
     %2 = spv.AccessChain %arg0[%arg1, %arg2] : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @load_store_zero_rank_float(%arg0: !spv.ptr<!spv.struct<!spv.array<1 x f32 [4]> [0]>, StorageBuffer>, %arg1: !spv.ptr<!spv.struct<!spv.array<1 x f32 [4]> [0]>, StorageBuffer>) "None" {
     // CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<!spv.array<1 x f32 [4]> [0]>
     // CHECK-NEXT: [[VAL:%.*]] = spv.Load "StorageBuffer" [[LOAD_PTR]] : f32
     %0 = spv.constant 0 : i32
     %1 = spv.AccessChain %arg0[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x f32 [4]> [0]>, StorageBuffer>
     %2 = spv.Load "StorageBuffer" %1 : f32
 
     // CHECK: [[STORE_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<!spv.array<1 x f32 [4]> [0]>

mlir/test/Dialect/SPIRV/Serialization/module.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-// CHECK:      spv.module "Logical" "GLSL450" {
+// CHECK:      spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 // CHECK-NEXT:   spv.func @foo() "None" {
 // CHECK-NEXT:     spv.Return
 // CHECK-NEXT:   }
-// CHECK-NEXT: } attributes {major_version = 1 : i32, minor_version = 0 : i32}
+// CHECK-NEXT: }
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "None" {
      spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
-} attributes {
-  // CHECK: capabilities = ["Shader", "Float16"]
-  capabilities = ["Shader", "Float16"]
+// CHECK: v1.5
+spv.module Logical GLSL450 requires #spv.vce<v1.5, [Shader], []> {
+}
+
+// -----
+
+// CHECK: [Shader, Float16]
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader, Float16], []> {
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
-} attributes {
-  // CHECK: extensions = ["SPV_KHR_float_controls", "SPV_KHR_subgroup_vote"]
-  extensions = ["SPV_KHR_float_controls", "SPV_KHR_subgroup_vote"]
+// CHECK: [SPV_KHR_float_controls, SPV_KHR_subgroup_vote]
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_float_controls, SPV_KHR_subgroup_vote]> {
 }
 

mlir/test/Dialect/SPIRV/Serialization/non-uniform-ops.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @group_non_uniform_ballot
   spv.func @group_non_uniform_ballot(%predicate: i1) -> vector<4xi32> "None" {
     // CHECK: %{{.*}} = spv.GroupNonUniformBallot "Workgroup" %{{.*}}: vector<4xi32>
   %0 = spv.GroupNonUniformBallot "Workgroup" %predicate : vector<4xi32>
     spv.ReturnValue %0: vector<4xi32>
   }
 
   // CHECK-LABEL: @group_non_uniform_elect

mlir/test/Dialect/SPIRV/Serialization/phi.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 // Test branch with one block argument
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "None" {
 // CHECK:        %[[CST:.*]] = spv.constant 0
     %zero = spv.constant 0 : i32
 // CHECK-NEXT:   spv.Branch ^bb1(%[[CST]] : i32)
     spv.Branch ^bb1(%zero : i32)
 // CHECK-NEXT: ^bb1(%{{.*}}: i32):
   ^bb1(%arg0: i32):
    spv.Return
   }
 
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
-} attributes {
-  capabilities = ["Shader"]
 }
 
 // -----
 
 // Test branch with multiple block arguments
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "None" {
 // CHECK:        %[[ZERO:.*]] = spv.constant 0
     %zero = spv.constant 0 : i32
 // CHECK-NEXT:   %[[ONE:.*]] = spv.constant 1
     %one = spv.constant 1.0 : f32
 // CHECK-NEXT:   spv.Branch ^bb1(%[[ZERO]], %[[ONE]] : i32, f32)
     spv.Branch ^bb1(%zero, %one : i32, f32)
 
 // CHECK-NEXT: ^bb1(%{{.*}}: i32, %{{.*}}: f32):     // pred: ^bb0
   ^bb1(%arg0: i32, %arg1: f32):
    spv.Return
   }
 
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
-} attributes {
-  capabilities = ["Shader"]
 }
 
 // -----
 
 // Test using block arguments within branch
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "None" {
 // CHECK:        %[[CST0:.*]] = spv.constant 0
     %zero = spv.constant 0 : i32
 // CHECK-NEXT:   spv.Branch ^bb1(%[[CST0]] : i32)
     spv.Branch ^bb1(%zero : i32)
 
 // CHECK-NEXT: ^bb1(%[[ARG:.*]]: i32):
   ^bb1(%arg0: i32):
 // CHECK-NEXT:   %[[ADD:.*]] = spv.IAdd %[[ARG]], %[[ARG]] : i32
     %0 = spv.IAdd %arg0, %arg0 : i32
 // CHECK-NEXT:   %[[CST1:.*]] = spv.constant 0
 // CHECK-NEXT:   spv.Branch ^bb2(%[[CST1]], %[[ADD]] : i32, i32)
     spv.Branch ^bb2(%zero, %0 : i32, i32)
 
 // CHECK-NEXT: ^bb2(%{{.*}}: i32, %{{.*}}: i32):
   ^bb2(%arg1: i32, %arg2: i32):
    spv.Return
   }
 
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
-} attributes {
-  capabilities = ["Shader"]
 }
 
 // -----
 
 // Test block not following domination order
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "None" {
 // CHECK:        spv.Branch ^bb1
     spv.Branch ^bb1
 
 // CHECK-NEXT: ^bb1:
 // CHECK-NEXT:   %[[ZERO:.*]] = spv.constant 0
 // CHECK-NEXT:   %[[ONE:.*]] = spv.constant 1
 // CHECK-NEXT:   spv.Branch ^bb2(%[[ZERO]], %[[ONE]] : i32, f32)
^bb1:
     %one = spv.constant 1.0 : f32
     spv.Branch ^bb2(%zero, %one : i32, f32)
   }
 
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
-} attributes {
-  capabilities = ["Shader"]
 }
 
 // -----
 
 // Test multiple predecessors
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "None" {
     %var = spv.Variable : !spv.ptr<i32, Function>
 
 // CHECK:      spv.selection
     spv.selection {
       %true = spv.constant true
 // CHECK:        spv.BranchConditional %{{.*}}, ^bb1, ^bb2
       spv.BranchConditional %true, ^true, ^false
// CHECK-NEXT:   spv._merge
     }
     spv.Return
   }
 
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
-} attributes {
-  capabilities = ["Shader"]
 }
 
 // -----
 
 // Test nested loops with block arguments
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.globalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
   spv.globalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
   spv.func @fmul_kernel() "None" {
     %3 = spv.constant 12 : i32
     %4 = spv.constant 32 : i32
     %5 = spv.constant 4 : i32
     %6 = spv._address_of @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
     %7 = spv.Load "Input" %6 : vector<3xi32>

 // CHECK:     spv._merge
       spv._merge
     }
     spv.Return
   }
 
   spv.EntryPoint "GLCompute" @fmul_kernel, @__builtin_var_WorkgroupId__, @__builtin_var_NumWorkgroups__
   spv.ExecutionMode @fmul_kernel "LocalSize", 32, 1, 1
-} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}
+}

mlir/test/Dialect/SPIRV/Serialization/selection.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 // Selection with both then and else branches
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @selection(%cond: i1) -> () "None" {
 // CHECK:        spv.Branch ^bb1
 // CHECK-NEXT: ^bb1:
     %zero = spv.constant 0: i32
     %one = spv.constant 1: i32
     %two = spv.constant 2: i32
     %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
// CHECK-NEXT:     spv._merge
     spv.Return
   }
 
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
   spv.ExecutionMode @main "LocalSize", 1, 1, 1
-} attributes {
-  capabilities = ["Shader"]
 }
 
 // -----
 
 // Selection with only then branch
 // Selection in function entry block
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 // CHECK:      spv.func @selection(%[[ARG:.*]]: i1
   spv.func @selection(%cond: i1) -> (i32) "None" {
 // CHECK:        spv.Branch ^bb1
 // CHECK-NEXT: ^bb1:
 // CHECK-NEXT:   spv.selection
     spv.selection {
 // CHECK-NEXT: spv.BranchConditional %[[ARG]], ^bb1, ^bb2
       spv.BranchConditional %cond, ^then, ^merge
// CHECK-NEXT:     spv._merge
     spv.ReturnValue  %one : i32
   }
 
   spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
   spv.ExecutionMode @main "LocalSize", 1, 1, 1
-} attributes {
-  capabilities = ["Shader"]
 }
 

mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: spv.specConstant @sc_true = true
   spv.specConstant @sc_true = true
   // CHECK: spv.specConstant @sc_false spec_id(1) = false
   spv.specConstant @sc_false spec_id(1) = false
 
   // CHECK: spv.specConstant @sc_int = -5 : i32
   spv.specConstant @sc_int = -5 : i32
 

mlir/test/Dialect/SPIRV/Serialization/struct.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Input>
   spv.globalVariable @var0 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Input>
 
   // CHECK: !spv.ptr<!spv.struct<f32 [0], !spv.struct<f32 [0], !spv.array<16 x f32 [4]> [4]> [4]>, Input>
   spv.globalVariable @var1 bind(0, 2) : !spv.ptr<!spv.struct<f32 [0], !spv.struct<f32 [0], !spv.array<16 x f32 [4]> [4]> [4]>, Input>
 
   // CHECK: !spv.ptr<!spv.struct<f32 [0], i32 [4], f64 [8], i64 [16], f32 [24], i32 [30], f32 [34], i32 [38]>, StorageBuffer>
   spv.globalVariable @var2 : !spv.ptr<!spv.struct<f32 [0], i32 [4], f64 [8], i64 [16], f32 [24], i32 [30], f32 [34], i32 [38]>, StorageBuffer>

mlir/test/Dialect/SPIRV/Serialization/terminator.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @ret
   spv.func @ret() -> () "None" {
     // CHECK: spv.Return
     spv.Return
   }
 
   // CHECK-LABEL: @ret_val
   spv.func @ret_val() -> (i32) "None" {

mlir/test/Dialect/SPIRV/Serialization/undef.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "None" {
     // CHECK: {{%.*}} = spv.undef : f32
     // CHECK-NEXT: {{%.*}} = spv.undef : f32
     %0 = spv.undef : f32
     %1 = spv.undef : f32
     %2 = spv.FAdd %0, %1 : f32
     // CHECK: {{%.*}} = spv.undef : vector<4xi32>
     %3 = spv.undef : vector<4xi32>
     %4 = spv.CompositeExtract %3[1 : i32] : vector<4xi32>
     // CHECK: {{%.*}} = spv.undef : !spv.array<4 x !spv.array<4 x i32>>
     %5 = spv.undef : !spv.array<4x!spv.array<4xi32>>
     %6 = spv.CompositeExtract %5[1 : i32, 2 : i32] : !spv.array<4x!spv.array<4xi32>>
     // CHECK: {{%.*}} = spv.undef : !spv.ptr<!spv.struct<f32>, StorageBuffer>
     %7 = spv.undef : !spv.ptr<!spv.struct<f32>, StorageBuffer>
     %8 = spv.constant 0 : i32
     %9 = spv.AccessChain %7[%8] : !spv.ptr<!spv.struct<f32>, StorageBuffer>
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: spv.func {{@.*}}
   spv.func @ignore_unused_undef() -> () "None" {
     // CHECK-NEXT: spv.Return
     %0 = spv.undef : f32
     spv.Return
   }
 }

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

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 // CHECK-LABEL: spv.module
-spv.module "Logical" "GLSL450" {
+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>
   // CHECK-DAG: spv.globalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
   spv.globalVariable @__builtin_var_LocalInvocationId__ built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
   // CHECK-DAG: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   spv.globalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
attributes  {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
     %42 = spv.FAdd %39, %41 : f32
     // CHECK: spv.AccessChain [[ARG2]]
     %43 = spv.AccessChain %arg2[%c0, %36, %37] : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32>>>, StorageBuffer>
     spv.Store "StorageBuffer" %43, %42 : f32
     spv.Return
   }
   // CHECK: spv.EntryPoint "GLCompute" [[FN]], [[WORKGROUPID]], [[LOCALINVOCATIONID]], [[NUMWORKGROUPS]], [[WORKGROUPSIZE]]
   // CHECK-NEXT: spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
-} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}
+}

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

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 // CHECK-LABEL: spv.module
-spv.module "Logical" "GLSL450" {
+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 [4]> [0]>, StorageBuffer>
   // CHECK:    spv.func [[FN:@.*]]()
   spv.func @kernel(%arg0: f32
                 {spv.interface_var_abi = {binding = 0 : i32,
                                           descriptor_set = 0 : i32,
                                           storage_class = 12 : i32}},
                  %arg1: !spv.ptr<!spv.struct<!spv.array<12 x f32>>, StorageBuffer>
                  {spv.interface_var_abi = {binding = 1 : i32,
                                            descriptor_set = 0 : i32,
                                            storage_class = 12 : i32}}) "None"
   attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
     // CHECK: [[ARG1:%.*]] = spv._address_of [[VAR1]]
     // CHECK: [[ADDRESSARG0:%.*]] = spv._address_of [[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
-} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}
+}

mlir/test/Dialect/SPIRV/Transforms/inlining.mlir

 // RUN: mlir-opt %s -split-input-file -pass-pipeline='spv.module(inline)' -mlir-disable-inline-simplify | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @callee() "None" {
     spv.Return
   }
 
   // CHECK-LABEL: @calling_single_block_ret_func
   spv.func @calling_single_block_ret_func() "None" {
     // CHECK-NEXT: spv.Return
     spv.FunctionCall @callee() : () -> ()
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @callee() -> i32 "None" {
     %0 = spv.constant 42 : i32
     spv.ReturnValue %0 : i32
   }
 
   // CHECK-LABEL: @calling_single_block_retval_func
   spv.func @calling_single_block_retval_func() -> i32 "None" {
     // CHECK-NEXT: %[[CST:.*]] = spv.constant 42
     %0 = spv.FunctionCall @callee() : () -> (i32)
     // CHECK-NEXT: spv.ReturnValue %[[CST]]
     spv.ReturnValue %0 : i32
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.globalVariable @data bind(0, 0) : !spv.ptr<!spv.struct<!spv.rtarray<i32> [0]>, StorageBuffer>
   spv.func @callee() "None" {
     %0 = spv._address_of @data : !spv.ptr<!spv.struct<!spv.rtarray<i32> [0]>, StorageBuffer>
     %1 = spv.constant 0: i32
     %2 = spv.AccessChain %0[%1, %1] : !spv.ptr<!spv.struct<!spv.rtarray<i32> [0]>, StorageBuffer>
     spv.Branch ^next
 
   ^next:
spv.func @calling_multi_block_ret_func() "None" {
     spv.Return
   }
 }
 
 // TODO: calling_multi_block_retval_func
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @callee(%cond : i1) -> () "None" {
     spv.selection {
       spv.BranchConditional %cond, ^then, ^merge
     ^then:
       spv.Return
     ^merge:
       spv._merge
     }
     spv.Return
   }
 
   // CHECK-LABEL: @calling_selection_ret_func
   spv.func @calling_selection_ret_func() "None" {
     %0 = spv.constant true
     // CHECK: spv.FunctionCall
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @callee(%cond : i1) -> () "None" {
     spv.selection {
       spv.BranchConditional %cond, ^then, ^merge
     ^then:
       spv.Branch ^merge
     ^merge:
       spv._merge
     }
spv.func @calling_selection_no_ret_func() "None" {
     // CHECK-NEXT:   spv._merge
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @callee(%cond : i1) -> () "None" {
     spv.loop {
       spv.Branch ^header
     ^header:
       spv.BranchConditional %cond, ^body, ^merge
     ^body:
       spv.Return
     ^continue:
spv.func @calling_loop_ret_func() "None" {
     // CHECK: spv.FunctionCall
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @callee(%cond : i1) -> () "None" {
     spv.loop {
       spv.Branch ^header
     ^header:
       spv.BranchConditional %cond, ^body, ^merge
     ^body:
       spv.Branch ^continue
     ^continue:
spv.func @calling_loop_no_ret_func() "None" {
     // CHECK-NEXT:   spv._merge
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.globalVariable @arg_0 bind(0, 0) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
   spv.globalVariable @arg_1 bind(0, 1) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
 
   // CHECK: @inline_into_selection_region
   spv.func @inline_into_selection_region() "None" {
     %1 = spv.constant 0 : i32
     // CHECK-DAG: [[ADDRESS_ARG0:%.*]] = spv._address_of @arg_0
     // CHECK-DAG: [[ADDRESS_ARG1:%.*]] = spv._address_of @arg_1
spv.func @inline_into_selection_region() "None" {
     spv.Return
   }
   spv.func @atomic_add(%arg0: i32, %arg1: !spv.ptr<i32, StorageBuffer>) "None" {
     %0 = spv.AtomicIAdd "Device" "AcquireRelease" %arg1, %arg0 : !spv.ptr<i32, StorageBuffer>
     spv.Return
   }
   spv.EntryPoint "GLCompute" @inline_into_selection_region
   spv.ExecutionMode @inline_into_selection_region "LocalSize", 32, 1, 1
-} attributes {capabilities = ["Shader"], extensions = ["SPV_KHR_storage_buffer_storage_class"]}
+}
 
 // TODO: Add tests for inlining structured control flow into
 // structured control flow.

mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir

 // RUN: mlir-opt -decorate-spirv-composite-type-layout -split-input-file -verify-diagnostics %s -o - | FileCheck %s
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // CHECK: spv.globalVariable @var0 bind(0, 1) : !spv.ptr<!spv.struct<i32 [0], !spv.struct<f32 [0], i32 [4]> [4], f32 [12]>, Uniform>
   spv.globalVariable @var0 bind(0,1) : !spv.ptr<!spv.struct<i32, !spv.struct<f32, i32>, f32>, Uniform>
 
   // CHECK: spv.globalVariable @var1 bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<64 x i32 [4]> [0], f32 [256]>, StorageBuffer>
   spv.globalVariable @var1 bind(0,2) : !spv.ptr<!spv.struct<!spv.array<64xi32>, f32>, StorageBuffer>
 
   // CHECK: spv.globalVariable @var2 bind(1, 0) : !spv.ptr<!spv.struct<!spv.struct<!spv.array<64 x i32 [4]> [0], f32 [256]> [0], i32 [260]>, StorageBuffer>
   spv.globalVariable @var2 bind(1,0) : !spv.ptr<!spv.struct<!spv.struct<!spv.array<64xi32>, f32>, i32>, StorageBuffer>
spv.func @kernel() -> () "None" {
     // CHECK:  {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.struct<i32 [0], !spv.struct<f32 [0], i32 [4]> [4], f32 [12]>, Uniform>
     %1 = spv.AccessChain %0[%c0] : !spv.ptr<!spv.struct<i32, !spv.struct<f32, i32>, f32>, Uniform>
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // CHECK: spv.globalVariable @var0 : !spv.ptr<!spv.struct<!spv.struct<!spv.struct<!spv.struct<!spv.struct<i1 [0], i1 [1], f64 [8]> [0], i1 [16]> [0], i1 [24]> [0], i1 [32]> [0], i1 [40]>, Uniform>
   spv.globalVariable @var0 : !spv.ptr<!spv.struct<!spv.struct<!spv.struct<!spv.struct<!spv.struct<i1, i1, f64>, i1>, i1>, i1>, i1>, Uniform>
 
   // CHECK: spv.globalVariable @var1 : !spv.ptr<!spv.struct<!spv.struct<i16 [0], !spv.struct<i1 [0], f64 [8]> [8], f32 [24]> [0], f32 [32]>, Uniform>
   spv.globalVariable @var1 : !spv.ptr<!spv.struct<!spv.struct<i16, !spv.struct<i1, f64>, f32>, f32>, Uniform>
 
   // CHECK: spv.globalVariable @var2 : !spv.ptr<!spv.struct<!spv.struct<i16 [0], !spv.struct<i1 [0], !spv.array<16 x !spv.array<16 x i64 [8]> [128]> [8]> [8], f32 [2064]> [0], f32 [2072]>, Uniform>
   spv.globalVariable @var2 : !spv.ptr<!spv.struct<!spv.struct<i16, !spv.struct<i1, !spv.array<16x!spv.array<16xi64>>>, f32>, f32>, Uniform>
spv.module Logical GLSL450 {
   spv.globalVariable @var6 : !spv.ptr<!spv.struct<i1, !spv.struct<i64, i32, i16, i8, i1>, i1>, Uniform>
 
   // CHECK: spv.globalVariable @var7 : !spv.ptr<!spv.struct<i1 [0], !spv.struct<!spv.struct<i1 [0], i64 [8]> [0], i1 [16]> [8], i1 [32]>, Uniform>
   spv.globalVariable @var7 : !spv.ptr<!spv.struct<i1, !spv.struct<!spv.struct<i1, i64>, i1>, i1>, Uniform>
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // CHECK: spv.globalVariable @var0 : !spv.ptr<!spv.struct<vector<2xi32> [0], f32 [8]>, StorageBuffer>
   spv.globalVariable @var0 : !spv.ptr<!spv.struct<vector<2xi32>, f32>, StorageBuffer>
 
   // CHECK: spv.globalVariable @var1 : !spv.ptr<!spv.struct<vector<3xi32> [0], f32 [12]>, StorageBuffer>
   spv.globalVariable @var1 : !spv.ptr<!spv.struct<vector<3xi32>, f32>, StorageBuffer>
 
   // CHECK: spv.globalVariable @var2 : !spv.ptr<!spv.struct<vector<4xi32> [0], f32 [16]>, StorageBuffer>
   spv.globalVariable @var2 : !spv.ptr<!spv.struct<vector<4xi32>, f32>, StorageBuffer>
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // CHECK: spv.globalVariable @emptyStructAsMember : !spv.ptr<!spv.struct<!spv.struct<> [0]>, StorageBuffer>
   spv.globalVariable @emptyStructAsMember : !spv.ptr<!spv.struct<!spv.struct<>>, StorageBuffer>
 
   // CHECK: spv.globalVariable @arrayType : !spv.ptr<!spv.array<4 x !spv.array<4 x f32>>, StorageBuffer>
   spv.globalVariable @arrayType : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, StorageBuffer>
 
   // CHECK: spv.globalVariable @InputStorage : !spv.ptr<!spv.struct<!spv.array<256 x f32>>, Input>
   spv.globalVariable @InputStorage : !spv.ptr<!spv.struct<!spv.array<256xf32>>, Input>
 
   // CHECK: spv.globalVariable @customLayout : !spv.ptr<!spv.struct<f32 [256], i32 [512]>, Uniform>
   spv.globalVariable @customLayout : !spv.ptr<!spv.struct<f32 [256], i32 [512]>, Uniform>
 
   // CHECK:  spv.globalVariable @emptyStruct : !spv.ptr<!spv.struct<>, Uniform>
   spv.globalVariable @emptyStruct : !spv.ptr<!spv.struct<>, Uniform>
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // CHECK: spv.globalVariable @var0 : !spv.ptr<!spv.struct<i32 [0]>, PushConstant>
   spv.globalVariable @var0 : !spv.ptr<!spv.struct<i32>, PushConstant>
   // CHECK: spv.globalVariable @var1 : !spv.ptr<!spv.struct<i32 [0]>, PhysicalStorageBuffer>
   spv.globalVariable @var1 : !spv.ptr<!spv.struct<i32>, PhysicalStorageBuffer>
 }

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: vce_triple = #spv.vce<v1.0, [Shader], []>
-spv.module "Logical" "GLSL450" {
-  spv.func @iadd(%val : i32) -> i32 "None" {
-    %0 = spv.IAdd %val, %val: i32
-    spv.ReturnValue %0: i32
-  }
-} attributes {
+// CHECK: requires #spv.vce<v1.0, [Shader], []>
+spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.5, [Shader], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+} {
+  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: vce_triple = #spv.vce<v1.3, [GroupNonUniformBallot, Shader], []>
-spv.module "Logical" "GLSL450" {
-  spv.func @group_non_uniform_ballot(%predicate : i1) -> vector<4xi32> "None" {
-    %0 = spv.GroupNonUniformBallot "Workgroup" %predicate : vector<4xi32>
-    spv.ReturnValue %0: vector<4xi32>
-  }
-} attributes {
+// 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], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+} {
+  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: vce_triple = #spv.vce<v1.0, [Shader], []>
-spv.module "Logical" "GLSL450" {
-  spv.func @iadd(%val : i32) -> i32 "None" {
-    %0 = spv.IAdd %val, %val: i32
-    spv.ReturnValue %0: i32
-  }
-} attributes {
+// 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], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+} {
+  spv.func @iadd(%val : i32) -> i32 "None" {
+    %0 = spv.IAdd %val, %val: i32
+    spv.ReturnValue %0: i32
+  }
 }
 
 // Test deducing implied capability.
 // AtomicStorage implies Shader.
 
-// CHECK: vce_triple = #spv.vce<v1.0, [Shader], []>
-spv.module "Logical" "GLSL450" {
-  spv.func @iadd(%val : i32) -> i32 "None" {
-    %0 = spv.IAdd %val, %val: i32
-    spv.ReturnValue %0: i32
-  }
-} attributes {
+// CHECK: requires #spv.vce<v1.0, [Shader], []>
+spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [AtomicStorage], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+} {
+  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: vce_triple = #spv.vce<v1.3, [GroupNonUniformArithmetic, Shader], []>
-spv.module "Logical" "GLSL450" {
-  spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" {
-    %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32
-    spv.ReturnValue %0: i32
-  }
-} attributes {
+// 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], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
-}
-
-// CHECK: vce_triple = #spv.vce<v1.3, [GroupNonUniformClustered, GroupNonUniformBallot, Shader], []>
-spv.module "Logical" "GLSL450" {
+} {
   spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" {
     %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32
     spv.ReturnValue %0: i32
   }
-} attributes {
+}
+
+// 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], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+} {
+  spv.func @group_non_uniform_iadd(%val : i32) -> i32 "None" {
+    %0 = spv.GroupNonUniformIAdd "Subgroup" "Reduce" %val : i32
+    spv.ReturnValue %0: i32
+  }
 }
 
 //===----------------------------------------------------------------------===//
 // Extension
 //===----------------------------------------------------------------------===//
 
 // Test deducing minimal extensions.
 // spv.SubgroupBallotKHR requires the SPV_KHR_shader_ballot extension.
 
-// CHECK: vce_triple = #spv.vce<v1.0, [SubgroupBallotKHR, Shader], [SPV_KHR_shader_ballot]>
-spv.module "Logical" "GLSL450" {
-  spv.func @subgroup_ballot(%predicate : i1) -> vector<4xi32> "None" {
-    %0 = spv.SubgroupBallotKHR %predicate: vector<4xi32>
-    spv.ReturnValue %0: vector<4xi32>
-  }
-} attributes {
+// 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]>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+} {
+  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: vce_triple = #spv.vce<v1.0, [VulkanMemoryModel], [SPV_KHR_vulkan_memory_model]>
-spv.module "Logical" "Vulkan" {
-  spv.func @iadd(%val : i32) -> i32 "None" {
-    %0 = spv.IAdd %val, %val: i32
-    spv.ReturnValue %0: i32
-  }
-} attributes {
+// 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], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+} {
+  spv.func @iadd(%val : i32) -> i32 "None" {
+    %0 = spv.IAdd %val, %val: i32
+    spv.ReturnValue %0: i32
+  }
 }

mlir/test/Dialect/SPIRV/availability.mlir

 }
 
 // CHECK-LABEL: module_logical_glsl450
 func @module_logical_glsl450() {
   // CHECK: spv.module min version: v1.0
   // CHECK: spv.module max version: v1.5
   // CHECK: spv.module extensions: [ ]
   // CHECK: spv.module capabilities: [ [Shader] ]
-  spv.module "Logical" "GLSL450" { }
+  spv.module Logical GLSL450 { }
   return
 }
 
 // CHECK-LABEL: module_physical_storage_buffer64_vulkan
 func @module_physical_storage_buffer64_vulkan() {
   // CHECK: spv.module min version: v1.0
   // CHECK: spv.module max version: v1.5
   // CHECK: spv.module extensions: [ [SPV_EXT_physical_storage_buffer, SPV_KHR_physical_storage_buffer] [SPV_KHR_vulkan_memory_model] ]
   // CHECK: spv.module capabilities: [ [PhysicalStorageBufferAddresses] [VulkanMemoryModel] ]
-  spv.module "PhysicalStorageBuffer64" "Vulkan" { }
+  spv.module PhysicalStorageBuffer64 Vulkan { }
   return
 }

mlir/test/Dialect/SPIRV/control-flow-ops.mlir

 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.FunctionCall
 //===----------------------------------------------------------------------===//
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @fmain(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>, %arg2 : i32) -> i32 "None" {
     // CHECK: {{%.*}} = spv.FunctionCall @f_0({{%.*}}, {{%.*}}) : (vector<4xf32>, vector<4xf32>) -> vector<4xf32>
     %0 = spv.FunctionCall @f_0(%arg0, %arg1) : (vector<4xf32>, vector<4xf32>) -> vector<4xf32>
     // CHECK: spv.FunctionCall @f_1({{%.*}}, {{%.*}}) : (vector<4xf32>, vector<4xf32>) -> ()
     spv.FunctionCall @f_1(%0, %arg1) : (vector<4xf32>, vector<4xf32>) ->  ()
     // CHECK: spv.FunctionCall @f_2() : () -> ()
     spv.FunctionCall @f_2() : () -> ()
     // CHECK: {{%.*}} = spv.FunctionCall @f_3({{%.*}}) : (i32) -> i32

 func @caller() {
   // CHECK: spv.FunctionCall
   spv.FunctionCall @callee() : () -> ()
   spv.Return
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @f_invalid_result_type(%arg0 : i32, %arg1 : i32) -> () "None" {
     // expected-error @+1 {{expected callee function to have 0 or 1 result, but provided 2}}
     %0:2 = spv.FunctionCall @f_invalid_result_type(%arg0, %arg1) : (i32, i32) -> (i32, i32)
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @f_result_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" {
     // expected-error @+1 {{has incorrect number of results has for callee: expected 0, but provided 1}}
     %1 = spv.FunctionCall @f_result_type_mismatch(%arg0, %arg0) : (i32, i32) -> (i32)
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" {
     // expected-error @+1 {{has incorrect number of operands for callee: expected 2, but provided 1}}
     spv.FunctionCall @f_type_mismatch(%arg0) : (i32) -> ()
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" {
     %0 = spv.constant 2.0 : f32
     // expected-error @+1 {{operand type mismatch: expected operand type 'i32', but provided 'f32' for operand number 1}}
     spv.FunctionCall @f_type_mismatch(%arg0, %0) : (i32, f32) -> ()
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> i32 "None" {
     %cst = spv.constant 0: i32
     // expected-error @+1 {{result type mismatch: expected 'i32', but provided 'f32'}}
     %0 = spv.FunctionCall @f_type_mismatch(%arg0, %arg0) : (i32, i32) -> f32
     spv.ReturnValue %cst: i32
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @f_foo(%arg0 : i32, %arg1 : i32) -> i32 "None" {
     // expected-error @+1 {{op callee function 'f_undefined' not found in nearest symbol table}}
     %0 = spv.FunctionCall @f_undefined(%arg0, %arg0) : (i32, i32) -> i32
     spv.ReturnValue %0: i32
   }
 }
 
 // -----

 "foo.function"() ({
   // expected-error @+1 {{op must appear in a function-like op's block}}
   spv.Return
 })  : () -> ()
 
 // -----
 
 // Return mismatches function signature
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @work() -> (i32) "None" {
     // expected-error @+1 {{cannot be used in functions returning value}}
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @in_nested_region(%cond: i1) -> (i32) "None" {
     spv.selection {
       spv.BranchConditional %cond, ^then, ^merge
     ^then:
       // expected-error @+1 {{cannot be used in functions returning value}}
       spv.Return
     ^merge:
       spv._merge

 "foo.function"() ({
   %0 = spv.constant true
   // expected-error @+1 {{op must appear in a function-like op's block}}
   spv.ReturnValue %0 : i1
 })  : () -> ()
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @value_count_mismatch() -> () "None" {
     %0 = spv.constant 42 : i32
     // expected-error @+1 {{op returns 1 value but enclosing function requires 0 results}}
     spv.ReturnValue %0 : i32
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @value_type_mismatch() -> (f32) "None" {
     %0 = spv.constant 42 : i32
     // expected-error @+1 {{return value's type ('i32') mismatch with function's result type ('f32')}}
     spv.ReturnValue %0 : i32
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @in_nested_region(%cond: i1) -> () "None" {
     spv.selection {
       spv.BranchConditional %cond, ^then, ^merge
     ^then:
       %cst = spv.constant 0: i32
       // expected-error @+1 {{op returns 1 value but enclosing function requires 0 results}}
       spv.ReturnValue %cst: i32
     ^merge:

mlir/test/Dialect/SPIRV/ops.mlir

 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.ExecutionMode
 //===----------------------------------------------------------------------===//
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    spv.EntryPoint "GLCompute" @do_nothing
    // CHECK: spv.ExecutionMode {{@.*}} "ContractionOff"
    spv.ExecutionMode @do_nothing "ContractionOff"
 }
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    spv.EntryPoint "GLCompute" @do_nothing
    // CHECK: spv.ExecutionMode {{@.*}} "LocalSizeHint", 3, 4, 5
    spv.ExecutionMode @do_nothing "LocalSizeHint", 3, 4, 5
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    spv.EntryPoint "GLCompute" @do_nothing
    // expected-error @+1 {{custom op 'spv.ExecutionMode' invalid execution_mode attribute specification: "GLCompute"}}
    spv.ExecutionMode @do_nothing "GLCompute", 3, 4, 5
 }
 
func @aligned_load_incorrect_attributes() -> () {
   %0 = spv.Variable : !spv.ptr<f32, Function>
   // expected-error @+1 {{expected ']'}}
   %1 = spv.Load "Function" %0 ["Aligned", 4, 23] : f32
   return
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.globalVariable @var0 : !spv.ptr<f32, Input>
   // CHECK_LABEL: @simple_load
   spv.func @simple_load() -> () "None" {
     // CHECK: spv.Load "Input" {{%.*}} : f32
     %0 = spv._address_of @var0 : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.Return
   }
func @aligned_store_incorrect_attributes(%arg0 : f32) -> () {
   %0 = spv.Variable : !spv.ptr<f32, Function>
   // expected-error @+1 {{expected ']'}}
   spv.Store  "Function" %0, %arg0 ["Aligned", 4, 23] : f32
   return
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.globalVariable @var0 : !spv.ptr<f32, Input>
   spv.func @simple_store(%arg0 : f32) -> () "None" {
     %0 = spv._address_of @var0 : !spv.ptr<f32, Input>
     // CHECK: spv.Store  "Input" {{%.*}}, {{%.*}} : f32
     spv.Store  "Input" %0, %arg0 : f32
     spv.Return
   }
 }
func @variable_init_normal_constant() -> () {
   %0 = spv.constant 4.0 : f32
   // CHECK: spv.Variable init(%0) : !spv.ptr<f32, Function>
   %1 = spv.Variable init(%0) : !spv.ptr<f32, Function>
   return
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.globalVariable @global : !spv.ptr<f32, Workgroup>
   spv.func @variable_init_global_variable() -> () "None" {
     %0 = spv._address_of @global : !spv.ptr<f32, Workgroup>
     // CHECK: spv.Variable init({{.*}}) : !spv.ptr<!spv.ptr<f32, Workgroup>, Function>
     %1 = spv.Variable init(%0) : !spv.ptr<!spv.ptr<f32, Workgroup>, Function>
     spv.Return
   }
-} attributes {
-  capability = ["VariablePointers"],
-  extension = ["SPV_KHR_variable_pointers"]
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.specConstant @sc = 42 : i32
   // CHECK-LABEL: @variable_init_spec_constant
   spv.func @variable_init_spec_constant() -> () "None" {
     %0 = spv._reference_of @sc : i32
     // CHECK: spv.Variable init(%0) : !spv.ptr<i32, Function>
     %1 = spv.Variable init(%0) : !spv.ptr<i32, Function>
     spv.Return
   }

mlir/test/Dialect/SPIRV/structure-ops.mlir

 // RUN: mlir-opt -split-input-file -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spv._address_of
 //===----------------------------------------------------------------------===//
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.globalVariable @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
   spv.func @access_chain() -> () "None" {
     %0 = spv.constant 1: i32
     // CHECK: [[VAR1:%.*]] = spv._address_of @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4 x f32>>, Input>
     // CHECK-NEXT: spv.AccessChain [[VAR1]][{{.*}}, {{.*}}] : !spv.ptr<!spv.struct<f32, !spv.array<4 x f32>>, Input>
     %1 = spv._address_of @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
     %2 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
     spv.Return
   }
 }
 
 // -----
 
 // Allow taking address of global variables in other module-like ops
 spv.globalVariable @var : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
 func @address_of() -> () {
   // CHECK: spv._address_of @var
   %1 = spv._address_of @var : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
   return
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.globalVariable @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{expected spv.globalVariable symbol}}
     %0 = spv._address_of @var2 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.globalVariable @var1 : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Input>
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{result type mismatch with the referenced global variable's type}}
     %0 = spv._address_of @var1 : !spv.ptr<f32, Input>
   }
 }
 
 // -----

 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.EntryPoint
 //===----------------------------------------------------------------------===//
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    // CHECK: spv.EntryPoint "GLCompute" @do_nothing
    spv.EntryPoint "GLCompute" @do_nothing
 }
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
    spv.globalVariable @var2 : !spv.ptr<f32, Input>
    spv.globalVariable @var3 : !spv.ptr<f32, Output>
    spv.func @do_something(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) -> () "None" {
      %1 = spv.Load "Input" %arg0 : f32
      spv.Store "Output" %arg1, %1 : f32
      spv.Return
    }
    // CHECK: spv.EntryPoint "GLCompute" @do_something, @var2, @var3
    spv.EntryPoint "GLCompute" @do_something, @var2, @var3
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    // expected-error @+1 {{invalid kind of attribute specified}}
    spv.EntryPoint "GLCompute" "do_nothing"
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    // expected-error @+1 {{function 'do_something' not found in 'spv.module'}}
    spv.EntryPoint "GLCompute" @do_something
 }
 
 /// TODO(ravishankarm) : Add a test that verifies an error is thrown
 /// when interface entries of EntryPointOp are not
 /// spv.Variables. There is currently no other op that has a spv.ptr
 /// return type
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      // expected-error @+1 {{op must appear in a module-like op's block}}
      spv.EntryPoint "GLCompute" @do_something
    }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    spv.EntryPoint "GLCompute" @do_nothing
    // expected-error @+1 {{duplicate of a previous EntryPointOp}}
    spv.EntryPoint "GLCompute" @do_nothing
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    spv.EntryPoint "GLCompute" @do_nothing
-   // expected-error @+1 {{custom op 'spv.EntryPoint' invalid execution_model attribute specification: "ContractionOff"}}
+   // expected-error @+1 {{'spv.EntryPoint' invalid execution_model attribute specification: "ContractionOff"}}
    spv.EntryPoint "ContractionOff" @do_nothing
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.func
 //===----------------------------------------------------------------------===//

 // -----
 
 // expected-error @+1 {{expected SSA identifier}}
 spv.func @cannot_have_variadic_arguments(%arg: i32, ...) "None"
 
 // -----
 
 // Nested function
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @outer_func() -> () "None" {
     // expected-error @+1 {{must appear in a module-like op's block}}
     spv.func @inner_func() -> () "None" {
       spv.Return
     }
     spv.Return
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.globalVariable
 //===----------------------------------------------------------------------===//
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // CHECK: spv.globalVariable @var0 : !spv.ptr<f32, Input>
   spv.globalVariable @var0 : !spv.ptr<f32, Input>
 }
 
 // TODO: Fix test case after initialization with normal constant is addressed
-// spv.module "Logical" "GLSL450" {
+// spv.module Logical GLSL450 {
 //   %0 = spv.constant 4.0 : f32
 //   // CHECK1: spv.Variable init(%0) : !spv.ptr<f32, Private>
 //   spv.globalVariable @var1 init(%0) : !spv.ptr<f32, Private>
 // }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.specConstant @sc = 4.0 : f32
   // CHECK: spv.globalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
   spv.globalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
 }
 
 // -----
 
 // Allow initializers coming from other module-like ops
 spv.specConstant @sc = 4.0 : f32
 // CHECK: spv.globalVariable @var initializer(@sc)
 spv.globalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // CHECK: spv.globalVariable @var0 bind(1, 2) : !spv.ptr<f32, Uniform>
   spv.globalVariable @var0 bind(1, 2) : !spv.ptr<f32, Uniform>
 }
 
 // TODO: Fix test case after initialization with constant is addressed
-// spv.module "Logical" "GLSL450" {
+// spv.module Logical GLSL450 {
 //   %0 = spv.constant 4.0 : f32
 //   // CHECK1: spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
 //   spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
 // }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // CHECK: spv.globalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
   spv.globalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
   // CHECK: spv.globalVariable @var2 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
   spv.globalVariable @var2 {built_in = "GlobalInvocationID"} : !spv.ptr<vector<3xi32>, Input>
 }
 
 // -----
 
 // Allow in other module-like ops
 module {
   // CHECK: spv.globalVariable
   spv.globalVariable @var0 : !spv.ptr<f32, Input>
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // expected-error @+1 {{expected spv.ptr type}}
   spv.globalVariable @var0 : f32
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // expected-error @+1 {{op initializer must be result of a spv.specConstant or spv.globalVariable op}}
   spv.globalVariable @var0 initializer(@var1) : !spv.ptr<f32, Private>
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // expected-error @+1 {{storage class cannot be 'Generic'}}
   spv.globalVariable @var0 : !spv.ptr<f32, Generic>
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @foo() "None" {
     // expected-error @+1 {{op must appear in a module-like op's block}}
     spv.globalVariable @var0 : !spv.ptr<f32, Input>
     spv.Return
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.module
 //===----------------------------------------------------------------------===//
 
 // Module without capability and extension
-// CHECK: spv.module "Logical" "GLSL450"
-spv.module "Logical" "GLSL450" { }
+// CHECK: spv.module Logical GLSL450
+spv.module Logical GLSL450 { }
 
-// Module with capability and extension
-// CHECK: attributes {capability = ["Shader"], extension = ["SPV_KHR_16bit_storage"]}
-spv.module "Logical" "GLSL450" { } attributes {
-  capability = ["Shader"],
-  extension = ["SPV_KHR_16bit_storage"]
-}
+
+// Module with (version, capabilities, extensions) triple
+// CHECK: spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]>
+spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> { }
+
+// Module with additional attributes
+// CHECK: spv.module Logical GLSL450 attributes {foo = "bar"}
+spv.module Logical GLSL450 attributes {foo = "bar"} { }
+
+// Module with VCE triple and additional attributes
+// CHECK: spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> attributes {foo = "bar"}
+spv.module Logical GLSL450
+  requires #spv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]>
+  attributes {foo = "bar"} { }
 
 // Module with explicit spv._module_end
 // CHECK: spv.module
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv._module_end
 }
 
 // Module with function
 // CHECK: spv.module
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @do_nothing() -> () "None" {
     spv.Return
   }
 }
 
 // -----
 
 // Missing addressing model
-// expected-error@+1 {{custom op 'spv.module' expected addressing_model attribute specified as string}}
+// expected-error@+1 {{'spv.module' expected valid keyword}}
 spv.module { }
 
 // -----
 
 // Wrong addressing model
-// expected-error@+1 {{custom op 'spv.module' invalid addressing_model attribute specification: "Physical"}}
-spv.module "Physical" { }
+// expected-error@+1 {{'spv.module' invalid addressing_model attribute specification: Physical}}
+spv.module Physical { }
 
 // -----
 
 // Missing memory model
-// expected-error@+1 {{custom op 'spv.module' expected memory_model attribute specified as string}}
-spv.module "Logical" { }
+// expected-error@+1 {{'spv.module' expected valid keyword}}
+spv.module Logical { }
 
 // -----
 
 // Wrong memory model
-// expected-error@+1 {{custom op 'spv.module' invalid memory_model attribute specification: "Bla"}}
-spv.module "Logical" "Bla" { }
+// expected-error@+1 {{'spv.module' invalid memory_model attribute specification: Bla}}
+spv.module Logical Bla { }
 
 // -----
 
 // Module with multiple blocks
 // expected-error @+1 {{expects region #0 to have 0 or 1 blocks}}
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
 ^first:
   spv.Return
 ^second:
   spv.Return
 }
 
 // -----
 
 // Module with wrong terminator
 // expected-error@+2 {{expects regions to end with 'spv._module_end'}}
 // expected-note@+1 {{in custom textual format, the absence of terminator implies 'spv._module_end'}}
 "spv.module"() ({
   %0 = spv.constant true
 }) {addressing_model = 0 : i32, memory_model = 1 : i32} : () -> ()
 
 // -----
 
 // Use non SPIR-V op inside module
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // expected-error @+1 {{'spv.module' can only contain spv.* ops}}
   "dialect.op"() : () -> ()
 }
 
 // -----
 
 // Use non SPIR-V op inside function
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @do_nothing() -> () "None" {
     // expected-error @+1 {{functions in 'spv.module' can only contain spv.* ops}}
     "dialect.op"() : () -> ()
   }
 }
 
 // -----
 
 // Use external function
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // expected-error @+1 {{'spv.module' cannot contain external functions}}
   spv.func @extern() -> () "None"
 }
 
 // -----
 
-// expected-error @+1 {{uses unknown capability: MyAwesomeCapability}}
-spv.module "Logical" "GLSL450" {
-} attributes {
-  capabilities = ["MyAwesomeCapability"]
-}
-
-// -----
-
-// expected-error @+1 {{uses unknown extension: MyAwesomeExtension}}
-spv.module "Logical" "GLSL450" {
-} attributes {
-  extensions = ["MyAwesomeExtension"]
-}
-
-// -----
-
 //===----------------------------------------------------------------------===//
 // spv._module_end
 //===----------------------------------------------------------------------===//
 
 func @module_end_not_in_module() -> () {
   // expected-error @+1 {{op must appear in a module-like op's block}}
   spv._module_end
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv._reference_of
 //===----------------------------------------------------------------------===//
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.specConstant @sc1 = false
   spv.specConstant @sc2 = 42 : i64
   spv.specConstant @sc3 = 1.5 : f32
 
   // CHECK-LABEL: @reference
   spv.func @reference() -> i1 "None" {
     // CHECK: spv._reference_of @sc1 : i1
     %0 = spv._reference_of @sc1 : i1

 func @reference_of() {
   // CHECK: spv._reference_of @sc
   %0 = spv._reference_of @sc : i32
   return
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{expected spv.specConstant symbol}}
     %0 = spv._reference_of @sc : i32
     spv.Return
   }
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   spv.specConstant @sc = 42 : i32
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{result type mismatch with the referenced specialization constant's type}}
     %0 = spv._reference_of @sc : f32
     spv.Return
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.specConstant
 //===----------------------------------------------------------------------===//
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // CHECK: spv.specConstant @sc1 = false
   spv.specConstant @sc1 = false
   // CHECK: spv.specConstant @sc2 spec_id(5) = 42 : i64
   spv.specConstant @sc2 spec_id(5) = 42 : i64
   // CHECK: spv.specConstant @sc3 = 1.500000e+00 : f32
   spv.specConstant @sc3 = 1.5 : f32
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // expected-error @+1 {{SpecId cannot be negative}}
   spv.specConstant @sc2 spec_id(-5) = 42 : i64
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // expected-error @+1 {{default value bitwidth disallowed}}
   spv.specConstant @sc = 15 : i4
 }
 
 // -----
 
-spv.module "Logical" "GLSL450" {
+spv.module Logical GLSL450 {
   // expected-error @+1 {{default value can only be a bool, integer, or float scalar}}
   spv.specConstant @sc = dense<[2, 3]> : vector<2xi32>
 }
 
 // -----
 
 func @use_in_function() -> () {
   // expected-error @+1 {{op must appear in a module-like op's block}}
   spv.specConstant @sc = false
   return
 }

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

func @subgroup_ballot_missing_extension(%predicate: i1) -> vector<4xi32> attributes {
   %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
   return %0: vector<4xi32>
 }
 
 // CHECK-LABEL: @module_suitable_extension1
 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]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
-  // CHECK: spv.module "PhysicalStorageBuffer64" "Vulkan"
+  // CHECK: spv.module PhysicalStorageBuffer64 Vulkan
   "test.convert_to_module_op"() : () ->()
   return
 }
 
 // CHECK-LABEL: @module_suitable_extension2
 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]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
-  // CHECK: spv.module "PhysicalStorageBuffer64" "Vulkan"
+  // CHECK: spv.module PhysicalStorageBuffer64 Vulkan
   "test.convert_to_module_op"() : () -> ()
   return
 }
 
 // CHECK-LABEL: @module_missing_extension_mm
 func @module_missing_extension_mm() attributes {
   spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_physical_storage_buffer]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
func @module_missing_extension_am() attributes {
   return
 }
 
 // CHECK-LABEL: @module_implied_extension
 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], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
-  // CHECK: spv.module "PhysicalStorageBuffer64" "Vulkan"
+  // CHECK: spv.module PhysicalStorageBuffer64 Vulkan
   "test.convert_to_module_op"() : () -> ()
   return
 }

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} {
+module attributes {
+  gpu.container_module,
+  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, 128, 64]> : vector<3xi32>}>
+} {
   gpu.module @kernels {
     gpu.func @kernel_add(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
       attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32>}} {
       %0 = "gpu.block_id"() {dimension = "x"} : () -> index
       %1 = load %arg0[%0] : memref<8xf32>
       %2 = load %arg1[%0] : memref<8xf32>
       %3 = addf %1, %2 : f32
       store %3, %arg2[%0] : memref<8xf32>

mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp

static LogicalResult runMLIRPasses(ModuleOp module) {
   PassManager passManager(module.getContext());
   applyPassManagerCLOptions(passManager);
 
   passManager.addPass(createGpuKernelOutliningPass());
   passManager.addPass(createLegalizeStdOpsForSPIRVLoweringPass());
   passManager.addPass(createConvertGPUToSPIRVPass());
   OpPassManager &modulePM = passManager.nest<spirv::ModuleOp>();
   modulePM.addPass(spirv::createLowerABIAttributesPass());
+  modulePM.addPass(spirv::createUpdateVersionCapabilityExtensionPass());
   passManager.addPass(createConvertGpuLaunchFuncToVulkanLaunchFuncPass());
   passManager.addPass(createLowerToLLVMPass());
   passManager.addPass(createConvertVulkanLaunchFuncToVulkanCallsPass());
   return passManager.run(module);
 }
 
 int main(int argc, char **argv) {
   llvm::llvm_shutdown_obj x;

mlir/unittests/Dialect/SPIRV/DeserializationTest.cpp

void expectDiagnostic(StringRef errorMessage) {
     EXPECT_THAT(diagnostic->str(), StrEq(std::string(errorMessage)));
   }
 
   //===--------------------------------------------------------------------===//
   // SPIR-V builder methods
   //===--------------------------------------------------------------------===//
 
   /// Adds the SPIR-V module header to `binary`.
-  void addHeader() { spirv::appendModuleHeader(binary, /*idBound=*/0); }
+  void addHeader() {
+    spirv::appendModuleHeader(binary, spirv::Version::V_1_0, /*idBound=*/0);
+  }
 
   /// Adds the SPIR-V instruction into `binary`.
   void addInstruction(spirv::Opcode op, ArrayRef<uint32_t> operands) {
     uint32_t wordCount = 1 + operands.size();
     binary.push_back(spirv::getPrefixedOpcode(wordCount, op));
     binary.append(operands.begin(), operands.end());
   }
 

mlir/unittests/Dialect/SPIRV/SerializationTest.cpp

 //===- SerializationTest.cpp - SPIR-V Serialization Tests -----------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains corner case tests for the SPIR-V serializer that are not
 // covered by normal serialization and deserialization roundtripping.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SPIRV/Serialization.h"
+#include "mlir/Dialect/SPIRV/SPIRVAttributes.h"
 #include "mlir/Dialect/SPIRV/SPIRVBinaryUtils.h"
 #include "mlir/Dialect/SPIRV/SPIRVDialect.h"
 #include "mlir/Dialect/SPIRV/SPIRVOps.h"
 #include "mlir/Dialect/SPIRV/SPIRVTypes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Location.h"
 #include "mlir/IR/MLIRContext.h"
 #include "llvm/ADT/DenseSet.h"
void createModuleOp() {
     OperationState state(UnknownLoc::get(&context),
                          spirv::ModuleOp::getOperationName());
     state.addAttribute("addressing_model",
                        builder.getI32IntegerAttr(static_cast<uint32_t>(
                            spirv::AddressingModel::Logical)));
     state.addAttribute("memory_model",
                        builder.getI32IntegerAttr(
                            static_cast<uint32_t>(spirv::MemoryModel::GLSL450)));
+    state.addAttribute("vce_triple",
+                       spirv::VerCapExtAttr::get(
+                           spirv::Version::V_1_0, ArrayRef<spirv::Capability>(),
+                           ArrayRef<spirv::Extension>(), &context));
     spirv::ModuleOp::build(&builder, state);
     module = cast<spirv::ModuleOp>(Operation::create(state));
   }
 
   Type getFloatStructType() {
     OpBuilder opBuilder(module.body());
     llvm::SmallVector<Type, 1> elementTypes{opBuilder.getF32Type()};
     llvm::SmallVector<spirv::StructType::LayoutInfo, 1> layoutInfo{0};