[mlir][spirv] Use separate attribute for (version, capabilities, extensions)

[antiagainst]

Summary

We also need the (version, capabilities, extensions) triple on the
spv.module op. Thus far we have been using separate 'extensions'
and 'capabilities' attributes there and 'version' is missing. Creating
a separate attribute for the trip allows us to reuse the assembly
form and verification.

[denis13]

LGTM! Thanks!

[mravishankar]

This is awesome! Thanks!

mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h

59

Nit: can you just take an integer here?

[antiagainst]

mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h

59

This constructor's purpose is to avoid recreating the IntegerAttr if already having one. So I have it this one. Normally one should use the above.

mlir/docs/Dialects/SPIR-V.md

 APIs. These execution environments affect the availability of certain SPIR-V
 features. For example, a [Vulkan 1.1][VulkanSpirv] implementation must support
 the 1.0, 1.1, 1.2, and 1.3 versions of SPIR-V and the 1.0 version of the SPIR-V
 extended instructions for GLSL. Further Vulkan extensions may enable more SPIR-V
 instructions.
 
 SPIR-V compilation should also take into consideration of the execution
 environment, so we generate SPIR-V modules valid for the target environment.
-This is conveyed by the `spv.target_env` attribute. It should be of
-`#spv.target_env` attribute kind, which is defined as:
+This is conveyed by the `spv.target_env` (`spirv::TargetEnvAttr`) attribute. It
+should be of `#spv.target_env` attribute kind, which is defined as:
 
 ```
-spirv-version    ::= `V_1_0` | `V_1_1` | ...
+spirv-version    ::= `v1.0` | `v1.1` | ...
 spirv-extension  ::= `SPV_KHR_16bit_storage` | `SPV_EXT_physical_storage_buffer` | ...
 spirv-capability ::= `Shader` | `Kernel` | `GroupNonUniform` | ...
 
 spirv-extension-list     ::= `[` (spirv-extension-elements)? `]`
 spirv-extension-elements ::= spirv-extension (`,` spirv-extension)*
 
 spirv-capability-list     ::= `[` (spirv-capability-elements)? `]`
 spirv-capability-elements ::= spirv-capability (`,` spirv-capability)*
 
 spirv-resource-limits ::= dictionary-attribute
 
-spirv-target-env-attribute ::= `#` `spv.target_env` `<`
-                                  spirv-version `,`
-                                  spirv-extensions-list `,`
-                                  spirv-capability-list `,`
+spirv-vce-attribute ::= `#` `spv.vce` `<`
+                            spirv-extensions-list `>`
+
+spirv-target-env-attribute ::= `#` `spv.target_env` `<`
+                                  spirv-vce-attribute,
                                   spirv-resource-limits `>`
 ```
 
 The attribute has a few fields:
 
+*   A `#spv.vce` (`spirv::VerCapExtAttr`) attribute:
     *   The target SPIR-V version.
     *   A list of SPIR-V extensions for the target.
     *   A list of SPIR-V capabilities for the target.
 *   A dictionary of target resource limits (see the
     [Vulkan spec][VulkanResourceLimits] for explanation):
     *   `max_compute_workgroup_invocations`
     *   `max_compute_workgroup_size`
 
 For example,
 
 ```
 module attributes {
 spv.target_env = #spv.target_env<
-    V_1_3, [SPV_KHR_8bit_storage], [Shader, GroupNonUniform]
+    #spv.vce<v1.3, [Shader, GroupNonUniform], [SPV_KHR_8bit_storage]>,
     {
       max_compute_workgroup_invocations = 128 : i32,
       max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>
     }>
 } { ... }
 ```
 
 Dialect conversion framework will utilize the information in `spv.target_env`

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

class SPV_StrEnumAttr<string name, string description,
   ]>;
   let cppNamespace = "::mlir::spirv";
 }
 
 //===----------------------------------------------------------------------===//
 // SPIR-V availability definitions
 //===----------------------------------------------------------------------===//
 
