[mlir][spirv] Add a `spirv::InterfaceVarABIAttr`.

[denis13]

Summary

Add a proper dialect-specific attribute for interface variable ABI.

[denis13]

This patch is related to an issue https://bugs.llvm.org/show_bug.cgi?id=45392

[antiagainst]

Awesome! Thanks Denis! Just two nits. Feel free to submit after fixing them. :)

mlir/docs/Dialects/SPIR-V.md

892

integer-literal. same for the following

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

44

s/global/interface/

[denis13]

@antiagainst thanks for review!I'll fix it.

[denis13]

@mravishankar thanks for review!

mlir/docs/Dialects/SPIR-V.md

 ### Shader interface attributes
 
 The SPIR-V dialect defines [a few attributes][MlirSpirvAbi] for specifying these
 interfaces:
 
 *   `spv.entry_point_abi` is a struct attribute that should be attached to the
     entry function. It contains:
     *   `local_size` for specifying the local work group size for the dispatch.
-*   `spv.interface_var_abi` is a struct attribute that should be attached to
-    each operand and result of the entry function. It contains:
-    *   `descriptor_set` for specifying the descriptor set number for the
-        corresponding resource variable.
-    *   `binding` for specifying the binding number for the corresponding
-        resource variable.
-    *   `storage_class` for specifying the storage class for the corresponding
-        resource variable.
+*   `spv.interface_var_abi` is attribute that should be attached to each operand
+    and result of the entry function. It should be of `#spv.interface_var_abi`
+    attribute kind, which is defined as:
+
+```
+spv-storage-class     ::= `StorageBuffer` | ...
+spv-descriptor-set    ::= integer-literal
+spv-binding           ::= integer-literal
+spv-interface-var-abi ::= `#` `spv.interface_var_abi` `<(` spv-descriptor-set
+                          `,` spv-binding `)` (`,` spv-storage-class)? `>`
+```
+
+For example,
+
+```
+#spv.interface_var_abi<(0, 0), StorageBuffer>
+#spv.interface_var_abi<(0, 1)>
+```
+
+The attribute has a few fields:
+
+*   Descriptor set number for the corresponding resource variable.
+*   Binding number for the corresponding resource variable.
+*   Storage class for the corresponding resource variable.
 
 The SPIR-V dialect provides a [`LowerABIAttributesPass`][MlirSpirvPasses] for
 consuming these attributes and create SPIR-V module complying with the
 interface.
 
 ## Serialization and deserialization
 
 Although the main objective of the SPIR-V dialect is to act as a proper IR for

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

 //===- SPIRVAttributes.h - SPIR-V attribute declarations  -------*- C++ -*-===//
 //
 // 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 declares SPIR-V dialect specific attributes.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_DIALECT_SPIRV_SPIRVATTRIBUTES_H
 #define MLIR_DIALECT_SPIRV_SPIRVATTRIBUTES_H
 
+#include "mlir/Dialect/SPIRV/SPIRVTypes.h"
 #include "mlir/IR/Attributes.h"
 #include "mlir/Support/LLVM.h"
 
 namespace mlir {
 // Pull in SPIR-V attribute definitions for target and ABI.
 #include "mlir/Dialect/SPIRV/TargetAndABI.h.inc"
 
 namespace spirv {
 enum class Capability : uint32_t;
 enum class Extension;
 enum class Version : uint32_t;
 
 namespace detail {
+struct InterfaceVarABIAttributeStorage;
 struct TargetEnvAttributeStorage;
 struct VerCapExtAttributeStorage;
 } // namespace detail
 
 /// SPIR-V dialect-specific attribute kinds.
 namespace AttrKind {
 enum Kind {
-  TargetEnv = Attribute::FIRST_SPIRV_ATTR, /// Target environment
+  InterfaceVarABI = Attribute::FIRST_SPIRV_ATTR, /// Interface var ABI
+  TargetEnv,                                     /// Target environment
   VerCapExt, /// (version, extension, capability) triple
 };
 } // namespace AttrKind
 
+/// An attribute that specifies the information regarding the interface
+/// variable: descriptor set, binding, storage class.
+class InterfaceVarABIAttr
+    : public Attribute::AttrBase<InterfaceVarABIAttr, Attribute,
+                                 detail::InterfaceVarABIAttributeStorage> {
+public:
+  using Base::Base;
+
+  /// Gets a InterfaceVarABIAttr.
+  static InterfaceVarABIAttr get(uint32_t descirptorSet, uint32_t binding,
+                                 Optional<StorageClass> storageClass,
+                                 MLIRContext *context);
+  static InterfaceVarABIAttr get(IntegerAttr descriptorSet, IntegerAttr binding,
+                                 IntegerAttr storageClass);
+
+  /// Returns the attribute kind's name (without the 'spv.' prefix).
+  static StringRef getKindName();
+
+  /// Returns descriptor set.
+  uint32_t getDescriptorSet();
+
+  /// Returns binding.
+  uint32_t getBinding();
+
+  /// Returns `spirv::StorageClass`.
+  Optional<StorageClass> getStorageClass();
+
+  static bool kindof(unsigned kind) {
+    return kind == AttrKind::InterfaceVarABI;
+  }
+
+  static LogicalResult verifyConstructionInvariants(Location loc,
+                                                    IntegerAttr descriptorSet,
+                                                    IntegerAttr binding,
+                                                    IntegerAttr storageClass);
+};
+
 /// An attribute that specifies the SPIR-V (version, capabilities, extensions)
 /// triple.
 class VerCapExtAttr
     : public Attribute::AttrBase<VerCapExtAttr, Attribute,
                                  detail::VerCapExtAttributeStorage> {
 public:
   using Base::Base;
 

mlir/include/mlir/Dialect/SPIRV/TargetAndABI.td

 //
 //===----------------------------------------------------------------------===//
 
 #ifndef SPIRV_TARGET_AND_ABI
 #define SPIRV_TARGET_AND_ABI
 
 include "mlir/Dialect/SPIRV/SPIRVBase.td"
 
-// For arguments that eventually map to spv.globalVariable for the
-// shader interface, this attribute specifies the information regarding
-// the global variable:
-// 1) Descriptor Set.
-// 2) Binding number.
-// 3) Storage class.
-def SPV_InterfaceVarABIAttr : StructAttr<"InterfaceVarABIAttr", SPIRV_Dialect, [
-    StructFieldAttr<"descriptor_set", I32Attr>,
-    StructFieldAttr<"binding", I32Attr>,
-    StructFieldAttr<"storage_class", OptionalAttr<SPV_StorageClassAttr>>
-]>;
-
 // For entry functions, this attribute specifies information related to entry
 // points in the generated SPIR-V module:
 // 1) WorkGroup Size.
 def SPV_EntryPointABIAttr : StructAttr<"EntryPointABIAttr", SPIRV_Dialect, [
     StructFieldAttr<"local_size", I32ElementsAttr>
 ]>;
 
 def SPV_ExtensionArrayAttr : TypedArrayAttrBase<

mlir/lib/Dialect/SPIRV/SPIRVAttributes.cpp

 #include "mlir/Dialect/SPIRV/TargetAndABI.cpp.inc"
 
 //===----------------------------------------------------------------------===//
 // Attribute storage classes
 //===----------------------------------------------------------------------===//
 
 namespace spirv {
 namespace detail {
+
+struct InterfaceVarABIAttributeStorage : public AttributeStorage {
+  using KeyTy = std::tuple<Attribute, Attribute, Attribute>;
+
+  InterfaceVarABIAttributeStorage(Attribute descriptorSet, Attribute binding,
+                                  Attribute storageClass)
+      : descriptorSet(descriptorSet), binding(binding),
+        storageClass(storageClass) {}
+
+  bool operator==(const KeyTy &key) const {
+    return std::get<0>(key) == descriptorSet && std::get<1>(key) == binding &&
+           std::get<2>(key) == storageClass;
+  }
+
+  static InterfaceVarABIAttributeStorage *
+  construct(AttributeStorageAllocator &allocator, const KeyTy &key) {
+    return new (allocator.allocate<InterfaceVarABIAttributeStorage>())
+        InterfaceVarABIAttributeStorage(std::get<0>(key), std::get<1>(key),
+                                        std::get<2>(key));
+  }
+
+  Attribute descriptorSet;
+  Attribute binding;
+  Attribute storageClass;
+};
+
 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 {
struct TargetEnvAttributeStorage : public AttributeStorage {
   Attribute triple;
   Attribute limits;
 };
 } // namespace detail
 } // namespace spirv
 } // namespace mlir
 
 //===----------------------------------------------------------------------===//
+// InterfaceVarABIAttr
+//===----------------------------------------------------------------------===//
+
+spirv::InterfaceVarABIAttr
+spirv::InterfaceVarABIAttr::get(uint32_t descriptorSet, uint32_t binding,
+                                Optional<spirv::StorageClass> storageClass,
+                                MLIRContext *context) {
+  Builder b(context);
+  auto descriptorSetAttr = b.getI32IntegerAttr(descriptorSet);
+  auto bindingAttr = b.getI32IntegerAttr(binding);
+  auto storageClassAttr =
+      storageClass ? b.getI32IntegerAttr(static_cast<uint32_t>(*storageClass))
+                   : IntegerAttr();
+  return get(descriptorSetAttr, bindingAttr, storageClassAttr);
+}
+
+spirv::InterfaceVarABIAttr
+spirv::InterfaceVarABIAttr::get(IntegerAttr descriptorSet, IntegerAttr binding,
+                                IntegerAttr storageClass) {
+  assert(descriptorSet && binding);
+  MLIRContext *context = descriptorSet.getContext();
+  return Base::get(context, spirv::AttrKind::InterfaceVarABI, descriptorSet,
+                   binding, storageClass);
+}
+
+StringRef spirv::InterfaceVarABIAttr::getKindName() {
+  return "interface_var_abi";
+}
+
+uint32_t spirv::InterfaceVarABIAttr::getBinding() {
+  return getImpl()->binding.cast<IntegerAttr>().getInt();
+}
+
+uint32_t spirv::InterfaceVarABIAttr::getDescriptorSet() {
+  return getImpl()->descriptorSet.cast<IntegerAttr>().getInt();
+}
+
+Optional<spirv::StorageClass> spirv::InterfaceVarABIAttr::getStorageClass() {
+  if (getImpl()->storageClass)
+    return static_cast<spirv::StorageClass>(
+        getImpl()->storageClass.cast<IntegerAttr>().getValue().getZExtValue());
+  return llvm::None;
+}
+
+LogicalResult spirv::InterfaceVarABIAttr::verifyConstructionInvariants(
+    Location loc, IntegerAttr descriptorSet, IntegerAttr binding,
+    IntegerAttr storageClass) {
+  if (!descriptorSet.getType().isSignlessInteger(32))
+    return emitError(loc, "expected 32-bit integer for descriptor set");
+
+  if (!binding.getType().isSignlessInteger(32))
+    return emitError(loc, "expected 32-bit integer for binding");
+
+  if (storageClass) {
+    if (auto storageClassAttr = storageClass.cast<IntegerAttr>()) {
+      auto storageClassValue =
+          spirv::symbolizeStorageClass(storageClassAttr.getInt());
+      if (!storageClassValue)
+        return emitError(loc, "unknown storage class");
+    } else {
+      return emitError(loc, "expected valid storage class");
+    }
+  }
+
+  return success();
+}
+
+//===----------------------------------------------------------------------===//
 // VerCapExtAttr
 //===----------------------------------------------------------------------===//
 
 spirv::VerCapExtAttr spirv::VerCapExtAttr::get(
     spirv::Version version, ArrayRef<spirv::Capability> capabilities,
     ArrayRef<spirv::Extension> extensions, MLIRContext *context) {
   Builder b(context);
 

mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp

 //===----------------------------------------------------------------------===//
 // SPIR-V Dialect
 //===----------------------------------------------------------------------===//
 
 SPIRVDialect::SPIRVDialect(MLIRContext *context)
     : Dialect(getDialectNamespace(), context) {
   addTypes<ArrayType, ImageType, PointerType, RuntimeArrayType, StructType>();
 
-  addAttributes<TargetEnvAttr, VerCapExtAttr>();
+  addAttributes<InterfaceVarABIAttr, 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::InterfaceVarABIAttr.
+static Attribute parseInterfaceVarABIAttr(DialectAsmParser &parser) {
+  if (parser.parseLess())
+    return {};
+
+  Builder &builder = parser.getBuilder();
+
+  if (parser.parseLParen())
+    return {};
+
+  IntegerAttr descriptorSetAttr;
+  {
+    auto loc = parser.getCurrentLocation();
+    uint32_t descriptorSet = 0;
+    auto descriptorSetParseResult = parser.parseOptionalInteger(descriptorSet);
+
+    if (!descriptorSetParseResult.hasValue() ||
+        failed(*descriptorSetParseResult)) {
+      parser.emitError(loc, "missing descriptor set");
+      return {};
+    }
+    descriptorSetAttr = builder.getI32IntegerAttr(descriptorSet);
+  }
+
+  if (parser.parseComma())
+    return {};
+
+  IntegerAttr bindingAttr;
+  {
+    auto loc = parser.getCurrentLocation();
+    uint32_t binding = 0;
+    auto bindingParseResult = parser.parseOptionalInteger(binding);
+
+    if (!bindingParseResult.hasValue() || failed(*bindingParseResult)) {
+      parser.emitError(loc, "missing binding");
+      return {};
+    }
+    bindingAttr = builder.getI32IntegerAttr(binding);
+  }
+
+  if (parser.parseRParen())
+    return {};
+
+  IntegerAttr storageClassAttr;
+  {
+    if (succeeded(parser.parseOptionalComma())) {
+      auto loc = parser.getCurrentLocation();
+      StringRef storageClass;
+      if (parser.parseKeyword(&storageClass))
+        return {};
+
+      if (auto storageClassSymbol =
+              spirv::symbolizeStorageClass(storageClass)) {
+        storageClassAttr = builder.getI32IntegerAttr(
+            static_cast<uint32_t>(*storageClassSymbol));
+      } else {
+        parser.emitError(loc, "unknown storage class: ") << storageClass;
+        return {};
+      }
+    }
+  }
+
+  if (parser.parseGreater())
+    return {};
+
+  return spirv::InterfaceVarABIAttr::get(descriptorSetAttr, bindingAttr,
+                                         storageClassAttr);
+}
+
 static Attribute parseVerCapExtAttr(DialectAsmParser &parser) {
   if (parser.parseLess())
     return {};
 
   Builder &builder = parser.getBuilder();
 
   IntegerAttr versionAttr;
   {
Attribute SPIRVDialect::parseAttribute(DialectAsmParser &parser,
   StringRef attrKind;
   if (parser.parseKeyword(&attrKind))
     return {};
 
   if (attrKind == spirv::TargetEnvAttr::getKindName())
     return parseTargetEnvAttr(parser);
   if (attrKind == spirv::VerCapExtAttr::getKindName())
     return parseVerCapExtAttr(parser);
+  if (attrKind == spirv::InterfaceVarABIAttr::getKindName())
+    return parseInterfaceVarABIAttr(parser);
 
   parser.emitError(parser.getNameLoc(), "unknown SPIR-V attribute kind: ")
       << attrKind;
   return {};
 }
 
 //===----------------------------------------------------------------------===//
 // Attribute Printing

 }
 
 static void print(spirv::TargetEnvAttr targetEnv, DialectAsmPrinter &printer) {
   printer << spirv::TargetEnvAttr::getKindName() << "<#spv.";
   print(targetEnv.getTripleAttr(), printer);
   printer << ", " << targetEnv.getResourceLimits() << ">";
 }
 
+static void print(spirv::InterfaceVarABIAttr interfaceVarABIAttr,
+                  DialectAsmPrinter &printer) {
+  printer << spirv::InterfaceVarABIAttr::getKindName() << "<("
+          << interfaceVarABIAttr.getDescriptorSet() << ", "
+          << interfaceVarABIAttr.getBinding() << ")";
+  auto storageClass = interfaceVarABIAttr.getStorageClass();
+  if (storageClass)
+    printer << ", " << spirv::stringifyStorageClass(*storageClass);
+  printer << ">";
+}
+
 void SPIRVDialect::printAttribute(Attribute attr,
                                   DialectAsmPrinter &printer) const {
   if (auto targetEnv = attr.dyn_cast<TargetEnvAttr>())
     print(targetEnv, printer);
   else if (auto vceAttr = attr.dyn_cast<VerCapExtAttr>())
     print(vceAttr, printer);
+  else if (auto interfaceVarABIAttr = attr.dyn_cast<InterfaceVarABIAttr>())
+    print(interfaceVarABIAttr, printer);
   else
     llvm_unreachable("unhandled SPIR-V attribute kind");
 }
 
 //===----------------------------------------------------------------------===//
 // Constant
 //===----------------------------------------------------------------------===//
 
static LogicalResult verifyRegionAttribute(Location loc, Type valueType,
 
   if (symbol != spirv::getInterfaceVarABIAttrName())
     return emitError(loc, "found unsupported '")
            << symbol << "' attribute on region argument";
 
   auto varABIAttr = attr.dyn_cast<spirv::InterfaceVarABIAttr>();
   if (!varABIAttr)
     return emitError(loc, "'")
-           << symbol
-           << "' attribute must be a dictionary attribute containing two or "
-              "three 32-bit integer attributes: 'descriptor_set', 'binding', "
-              "and optional 'storage_class'";
-  if (varABIAttr.storage_class() && !valueType.isIntOrIndexOrFloat())
+           << symbol << "' must be a spirv::InterfaceVarABIAttr";
+
+  if (varABIAttr.getStorageClass() && !valueType.isIntOrIndexOrFloat())
     return emitError(loc, "'") << symbol
                                << "' attribute cannot specify storage class "
                                   "when attaching to a non-scalar value";
 
   return success();
 }
 
 LogicalResult SPIRVDialect::verifyRegionArgAttribute(Operation *op,

mlir/lib/Dialect/SPIRV/TargetAndABI.cpp

 StringRef spirv::getInterfaceVarABIAttrName() {
   return "spv.interface_var_abi";
 }
 
 spirv::InterfaceVarABIAttr
 spirv::getInterfaceVarABIAttr(unsigned descriptorSet, unsigned binding,
                               Optional<spirv::StorageClass> storageClass,
                               MLIRContext *context) {
-  Type i32Type = IntegerType::get(32, context);
-  auto scAttr =
-      storageClass
-          ? IntegerAttr::get(i32Type, static_cast<int64_t>(*storageClass))
-          : IntegerAttr();
-  return spirv::InterfaceVarABIAttr::get(
-      IntegerAttr::get(i32Type, descriptorSet),
-      IntegerAttr::get(i32Type, binding), scAttr, context);
+  return spirv::InterfaceVarABIAttr::get(descriptorSet, binding, storageClass,
+                                         context);
 }
 
 StringRef spirv::getEntryPointABIAttrName() { return "spv.entry_point_abi"; }
 
 spirv::EntryPointABIAttr
 spirv::getEntryPointABIAttr(ArrayRef<int32_t> localSize, MLIRContext *context) {
   assert(localSize.size() == 3);
   return spirv::EntryPointABIAttr::get(

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

createGlobalVarForEntryPointArgument(OpBuilder &builder, spirv::FuncOp funcOp,
   std::string varName =
       funcOp.getName().str() + "_arg_" + std::to_string(argIndex);
 
   // Get the type of variable. If this is a scalar/vector type and has an ABI
   // info create a variable of type !spv.ptr<!spv.struct<elementType>>. If not
   // it must already be a !spv.ptr<!spv.struct<...>>.
   auto varType = funcOp.getType().getInput(argIndex);
   if (varType.cast<spirv::SPIRVType>().isScalarOrVector()) {
-    auto storageClass =
-        static_cast<spirv::StorageClass>(abiInfo.storage_class().getInt());
+    auto storageClass = abiInfo.getStorageClass();
+    if (!storageClass)
+      return nullptr;
     varType =
-        spirv::PointerType::get(spirv::StructType::get(varType), storageClass);
+        spirv::PointerType::get(spirv::StructType::get(varType), *storageClass);
   }
   auto varPtrType = varType.cast<spirv::PointerType>();
   auto varPointeeType = varPtrType.getPointeeType().cast<spirv::StructType>();
 
   // Set the offset information.
   varPointeeType =
       VulkanLayoutUtils::decorateType(varPointeeType).cast<spirv::StructType>();
   varType =
       spirv::PointerType::get(varPointeeType, varPtrType.getStorageClass());
 
   return builder.create<spirv::GlobalVariableOp>(
-      funcOp.getLoc(), varType, varName, abiInfo.descriptor_set().getInt(),
-      abiInfo.binding().getInt());
+      funcOp.getLoc(), varType, varName, abiInfo.getDescriptorSet(),
+      abiInfo.getBinding());
 }
 
 /// Gets the global variables that need to be specified as interface variable
 /// with an spv.EntryPointOp. Traverses the body of a entry function to do so.
 static LogicalResult
 getInterfaceVariables(spirv::FuncOp funcOp,
                       SmallVectorImpl<Attribute> &interfaceVars) {
   auto module = funcOp.getParentOfType<spirv::ModuleOp>();

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

module attributes {
   }
 
   // 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, stride=4> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32{{[}][}]}}
-    // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32, stride=4> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32{{[}][}]}}
-    // CHECK-SAME: [[ARG2:%.*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32, stride=4> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 2 : i32, descriptor_set = 0 : i32{{[}][}]}}
-    // CHECK-SAME: [[ARG3:%.*]]: i32 {spv.interface_var_abi = {binding = 3 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
-    // CHECK-SAME: [[ARG4:%.*]]: i32 {spv.interface_var_abi = {binding = 4 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
-    // CHECK-SAME: [[ARG5:%.*]]: i32 {spv.interface_var_abi = {binding = 5 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
-    // CHECK-SAME: [[ARG6:%.*]]: i32 {spv.interface_var_abi = {binding = 6 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
+    // CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32, stride=4> [0]>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 0)>}
+    // CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32, stride=4> [0]>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
+    // CHECK-SAME: [[ARG2:%.*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32, stride=4> [0]>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 2)>}
+    // CHECK-SAME: [[ARG3:%.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 3), StorageBuffer>}
+    // CHECK-SAME: [[ARG4:%.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 4), StorageBuffer>}
+    // CHECK-SAME: [[ARG5:%.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 5), StorageBuffer>}
+    // CHECK-SAME: [[ARG6:%.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 6), StorageBuffer>}
     gpu.func @load_store_kernel(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>, %arg3: index, %arg4: index, %arg5: index, %arg6: index)
       attributes {gpu.kernel, spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESSWORKGROUPID:%.*]] = spv._address_of [[WORKGROUPIDVAR]]
       // CHECK: [[WORKGROUPID:%.*]] = spv.Load "Input" [[ADDRESSWORKGROUPID]]
       // CHECK: [[WORKGROUPIDX:%.*]] = spv.CompositeExtract [[WORKGROUPID]]{{\[}}0 : i32{{\]}}
       // CHECK: [[ADDRESSLOCALINVOCATIONID:%.*]] = spv._address_of [[LOCALINVOCATIONIDVAR]]
       // CHECK: [[LOCALINVOCATIONID:%.*]] = spv.Load "Input" [[ADDRESSLOCALINVOCATIONID]]
       // CHECK: [[LOCALINVOCATIONIDX:%.*]] = spv.CompositeExtract [[LOCALINVOCATIONID]]{{\[}}0 : i32{{\]}}

mlir/test/Conversion/GPUToSPIRV/simple.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -split-input-file -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {gpu.container_module} {
   gpu.module @kernels {
     // CHECK:       spv.module Logical GLSL450 {
     // CHECK-LABEL: spv.func @basic_module_structure
-    // CHECK-SAME: {{%.*}}: f32 {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32{{[}][}]}}
-    // CHECK-SAME: {{%.*}}: !spv.ptr<!spv.struct<!spv.array<12 x f32, stride=4> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : i32{{[}][}]}}
+    // CHECK-SAME: {{%.*}}: f32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>}
+    // CHECK-SAME: {{%.*}}: !spv.ptr<!spv.struct<!spv.array<12 x f32, stride=4> [0]>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
     // CHECK-SAME: spv.entry_point_abi = {local_size = dense<[32, 4, 1]> : vector<3xi32>}
     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
     }
   }
 

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

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
 // CHECK-LABEL: spv.module
 spv.module Logical GLSL450 {
   // CHECK-DAG:    spv.globalVariable [[VAR0:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<f32 [0]>, StorageBuffer>
   // CHECK-DAG:    spv.globalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<12 x f32, stride=4> [0]>, StorageBuffer>
   // CHECK:    spv.func [[FN:@.*]]()
   spv.func @kernel(
     %arg0: f32
-           {spv.interface_var_abi = {binding = 0 : i32,
-                                     descriptor_set = 0 : i32,
-                                     storage_class = 12 : i32}},
+           {spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>},
     %arg1: !spv.ptr<!spv.struct<!spv.array<12 x f32>>, StorageBuffer>
-           {spv.interface_var_abi = {binding = 1 : i32,
-                                     descriptor_set = 0 : i32}}) "None"
+           {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}) "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
 } // end spv.module
 
 } // end module

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

spv.module Logical GLSL450 {
   // CHECK-DAG: spv.globalVariable [[VAR2:@.*]] bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32, stride=4>, stride=16> [0]>, StorageBuffer>
   // CHECK-DAG: spv.globalVariable [[VAR3:@.*]] bind(0, 3) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
   // CHECK-DAG: spv.globalVariable [[VAR4:@.*]] bind(0, 4) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
   // CHECK-DAG: spv.globalVariable [[VAR5:@.*]] bind(0, 5) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
   // CHECK-DAG: spv.globalVariable [[VAR6:@.*]] bind(0, 6) : !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
   // CHECK: spv.func [[FN:@.*]]()
   spv.func @load_store_kernel(
     %arg0: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32>>>, StorageBuffer>
-    {spv.interface_var_abi = {binding = 0 : i32,
-                              descriptor_set = 0 : i32}},
+    {spv.interface_var_abi = #spv.interface_var_abi<(0, 0)>},
     %arg1: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32>>>, StorageBuffer>
-    {spv.interface_var_abi = {binding = 1 : i32,
-                              descriptor_set = 0 : i32}},
+    {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>},
     %arg2: !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32>>>, StorageBuffer>
-    {spv.interface_var_abi = {binding = 2 : i32,
-                              descriptor_set = 0 : i32}},
+    {spv.interface_var_abi = #spv.interface_var_abi<(0, 2)>},
     %arg3: i32
-    {spv.interface_var_abi = {binding = 3 : i32,
-                              descriptor_set = 0 : i32,
-                              storage_class = 12 : i32}},
+    {spv.interface_var_abi = #spv.interface_var_abi<(0, 3), StorageBuffer>},
     %arg4: i32
-    {spv.interface_var_abi = {binding = 4 : i32,
-                              descriptor_set = 0 : i32,
-                              storage_class = 12 : i32}},
+    {spv.interface_var_abi = #spv.interface_var_abi<(0, 4), StorageBuffer>},
     %arg5: i32
-    {spv.interface_var_abi = {binding = 5 : i32,
-                              descriptor_set = 0 : i32,
-                              storage_class = 12 : i32}},
+    {spv.interface_var_abi = #spv.interface_var_abi<(0, 5), StorageBuffer>},
     %arg6: i32
-    {spv.interface_var_abi = {binding = 6 : i32,
-                              descriptor_set = 0 : i32,
-                              storage_class = 12 : i32}}) "None"
+    {spv.interface_var_abi = #spv.interface_var_abi<(0, 6), StorageBuffer>}) "None"
   attributes  {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
     // CHECK: [[ADDRESSARG6:%.*]] = spv._address_of [[VAR6]]
     // CHECK: [[CONST6:%.*]] = spv.constant 0 : i32
     // CHECK: [[ARG6PTR:%.*]] = spv.AccessChain [[ADDRESSARG6]]{{\[}}[[CONST6]]
     // CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG6PTR]]
     // CHECK: [[ADDRESSARG5:%.*]] = spv._address_of [[VAR5]]
     // CHECK: [[CONST5:%.*]] = spv.constant 0 : i32
     // CHECK: [[ARG5PTR:%.*]] = spv.AccessChain [[ADDRESSARG5]]{{\[}}[[CONST5]]

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

 } { return }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.interface_var_abi
 //===----------------------------------------------------------------------===//
 
-// expected-error @+1 {{'spv.interface_var_abi' attribute must be a dictionary attribute containing two or three 32-bit integer attributes: 'descriptor_set', 'binding', and optional 'storage_class'}}
+// expected-error @+1 {{'spv.interface_var_abi' must be a spirv::InterfaceVarABIAttr}}
 func @interface_var(
   %arg0 : f32 {spv.interface_var_abi = 64}
 ) { return }
 
 // -----
 
-// expected-error @+1 {{'spv.interface_var_abi' attribute must be a dictionary attribute containing two or three 32-bit integer attributes: 'descriptor_set', 'binding', and optional 'storage_class'}}
 func @interface_var(
-  %arg0 : f32 {spv.interface_var_abi = {binding = 0: i32}}
+// expected-error @+1 {{missing descriptor set}}
+  %arg0 : f32 {spv.interface_var_abi = #spv.interface_var_abi<()>}
 ) { return }
 
 // -----
 
-// CHECK: {spv.interface_var_abi = {binding = 0 : i32, descriptor_set = 0 : i32, storage_class = 12 : i32}}
 func @interface_var(
-  %arg0 : f32 {spv.interface_var_abi = {binding = 0 : i32,
-                                        descriptor_set = 0 : i32,
-                                        storage_class = 12 : i32}}
+// expected-error @+1 {{missing binding}}
+  %arg0 : f32 {spv.interface_var_abi = #spv.interface_var_abi<(1,)>}
+) { return }
+
+// -----
+
+func @interface_var(
+// expected-error @+1 {{unknown storage class: }}
+  %arg0 : f32 {spv.interface_var_abi = #spv.interface_var_abi<(1,2), Foo>}
+) { return }
+
+// -----
+
+// CHECK: {spv.interface_var_abi = #spv.interface_var_abi<(0, 1), Uniform>}
+func @interface_var(
+    %arg0 : f32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 1), Uniform>}
+) { return }
+
+// -----
+
+// CHECK: {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
+func @interface_var(
+    %arg0 : f32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
 ) { return }
 
 // -----
 
 // expected-error @+1 {{'spv.interface_var_abi' attribute cannot specify storage class when attaching to a non-scalar value}}
 func @interface_var(
-  %arg0 : memref<4xf32> {spv.interface_var_abi = {binding = 0 : i32,
-                                                  descriptor_set = 0 : i32,
-                                                  storage_class = 12 : i32}}
+  %arg0 : memref<4xf32> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1), Uniform>}
 ) { return }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.target_env
 //===----------------------------------------------------------------------===//