-def SPV_V_1_0 : I32EnumAttrCase<"V_1_0", 0>;
-def SPV_V_1_1 : I32EnumAttrCase<"V_1_1", 1>;
-def SPV_V_1_2 : I32EnumAttrCase<"V_1_2", 2>;
-def SPV_V_1_3 : I32EnumAttrCase<"V_1_3", 3>;
-def SPV_V_1_4 : I32EnumAttrCase<"V_1_4", 4>;
-def SPV_V_1_5 : I32EnumAttrCase<"V_1_5", 5>;
+def SPV_V_1_0 : I32EnumAttrCase<"V_1_0", 0, "v1.0">;
+def SPV_V_1_1 : I32EnumAttrCase<"V_1_1", 1, "v1.1">;
+def SPV_V_1_2 : I32EnumAttrCase<"V_1_2", 2, "v1.2">;
+def SPV_V_1_3 : I32EnumAttrCase<"V_1_3", 3, "v1.3">;
+def SPV_V_1_4 : I32EnumAttrCase<"V_1_4", 4, "v1.4">;
+def SPV_V_1_5 : I32EnumAttrCase<"V_1_5", 5, "v1.5">;
 
 def SPV_VersionAttr : SPV_I32EnumAttr<"Version", "valid SPIR-V version", [
     SPV_V_1_0, SPV_V_1_1, SPV_V_1_2, SPV_V_1_3, SPV_V_1_4, SPV_V_1_5]>;
 
 class MinVersion<I32EnumAttrCase min> : MinVersionBase<
     "QueryMinVersionInterface", SPV_VersionAttr, min> {
   let interfaceDescription = [{
     Querying interface for minimal required SPIR-V version.

mlir/include/mlir/Dialect/SPIRV/TargetAndABI.h

 namespace spirv {
 enum class Capability : uint32_t;
 enum class Extension;
 enum class StorageClass : uint32_t;
 enum class Version : uint32_t;
 
 namespace detail {
 struct TargetEnvAttributeStorage;
+struct VerCapExtAttributeStorage;
 } // namespace detail
 
 /// SPIR-V dialect-specific attribute kinds.
 // TODO(antiagainst): move to a more suitable place if we have more attributes.
 namespace AttrKind {
 enum Kind {
-  TargetEnv = Attribute::FIRST_SPIRV_ATTR,
+  TargetEnv = Attribute::FIRST_SPIRV_ATTR, /// Target environment
+  VerCapExt, /// (version, extension, capability) triple
 };
 } // namespace AttrKind
 
-/// An attribute that specifies the target version, allowed extensions and
-/// capabilities, and resource limits. These information describles a SPIR-V
-/// target environment.
-class TargetEnvAttr
-    : public Attribute::AttrBase<TargetEnvAttr, Attribute,
-                                 detail::TargetEnvAttributeStorage> {
+/// An attribute that specifies the SPIR-V (version, capabilities, extensions)
+/// triple.
+class VerCapExtAttr
+    : public Attribute::AttrBase<VerCapExtAttr, Attribute,
+                                 detail::VerCapExtAttributeStorage> {
 public:
   using Base::Base;
 
-  /// Gets a TargetEnvAttr instance.
-  static TargetEnvAttr get(Version version, ArrayRef<Extension> extensions,
-                           ArrayRef<Capability> capabilities,
-                           DictionaryAttr limits);
-  static TargetEnvAttr get(IntegerAttr version, ArrayAttr extensions,
-                           ArrayAttr capabilities, DictionaryAttr limits);
+  /// Gets a VerCapExtAttr instance.
+  static VerCapExtAttr get(Version version, ArrayRef<Capability> capabilities,
+                           ArrayRef<Extension> extensions,
+                           MLIRContext *context);
+  static VerCapExtAttr get(IntegerAttr version, ArrayAttr capabilities,
+                           ArrayAttr extensions);
 
   /// Returns the attribute kind's name (without the 'spv.' prefix).
   static StringRef getKindName();
 
-  /// Returns the target version.
+  /// Returns the version.
   Version getVersion();
 
   struct ext_iterator final
       : public llvm::mapped_iterator<ArrayAttr::iterator,
                                      Extension (*)(Attribute)> {
     explicit ext_iterator(ArrayAttr::iterator it);
   };
   using ext_range = llvm::iterator_range<ext_iterator>;
 
-  /// Returns the target extensions.
+  /// Returns the extensions.
   ext_range getExtensions();
-  /// Returns the target extensions as a string array attribute.
+  /// Returns the extensions as a string array attribute.
   ArrayAttr getExtensionsAttr();
 
   struct cap_iterator final
       : public llvm::mapped_iterator<ArrayAttr::iterator,
                                      Capability (*)(Attribute)> {
     explicit cap_iterator(ArrayAttr::iterator it);
   };
   using cap_range = llvm::iterator_range<cap_iterator>;
 
-  /// Returns the target capabilities.
+  /// Returns the capabilities.
   cap_range getCapabilities();
+  /// Returns the capabilities as an integer array attribute.
+  ArrayAttr getCapabilitiesAttr();
+
+  static bool kindof(unsigned kind) { return kind == AttrKind::VerCapExt; }
+
+  static LogicalResult verifyConstructionInvariants(Location loc,
+                                                    IntegerAttr version,
+                                                    ArrayAttr capabilities,
+                                                    ArrayAttr extensions);
+};
+
+/// An attribute that specifies the target version, allowed extensions and
+/// capabilities, and resource limits. These information describles a SPIR-V
+/// target environment.
+class TargetEnvAttr
+    : public Attribute::AttrBase<TargetEnvAttr, Attribute,
+                                 detail::TargetEnvAttributeStorage> {
+public:
+  using Base::Base;
+
+  /// Gets a TargetEnvAttr instance.
+  static TargetEnvAttr get(VerCapExtAttr triple, DictionaryAttr limits);
+
+  /// Returns the attribute kind's name (without the 'spv.' prefix).
+  static StringRef getKindName();
+
+  /// Returns the (version, capabilities, extensions) triple attribute.
+  VerCapExtAttr getTripleAttr();
+
+  /// Returns the target version.
+  Version getVersion();
+
+  /// Returns the target extensions.
+  VerCapExtAttr::ext_range getExtensions();
+  /// Returns the target extensions as a string array attribute.
+  ArrayAttr getExtensionsAttr();
+
+  /// Returns the target capabilities.
+  VerCapExtAttr::cap_range getCapabilities();
   /// Returns the target capabilities as an integer array attribute.
   ArrayAttr getCapabilitiesAttr();
 
   /// Returns the target resource limits.
   ResourceLimitsAttr getResourceLimits();
 
   static bool kindof(unsigned kind) { return kind == AttrKind::TargetEnv; }
 
   static LogicalResult verifyConstructionInvariants(Location loc,
-                                                    IntegerAttr version,
-                                                    ArrayAttr extensions,
-                                                    ArrayAttr capabilities,
+                                                    VerCapExtAttr triple,
                                                     DictionaryAttr limits);
 };
 
 /// Returns the attribute name for specifying argument ABI information.
 StringRef getInterfaceVarABIAttrName();
 
 /// Gets the InterfaceVarABIAttr given its fields.
 InterfaceVarABIAttr getInterfaceVarABIAttr(unsigned descriptorSet,

mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp

 //===----------------------------------------------------------------------===//
 // SPIR-V Dialect
 //===----------------------------------------------------------------------===//
 
 SPIRVDialect::SPIRVDialect(MLIRContext *context)
     : Dialect(getDialectNamespace(), context) {
   addTypes<ArrayType, ImageType, PointerType, RuntimeArrayType, StructType>();
 
-  addAttributes<TargetEnvAttr>();
+  addAttributes<TargetEnvAttr, VerCapExtAttr>();
 
   // Add SPIR-V ops.
   addOperations<
 #define GET_OP_LIST
 #include "mlir/Dialect/SPIRV/SPIRVOps.cpp.inc"
       >();
 
   addInterfaces<SPIRVInlinerInterface>();
static ParseResult parseKeywordList(
   } while (succeeded(parser.parseOptionalComma()));
 
   if (parser.parseRSquare())
     return failure();
 
   return success();
 }
 
-/// Parses a spirv::TargetEnvAttr.
-static Attribute parseTargetAttr(DialectAsmParser &parser) {
+static Attribute parseVerCapExtAttr(DialectAsmParser &parser) {
   if (parser.parseLess())
     return {};
 
   Builder &builder = parser.getBuilder();
 
   IntegerAttr versionAttr;
   {
     auto loc = parser.getCurrentLocation();
     StringRef version;
     if (parser.parseKeyword(&version) || parser.parseComma())
       return {};
 
     if (auto versionSymbol = spirv::symbolizeVersion(version)) {
       versionAttr =
           builder.getI32IntegerAttr(static_cast<uint32_t>(*versionSymbol));
     } else {
       parser.emitError(loc, "unknown version: ") << version;
       return {};
     }
   }
 
-  ArrayAttr extensionsAttr;
+  ArrayAttr capabilitiesAttr;
   {
-    SmallVector<Attribute, 1> extensions;
+    SmallVector<Attribute, 4> capabilities;
     llvm::SMLoc errorloc;
     StringRef errorKeyword;
 
-    auto processExtension = [&](llvm::SMLoc loc, StringRef extension) {
-      if (spirv::symbolizeExtension(extension)) {
-        extensions.push_back(builder.getStringAttr(extension));
+    auto processCapability = [&](llvm::SMLoc loc, StringRef capability) {
+      if (auto capSymbol = spirv::symbolizeCapability(capability)) {
+        capabilities.push_back(
+            builder.getI32IntegerAttr(static_cast<uint32_t>(*capSymbol)));
         return success();
       }
-      return errorloc = loc, errorKeyword = extension, failure();
+      return errorloc = loc, errorKeyword = capability, failure();
     };
-    if (parseKeywordList(parser, processExtension) || parser.parseComma()) {
+    if (parseKeywordList(parser, processCapability) || parser.parseComma()) {
       if (!errorKeyword.empty())
-        parser.emitError(errorloc, "unknown extension: ") << errorKeyword;
+        parser.emitError(errorloc, "unknown capability: ") << errorKeyword;
       return {};
     }
 
-    extensionsAttr = builder.getArrayAttr(extensions);
+    capabilitiesAttr = builder.getArrayAttr(capabilities);
   }
 
-  ArrayAttr capabilitiesAttr;
+  ArrayAttr extensionsAttr;
   {
-    SmallVector<Attribute, 4> capabilities;
+    SmallVector<Attribute, 1> extensions;
     llvm::SMLoc errorloc;
     StringRef errorKeyword;
 
-    auto processCapability = [&](llvm::SMLoc loc, StringRef capability) {
-      if (auto capSymbol = spirv::symbolizeCapability(capability)) {
-        capabilities.push_back(
-            builder.getI32IntegerAttr(static_cast<uint32_t>(*capSymbol)));
+    auto processExtension = [&](llvm::SMLoc loc, StringRef extension) {
+      if (spirv::symbolizeExtension(extension)) {
+        extensions.push_back(builder.getStringAttr(extension));
         return success();
       }
-      return errorloc = loc, errorKeyword = capability, failure();
+      return errorloc = loc, errorKeyword = extension, failure();
     };
-    if (parseKeywordList(parser, processCapability) || parser.parseComma()) {
+    if (parseKeywordList(parser, processExtension)) {
       if (!errorKeyword.empty())
-        parser.emitError(errorloc, "unknown capability: ") << errorKeyword;
+        parser.emitError(errorloc, "unknown extension: ") << errorKeyword;
       return {};
     }
 
-    capabilitiesAttr = builder.getArrayAttr(capabilities);
+    extensionsAttr = builder.getArrayAttr(extensions);
   }
 
+  if (parser.parseGreater())
+    return {};
+
+  return spirv::VerCapExtAttr::get(versionAttr, capabilitiesAttr,
+                                   extensionsAttr);
+}
+
+/// Parses a spirv::TargetEnvAttr.
+static Attribute parseTargetEnvAttr(DialectAsmParser &parser) {
+  if (parser.parseLess())
+    return {};
+
+  spirv::VerCapExtAttr tripleAttr;
+  if (parser.parseAttribute(tripleAttr) || parser.parseComma())
+    return {};
+
   DictionaryAttr limitsAttr;
   {
     auto loc = parser.getCurrentLocation();
     if (parser.parseAttribute(limitsAttr))
       return {};
 
     if (!limitsAttr.isa<spirv::ResourceLimitsAttr>()) {
       parser.emitError(
           loc,
           "limits must be a dictionary attribute containing two 32-bit integer "
           "attributes 'max_compute_workgroup_invocations' and "
           "'max_compute_workgroup_size'");
       return {};
     }
   }
 
   if (parser.parseGreater())
     return {};
 
-  return spirv::TargetEnvAttr::get(versionAttr, extensionsAttr,
-                                   capabilitiesAttr, limitsAttr);
+  return spirv::TargetEnvAttr::get(tripleAttr, limitsAttr);
 }
 
 Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser,
                                        Type type) const {
   // SPIR-V attributes are dictionaries so they do not have type.
   if (type) {
     parser.emitError(parser.getNameLoc(), "unexpected type");
     return {};
   }
 
   // Parse the kind keyword first.
   StringRef attrKind;
   if (parser.parseKeyword(&attrKind))
     return {};
 
   if (attrKind == spirv::TargetEnvAttr::getKindName())
-    return parseTargetAttr(parser);
+    return parseTargetEnvAttr(parser);
+  if (attrKind == spirv::VerCapExtAttr::getKindName())
+    return parseVerCapExtAttr(parser);
 
   parser.emitError(parser.getNameLoc(), "unknown SPIR-V attriubte kind: ")
       << attrKind;
   return {};
 }
 
 //===----------------------------------------------------------------------===//
 // Attribute Printing
 //===----------------------------------------------------------------------===//
 
-static void print(spirv::TargetEnvAttr targetEnv, DialectAsmPrinter &printer) {
+static void print(spirv::VerCapExtAttr triple, DialectAsmPrinter &printer) {
   auto &os = printer.getStream();
-  printer << spirv::TargetEnvAttr::getKindName() << "<"
-          << spirv::stringifyVersion(targetEnv.getVersion()) << ", [";
-  interleaveComma(targetEnv.getExtensionsAttr(), os, [&](Attribute attr) {
-    os << attr.cast<StringAttr>().getValue();
+  printer << spirv::VerCapExtAttr::getKindName() << "<"
+          << spirv::stringifyVersion(triple.getVersion()) << ", [";
+  interleaveComma(triple.getCapabilities(), os, [&](spirv::Capability cap) {
+    os << spirv::stringifyCapability(cap);
   });
   printer << "], [";
-  interleaveComma(targetEnv.getCapabilities(), os, [&](spirv::Capability cap) {
-    os << spirv::stringifyCapability(cap);
+  interleaveComma(triple.getExtensionsAttr(), os, [&](Attribute attr) {
+    os << attr.cast<StringAttr>().getValue();
   });
-  printer << "], " << targetEnv.getResourceLimits() << ">";
+  printer << "]>";
+}
+
+static void print(spirv::TargetEnvAttr targetEnv, DialectAsmPrinter &printer) {
+  printer << spirv::TargetEnvAttr::getKindName() << "<#spv.";
+  print(targetEnv.getTripleAttr(), printer);
+  printer << ", " << targetEnv.getResourceLimits() << ">";
 }
 
 void SPIRVDialect::printAttribute(Attribute attr,
                                   DialectAsmPrinter &printer) const {
   if (auto targetEnv = attr.dyn_cast<TargetEnvAttr>())
     print(targetEnv, printer);
+  else if (auto vceAttr = attr.dyn_cast<VerCapExtAttr>())
+    print(vceAttr, printer);
   else
     llvm_unreachable("unhandled SPIR-V attribute kind");
 }
 
 //===----------------------------------------------------------------------===//
 // Constant
 //===----------------------------------------------------------------------===//
 
 Operation *SPIRVDialect::materializeConstant(OpBuilder &builder,
                                              Attribute value, Type type,
                                              Location loc) {
-  if (!ConstantOp::isBuildableWith(type))
+  if (!spirv::ConstantOp::isBuildableWith(type))
     return nullptr;
 
   return builder.create<spirv::ConstantOp>(loc, type, value);
 }
 
 //===----------------------------------------------------------------------===//
 // Shader Interface ABI
 //===----------------------------------------------------------------------===//
 
 LogicalResult SPIRVDialect::verifyOperationAttribute(Operation *op,
                                                      NamedAttribute attribute) {
   StringRef symbol = attribute.first.strref();
   Attribute attr = attribute.second;
 
   // TODO(antiagainst): figure out a way to generate the description from the
   // StructAttr definition.
   if (symbol == spirv::getEntryPointABIAttrName()) {
     if (!attr.isa<spirv::EntryPointABIAttr>())
       return op->emitError("'")
              << symbol
              << "' attribute must be a dictionary attribute containing one "
                 "32-bit integer elements attribute: 'local_size'";
   } else if (symbol == spirv::getTargetEnvAttrName()) {
     if (!attr.isa<spirv::TargetEnvAttr>())
-      return op->emitError("'")
-             << symbol
-             << "' must be a dictionary attribute containing one 32-bit "
-                "integer attribute 'version', one string array attribute "
-                "'extensions', one 32-bit integer array attribute "
-                "'capabilities', and one dictionary attribute 'limits'";
+      return op->emitError("'") << symbol << "' must be a spirv::TargetEnvAttr";
   } else {
     return op->emitError("found unsupported '")
            << symbol << "' attribute on operation";
   }
 
   return success();
 }
 

mlir/lib/Dialect/SPIRV/TargetAndABI.cpp

 //===- TargetAndABI.cpp - SPIR-V target and ABI 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SPIRV/TargetAndABI.h"
 #include "mlir/Dialect/SPIRV/SPIRVTypes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/FunctionSupport.h"
 #include "mlir/IR/Operation.h"
 
 using namespace mlir;
 
+//===----------------------------------------------------------------------===//
+// DictionaryDict derived attributes
+//===----------------------------------------------------------------------===//
+
 namespace mlir {
 #include "mlir/Dialect/SPIRV/TargetAndABI.cpp.inc"
 
+//===----------------------------------------------------------------------===//
+// Attribute storage classes
+//===----------------------------------------------------------------------===//
+
 namespace spirv {
 namespace detail {
+struct VerCapExtAttributeStorage : public AttributeStorage {
+  using KeyTy = std::tuple<Attribute, Attribute, Attribute>;
+
+  VerCapExtAttributeStorage(Attribute version, Attribute capabilities,
+                            Attribute extensions)
+      : version(version), capabilities(capabilities), extensions(extensions) {}
+
+  bool operator==(const KeyTy &key) const {
+    return std::get<0>(key) == version && std::get<1>(key) == capabilities &&
+           std::get<2>(key) == extensions;
+  }
+
+  static VerCapExtAttributeStorage *
+  construct(AttributeStorageAllocator &allocator, const KeyTy &key) {
+    return new (allocator.allocate<VerCapExtAttributeStorage>())
+        VerCapExtAttributeStorage(std::get<0>(key), std::get<1>(key),
+                                  std::get<2>(key));
+  }
+
+  Attribute version;
+  Attribute capabilities;
+  Attribute extensions;
+};
+
 struct TargetEnvAttributeStorage : public AttributeStorage {
-  using KeyTy = std::tuple<Attribute, Attribute, Attribute, Attribute>;
+  using KeyTy = std::pair<Attribute, Attribute>;
 
-  TargetEnvAttributeStorage(Attribute version, Attribute extensions,
-                            Attribute capabilities, Attribute limits)
-      : version(version), extensions(extensions), capabilities(capabilities),
-        limits(limits) {}
+  TargetEnvAttributeStorage(Attribute triple, Attribute limits)
+      : triple(triple), limits(limits) {}
 
   bool operator==(const KeyTy &key) const {
-    return std::get<0>(key) == version && std::get<1>(key) == extensions &&
-           std::get<2>(key) == capabilities && std::get<3>(key) == limits;
+    return key.first == triple && key.second == limits;
   }
 
   static TargetEnvAttributeStorage *
   construct(AttributeStorageAllocator &allocator, const KeyTy &key) {
     return new (allocator.allocate<TargetEnvAttributeStorage>())
-        TargetEnvAttributeStorage(std::get<0>(key), std::get<1>(key),
-                                  std::get<2>(key), std::get<3>(key));
+        TargetEnvAttributeStorage(key.first, key.second);
   }
 
-  Attribute version;
-  Attribute extensions;
-  Attribute capabilities;
+  Attribute triple;
   Attribute limits;
 };
 } // namespace detail
 } // namespace spirv
 } // namespace mlir
 
-spirv::TargetEnvAttr spirv::TargetEnvAttr::get(
-    spirv::Version version, ArrayRef<spirv::Extension> extensions,
-    ArrayRef<spirv::Capability> capabilities, DictionaryAttr limits) {
-  Builder b(limits.getContext());
+//===----------------------------------------------------------------------===//
+// VerCapExtAttr
+//===----------------------------------------------------------------------===//
 
-  auto versionAttr = b.getI32IntegerAttr(static_cast<uint32_t>(version));
+spirv::VerCapExtAttr spirv::VerCapExtAttr::get(
+    spirv::Version version, ArrayRef<spirv::Capability> capabilities,
+    ArrayRef<spirv::Extension> extensions, MLIRContext *context) {
+  Builder b(context);
 
-  SmallVector<Attribute, 4> extAttrs;
-  extAttrs.reserve(extensions.size());
-  for (spirv::Extension ext : extensions)
-    extAttrs.push_back(b.getStringAttr(spirv::stringifyExtension(ext)));
+  auto versionAttr = b.getI32IntegerAttr(static_cast<uint32_t>(version));
 
   SmallVector<Attribute, 4> capAttrs;
   capAttrs.reserve(capabilities.size());
   for (spirv::Capability cap : capabilities)
     capAttrs.push_back(b.getI32IntegerAttr(static_cast<uint32_t>(cap)));
 
-  return get(versionAttr, b.getArrayAttr(extAttrs), b.getArrayAttr(capAttrs),
-             limits);
+  SmallVector<Attribute, 4> extAttrs;
+  extAttrs.reserve(extensions.size());
+  for (spirv::Extension ext : extensions)
+    extAttrs.push_back(b.getStringAttr(spirv::stringifyExtension(ext)));
+
+  return get(versionAttr, b.getArrayAttr(capAttrs), b.getArrayAttr(extAttrs));
 }
 
-spirv::TargetEnvAttr spirv::TargetEnvAttr::get(IntegerAttr version,
-                                               ArrayAttr extensions,
+spirv::VerCapExtAttr spirv::VerCapExtAttr::get(IntegerAttr version,
                                                ArrayAttr capabilities,
-                                               DictionaryAttr limits) {
-  assert(version && extensions && capabilities && limits);
+                                               ArrayAttr extensions) {
+  assert(version && capabilities && extensions);
   MLIRContext *context = version.getContext();
-  return Base::get(context, spirv::AttrKind::TargetEnv, version, extensions,
-                   capabilities, limits);
+  return Base::get(context, spirv::AttrKind::VerCapExt, version, capabilities,
+                   extensions);
 }
 
-StringRef spirv::TargetEnvAttr::getKindName() { return "target_env"; }
+StringRef spirv::VerCapExtAttr::getKindName() { return "vce"; }
 
-spirv::Version spirv::TargetEnvAttr::getVersion() {
+spirv::Version spirv::VerCapExtAttr::getVersion() {
   return static_cast<spirv::Version>(
       getImpl()->version.cast<IntegerAttr>().getValue().getZExtValue());
 }
 
-spirv::TargetEnvAttr::ext_iterator::ext_iterator(ArrayAttr::iterator it)
+spirv::VerCapExtAttr::ext_iterator::ext_iterator(ArrayAttr::iterator it)
     : llvm::mapped_iterator<ArrayAttr::iterator,
                             spirv::Extension (*)(Attribute)>(
           it, [](Attribute attr) {
             return *symbolizeExtension(attr.cast<StringAttr>().getValue());
           }) {}
 
-spirv::TargetEnvAttr::ext_range spirv::TargetEnvAttr::getExtensions() {
+spirv::VerCapExtAttr::ext_range spirv::VerCapExtAttr::getExtensions() {
   auto range = getExtensionsAttr().getValue();
   return {ext_iterator(range.begin()), ext_iterator(range.end())};
 }
 
-ArrayAttr spirv::TargetEnvAttr::getExtensionsAttr() {
+ArrayAttr spirv::VerCapExtAttr::getExtensionsAttr() {
   return getImpl()->extensions.cast<ArrayAttr>();
 }
 
-spirv::TargetEnvAttr::cap_iterator::cap_iterator(ArrayAttr::iterator it)
+spirv::VerCapExtAttr::cap_iterator::cap_iterator(ArrayAttr::iterator it)
     : llvm::mapped_iterator<ArrayAttr::iterator,
                             spirv::Capability (*)(Attribute)>(
           it, [](Attribute attr) {
             return *symbolizeCapability(
                 attr.cast<IntegerAttr>().getValue().getZExtValue());
           }) {}
 
-spirv::TargetEnvAttr::cap_range spirv::TargetEnvAttr::getCapabilities() {
+spirv::VerCapExtAttr::cap_range spirv::VerCapExtAttr::getCapabilities() {
   auto range = getCapabilitiesAttr().getValue();
   return {cap_iterator(range.begin()), cap_iterator(range.end())};
 }
 
-ArrayAttr spirv::TargetEnvAttr::getCapabilitiesAttr() {
+ArrayAttr spirv::VerCapExtAttr::getCapabilitiesAttr() {
   return getImpl()->capabilities.cast<ArrayAttr>();
 }
 
-spirv::ResourceLimitsAttr spirv::TargetEnvAttr::getResourceLimits() {
-  return getImpl()->limits.cast<spirv::ResourceLimitsAttr>();
-}
-
-LogicalResult spirv::TargetEnvAttr::verifyConstructionInvariants(
-    Location loc, IntegerAttr version, ArrayAttr extensions,
-    ArrayAttr capabilities, DictionaryAttr limits) {
+LogicalResult spirv::VerCapExtAttr::verifyConstructionInvariants(
+    Location loc, IntegerAttr version, ArrayAttr capabilities,
+    ArrayAttr extensions) {
   if (!version.getType().isSignlessInteger(32))
     return emitError(loc, "expected 32-bit integer for version");
 
+  if (!llvm::all_of(capabilities.getValue(), [](Attribute attr) {
+        if (auto intAttr = attr.dyn_cast<IntegerAttr>())
+          if (spirv::symbolizeCapability(intAttr.getValue().getZExtValue()))
+            return true;
+        return false;
+      }))
+    return emitError(loc, "unknown capability in capability list");
+
   if (!llvm::all_of(extensions.getValue(), [](Attribute attr) {
         if (auto strAttr = attr.dyn_cast<StringAttr>())
           if (spirv::symbolizeExtension(strAttr.getValue()))
             return true;
         return false;
       }))
     return emitError(loc, "unknown extension in extension list");
 
-  if (!llvm::all_of(capabilities.getValue(), [](Attribute attr) {
-        if (auto intAttr = attr.dyn_cast<IntegerAttr>())
-          if (spirv::symbolizeCapability(intAttr.getValue().getZExtValue()))
-            return true;
-        return false;
-      }))
-    return emitError(loc, "unknown capability in capability list");
+  return success();
+}
+
+//===----------------------------------------------------------------------===//
+// TargetEnvAttr
+//===----------------------------------------------------------------------===//
+
+spirv::TargetEnvAttr spirv::TargetEnvAttr::get(spirv::VerCapExtAttr triple,
+                                               DictionaryAttr limits) {
+  assert(triple && limits && "expected valid triple and limits");
+  MLIRContext *context = triple.getContext();
+  return Base::get(context, spirv::AttrKind::TargetEnv, triple, limits);
+}
+
+StringRef spirv::TargetEnvAttr::getKindName() { return "target_env"; }
 
+spirv::VerCapExtAttr spirv::TargetEnvAttr::getTripleAttr() {
+  return getImpl()->triple.cast<spirv::VerCapExtAttr>();
+}
+
+spirv::Version spirv::TargetEnvAttr::getVersion() {
+  return getTripleAttr().getVersion();
+}
+
+spirv::VerCapExtAttr::ext_range spirv::TargetEnvAttr::getExtensions() {
+  return getTripleAttr().getExtensions();
+}
+
+ArrayAttr spirv::TargetEnvAttr::getExtensionsAttr() {
+  return getTripleAttr().getExtensionsAttr();
+}
+
+spirv::VerCapExtAttr::cap_range spirv::TargetEnvAttr::getCapabilities() {
+  return getTripleAttr().getCapabilities();
+}
+
+ArrayAttr spirv::TargetEnvAttr::getCapabilitiesAttr() {
+  return getTripleAttr().getCapabilitiesAttr();
+}
+
+spirv::ResourceLimitsAttr spirv::TargetEnvAttr::getResourceLimits() {
+  return getImpl()->limits.cast<spirv::ResourceLimitsAttr>();
+}
+
+LogicalResult spirv::TargetEnvAttr::verifyConstructionInvariants(
+    Location loc, spirv::VerCapExtAttr triple, DictionaryAttr limits) {
   if (!limits.isa<spirv::ResourceLimitsAttr>())
     return emitError(loc, "expected spirv::ResourceLimitsAttr for limits");
 
   return success();
 }
 
+//===----------------------------------------------------------------------===//
+// Utility functions
+//===----------------------------------------------------------------------===//
+
 StringRef spirv::getInterfaceVarABIAttrName() {
   return "spv.interface_var_abi";
 }
 
 spirv::InterfaceVarABIAttr
 spirv::getInterfaceVarABIAttr(unsigned descriptorSet, unsigned binding,
                               spirv::StorageClass storageClass,
                               MLIRContext *context) {
spirv::getDefaultResourceLimits(MLIRContext *context) {
   return spirv::ResourceLimitsAttr ::get(
       IntegerAttr::get(i32Type, 128),
       DenseIntElementsAttr::get<int32_t>(v3i32Type, {128, 128, 64}), context);
 }
 
 StringRef spirv::getTargetEnvAttrName() { return "spv.target_env"; }
 
 spirv::TargetEnvAttr spirv::getDefaultTargetEnv(MLIRContext *context) {
-  Builder builder(context);
-  return spirv::TargetEnvAttr::get(
-      builder.getI32IntegerAttr(static_cast<uint32_t>(spirv::Version::V_1_0)),
-      builder.getI32ArrayAttr({}),
-      builder.getI32ArrayAttr(
-          {static_cast<uint32_t>(spirv::Capability::Shader)}),
+  auto triple = spirv::VerCapExtAttr::get(spirv::Version::V_1_0,
+                                          {spirv::Capability::Shader},
+                                          ArrayRef<Extension>(), context);
+  return spirv::TargetEnvAttr::get(triple,
                                    spirv::getDefaultResourceLimits(context));
 }
 
 spirv::TargetEnvAttr spirv::lookupTargetEnvOrDefault(Operation *op) {
   if (auto attr = op->getAttrOfType<spirv::TargetEnvAttr>(
           spirv::getTargetEnvAttrName()))
     return attr;
   return getDefaultTargetEnv(op->getContext());
 }

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

#single_workgroup_reduction_trait = {
   indexing_maps = [
     affine_map<(i) -> (i)>,
     affine_map<(i) -> (0)>
   ]
 }
 
 module attributes {
   spv.target_env = #spv.target_env<
-    V_1_3, [], [Shader, GroupNonUniformArithmetic],
+    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>,
     {
       max_compute_workgroup_invocations = 128 : i32,
       max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>
     }>
 } {
 
 // CHECK:      spv.globalVariable
 // CHECK-SAME: built_in("LocalInvocationId")
#single_workgroup_reduction_trait = {
   indexing_maps = [
     affine_map<(i) -> (i)>,
     affine_map<(i) -> (0)>
   ]
 }
 
 module attributes {
   spv.target_env = #spv.target_env<
-    V_1_3, [], [Shader, GroupNonUniformArithmetic],
+    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>,
     {
       max_compute_workgroup_invocations = 128 : i32,
       max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>
     }>
 } {
 func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>) {
   // expected-error @+1 {{failed to legalize operation 'linalg.generic'}}
   linalg.generic #single_workgroup_reduction_trait %input, %output {
#single_workgroup_reduction_trait = {
   indexing_maps = [
     affine_map<(i) -> (i)>,
     affine_map<(i) -> (0)>
   ]
 }
 
 module attributes {
   spv.target_env = #spv.target_env<
-    V_1_3, [], [Shader, GroupNonUniformArithmetic],
+    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>,
     {
       max_compute_workgroup_invocations = 128 : i32,
       max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>
     }>
 } {
 func @single_workgroup_reduction(%input: memref<16xi32>, %output: memref<1xi32>) attributes {
   spv.entry_point_abi = {local_size = dense<[32, 1, 1]>: vector<3xi32>}
 } {
#single_workgroup_reduction_trait = {
   indexing_maps = [
     affine_map<(i, j) -> (i, j)>,
     affine_map<(i, j) -> (i)>
   ]
 }
 
 module attributes {
   spv.target_env = #spv.target_env<
-    V_1_3, [], [Shader, GroupNonUniformArithmetic],
+    #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>,
     {
       max_compute_workgroup_invocations = 128 : i32,
       max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>
     }>
 } {
 func @single_workgroup_reduction(%input: memref<16x8xi32>, %output: memref<16xi32>) attributes {
   spv.entry_point_abi = {local_size = dense<[16, 8, 1]>: vector<3xi32>}
 } {

mlir/test/Dialect/SPIRV/availability.mlir

 // RUN: mlir-opt -disable-pass-threading -test-spirv-op-availability %s | FileCheck %s
 
 // CHECK-LABEL: iadd
 func @iadd(%arg: i32) -> i32 {
-  // CHECK: min version: V_1_0
-  // CHECK: max version: V_1_5
+  // CHECK: min version: v1.0
+  // CHECK: max version: v1.5
   // CHECK: extensions: [ ]
   // CHECK: capabilities: [ ]
   %0 = spv.IAdd %arg, %arg: i32
   return %0: i32
 }
 
 // CHECK: atomic_compare_exchange_weak
 func @atomic_compare_exchange_weak(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 {
-  // CHECK: min version: V_1_0
-  // CHECK: max version: V_1_3
+  // CHECK: min version: v1.0
+  // CHECK: max version: v1.3
   // CHECK: extensions: [ ]
   // CHECK: capabilities: [ [Kernel] ]
   %0 = spv.AtomicCompareExchangeWeak "Workgroup" "Release" "Acquire" %ptr, %value, %comparator: !spv.ptr<i32, Workgroup>
   return %0: i32
 }
 
 // CHECK-LABEL: subgroup_ballot
 func @subgroup_ballot(%predicate: i1) -> vector<4xi32> {
-  // CHECK: min version: V_1_3
-  // CHECK: max version: V_1_5
+  // CHECK: min version: v1.3
+  // CHECK: max version: v1.5
   // CHECK: extensions: [ ]
   // CHECK: capabilities: [ [GroupNonUniformBallot] ]
   %0 = spv.GroupNonUniformBallot "Workgroup" %predicate : vector<4xi32>
   return %0: vector<4xi32>
 }
 
 // CHECK-LABEL: module_logical_glsl450
 func @module_logical_glsl450() {
-  // CHECK: spv.module min version: V_1_0
-  // CHECK: spv.module max version: V_1_5
+  // 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" { }
   return
 }
 
 // CHECK-LABEL: module_physical_storage_buffer64_vulkan
 func @module_physical_storage_buffer64_vulkan() {
-  // CHECK: spv.module min version: V_1_0
-  // CHECK: spv.module max version: V_1_5
+  // 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" { }
   return
 }

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

 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.target_env
 //===----------------------------------------------------------------------===//
 
-func @target_env_wrong_type() attributes {
-  // expected-error @+1 {{expected valid keyword}}
-  spv.target_env = #spv.target_env<64>
+func @target_env_missing_limits() attributes {
+  spv.target_env = #spv.target_env<
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
+    // expected-error @+1 {{limits must be a dictionary attribute containing two 32-bit integer attributes 'max_compute_workgroup_invocations' and 'max_compute_workgroup_size'}}
+    {max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
 } { return }
 
 // -----
 
-func @target_env_missing_fields() attributes {
-  // expected-error @+1 {{expected ','}}
-  spv.target_env = #spv.target_env<V_1_0>
+func @target_env_wrong_limits() attributes {
+  spv.target_env = #spv.target_env<
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
+    // expected-error @+1 {{limits must be a dictionary attribute containing two 32-bit integer attributes 'max_compute_workgroup_invocations' and 'max_compute_workgroup_size'}}
+    {max_compute_workgroup_invocations = 128 : i64, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
 } { return }
 
 // -----
 
-func @target_env_wrong_version() attributes {
-  // expected-error @+1 {{unknown version: V_x_y}}
-  spv.target_env = #spv.target_env<V_x_y, []>
+func @target_env() attributes {
+  // CHECK:      spv.target_env = #spv.target_env<
+  // CHECK-SAME:   #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
+  // CHECK-SAME:   {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
+    {
+      max_compute_workgroup_invocations = 128 : i32,
+      max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>
+    }>
 } { return }
 
 // -----
 
-func @target_env_wrong_extension_type() attributes {
-  // expected-error @+1 {{expected valid keyword}}
-  spv.target_env = #spv.target_env<V_1_0, [32: i32], [Shader]>
+func @target_env_extra_fields() attributes {
+  // expected-error @+6 {{expected '>'}}
+  spv.target_env = #spv.target_env<
+    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
+    {
+      max_compute_workgroup_invocations = 128 : i32,
+      max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>
+    },
+    more_stuff
+  >
 } { return }
 
 // -----
 
-func @target_env_wrong_extension() attributes {
-  // expected-error @+1 {{unknown extension: SPV_Something}}
-  spv.target_env = #spv.target_env<V_1_0, [SPV_Something], [Shader]>
+//===----------------------------------------------------------------------===//
+// spv.vce
+//===----------------------------------------------------------------------===//
+
+func @vce_wrong_type() attributes {
+  // expected-error @+1 {{expected valid keyword}}
+  vce = #spv.vce<64>
 } { return }
 
 // -----
 
-func @target_env_wrong_capability() attributes {
-  // expected-error @+1 {{unknown capability: Something}}
-  spv.target_env = #spv.target_env<V_1_0, [], [Something]>
+func @vce_missing_fields() attributes {
+  // expected-error @+1 {{expected ','}}
+  vce = #spv.vce<v1.0>
 } { return }
 
 // -----
 
-func @target_env_missing_limits() attributes {
-  spv.target_env = #spv.target_env<
-    V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader],
-    // expected-error @+1 {{limits must be a dictionary attribute containing two 32-bit integer attributes 'max_compute_workgroup_invocations' and 'max_compute_workgroup_size'}}
-    {max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
+func @vce_wrong_version() attributes {
+  // expected-error @+1 {{unknown version: V_x_y}}
+  vce = #spv.vce<V_x_y, []>
 } { return }
 
 // -----
 
-func @target_env_wrong_limits() attributes {
-  spv.target_env = #spv.target_env<
-    V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader],
-    // expected-error @+1 {{limits must be a dictionary attribute containing two 32-bit integer attributes 'max_compute_workgroup_invocations' and 'max_compute_workgroup_size'}}
-    {max_compute_workgroup_invocations = 128 : i64, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
+func @vce_wrong_extension_type() attributes {
+  // expected-error @+1 {{expected valid keyword}}
+  vce = #spv.vce<v1.0, [32: i32], [Shader]>
 } { return }
 
 // -----
 
-func @target_env() attributes {
+func @vce_wrong_extension() attributes {
+  // expected-error @+1 {{unknown extension: SPV_Something}}
+  vce = #spv.vce<v1.0, [Shader], [SPV_Something]>
+} { return }
 
-  // CHECK: spv.target_env = #spv.target_env<V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
-  spv.target_env = #spv.target_env<
-    V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader],
-    {
-      max_compute_workgroup_invocations = 128 : i32,
-      max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>
-    }>
+// -----
+
+func @vce_wrong_capability() attributes {
+  // expected-error @+1 {{unknown capability: Something}}
+  vce = #spv.vce<v1.0, [Something], []>
 } { return }
 
 // -----
 
-func @target_env_extra_fields() attributes {
-  // expected-error @+6 {{expected '>'}}
-  spv.target_env = #spv.target_env<
-    V_1_0, [SPV_KHR_storage_buffer_storage_class], [Shader],
-    {
-      max_compute_workgroup_invocations = 128 : i32,
-      max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>
-    },
-    more_stuff
-  >
+func @vce() attributes {
+  // CHECK: #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>
+  vce = #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>
 } { return }

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

 // which extensions are incorporated into SPIR-V 1.5.
 
 //===----------------------------------------------------------------------===//
 // MaxVersion
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @cmp_exchange_weak_suitable_version_capabilities
 func @cmp_exchange_weak_suitable_version_capabilities(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
-  spv.target_env = #spv.target_env<V_1_1, [], [Kernel, AtomicStorage], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.1, [Kernel, AtomicStorage], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: spv.AtomicCompareExchangeWeak "Workgroup" "AcquireRelease|AtomicCounterMemory" "Acquire"
   %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32)
   return %0: i32
 }
 
 // CHECK-LABEL: @cmp_exchange_weak_unsupported_version
 func @cmp_exchange_weak_unsupported_version(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
-  spv.target_env = #spv.target_env<V_1_4, [], [Kernel, AtomicStorage], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.4, [Kernel, AtomicStorage], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: test.convert_to_atomic_compare_exchange_weak_op
   %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32)
   return %0: i32
 }
 
 //===----------------------------------------------------------------------===//
 // MinVersion
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @group_non_uniform_ballot_suitable_version
 func @group_non_uniform_ballot_suitable_version(%predicate: i1) -> vector<4xi32> attributes {
-  spv.target_env = #spv.target_env<V_1_4, [], [GroupNonUniformBallot], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.4, [GroupNonUniformBallot], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: spv.GroupNonUniformBallot "Workgroup"
   %0 = "test.convert_to_group_non_uniform_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
   return %0: vector<4xi32>
 }
 
 // CHECK-LABEL: @group_non_uniform_ballot_unsupported_version
 func @group_non_uniform_ballot_unsupported_version(%predicate: i1) -> vector<4xi32> attributes {
-  spv.target_env = #spv.target_env<V_1_1, [], [GroupNonUniformBallot], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.1, [GroupNonUniformBallot], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: test.convert_to_group_non_uniform_ballot_op
   %0 = "test.convert_to_group_non_uniform_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
   return %0: vector<4xi32>
 }
 
 //===----------------------------------------------------------------------===//
 // Capability
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @cmp_exchange_weak_missing_capability_kernel
 func @cmp_exchange_weak_missing_capability_kernel(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
-  spv.target_env = #spv.target_env<V_1_3, [], [AtomicStorage], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.3, [AtomicStorage], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: test.convert_to_atomic_compare_exchange_weak_op
   %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32)
   return %0: i32
 }
 
 // CHECK-LABEL: @cmp_exchange_weak_missing_capability_atomic_storage
 func @cmp_exchange_weak_missing_capability_atomic_storage(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes {
-  spv.target_env = #spv.target_env<V_1_3, [], [Kernel], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.3, [Kernel], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: test.convert_to_atomic_compare_exchange_weak_op
   %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32)
   return %0: i32
 }
 
 // CHECK-LABEL: @subgroup_ballot_missing_capability
 func @subgroup_ballot_missing_capability(%predicate: i1) -> vector<4xi32> attributes {
-  spv.target_env = #spv.target_env<V_1_4, [SPV_KHR_shader_ballot], [], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.4, [], [SPV_KHR_shader_ballot]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: test.convert_to_subgroup_ballot_op
   %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
   return %0: vector<4xi32>
 }
 
 // CHECK-LABEL: @bit_reverse_directly_implied_capability
 func @bit_reverse_directly_implied_capability(%operand: i32) -> i32 attributes {
-  spv.target_env = #spv.target_env<V_1_0, [], [Geometry], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [Geometry], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: spv.BitReverse
   %0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32)
   return %0: i32
 }
 
 // CHECK-LABEL: @bit_reverse_recursively_implied_capability
 func @bit_reverse_recursively_implied_capability(%operand: i32) -> i32 attributes {
-  spv.target_env = #spv.target_env<V_1_0, [], [GeometryPointSize], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [GeometryPointSize], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: spv.BitReverse
   %0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32)
   return %0: i32
 }
 
 //===----------------------------------------------------------------------===//
 // Extension
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @subgroup_ballot_suitable_extension
 func @subgroup_ballot_suitable_extension(%predicate: i1) -> vector<4xi32> attributes {
-  spv.target_env = #spv.target_env<V_1_4, [SPV_KHR_shader_ballot], [SubgroupBallotKHR], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], [SPV_KHR_shader_ballot]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: spv.SubgroupBallotKHR
   %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
   return %0: vector<4xi32>
 }
 
 // CHECK-LABEL: @subgroup_ballot_missing_extension
 func @subgroup_ballot_missing_extension(%predicate: i1) -> vector<4xi32> attributes {
-  spv.target_env = #spv.target_env<V_1_4, [], [SubgroupBallotKHR], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: test.convert_to_subgroup_ballot_op
   %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>)
   return %0: vector<4xi32>
 }
 
 // CHECK-LABEL: @module_suitable_extension1
 func @module_suitable_extension1() attributes {
-  spv.target_env = #spv.target_env<V_1_0, [SPV_KHR_vulkan_memory_model, SPV_EXT_physical_storage_buffer], [VulkanMemoryModel, PhysicalStorageBufferAddresses], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  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"
   "test.convert_to_module_op"() : () ->()
   return
 }
 
 // CHECK-LABEL: @module_suitable_extension2
 func @module_suitable_extension2() attributes {
-  spv.target_env = #spv.target_env<V_1_0, [SPV_KHR_vulkan_memory_model, SPV_KHR_physical_storage_buffer], [VulkanMemoryModel, PhysicalStorageBufferAddresses], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  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"
   "test.convert_to_module_op"() : () -> ()
   return
 }
 
 // CHECK-LABEL: @module_missing_extension_mm
 func @module_missing_extension_mm() attributes {
-  spv.target_env = #spv.target_env<V_1_0, [SPV_KHR_physical_storage_buffer], [VulkanMemoryModel, PhysicalStorageBufferAddresses], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  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>}>
 } {
   // CHECK: test.convert_to_module_op
   "test.convert_to_module_op"() : () -> ()
   return
 }
 
 // CHECK-LABEL: @module_missing_extension_am
 func @module_missing_extension_am() attributes {
-  spv.target_env = #spv.target_env<V_1_0, [SPV_KHR_vulkan_memory_model], [VulkanMemoryModel, PhysicalStorageBufferAddresses], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
   // CHECK: test.convert_to_module_op
   "test.convert_to_module_op"() : () -> ()
   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<V_1_5, [], [VulkanMemoryModel, PhysicalStorageBufferAddresses], {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
+  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"
   "test.convert_to_module_op"() : () -> ()
   return
 }