[mlir][spirv] Improve stride support in array types

[antiagainst]

Summary

This commit added stride support in runtime array types. It also
adjusted the assembly form for the stride from [N] to stride=N.
This makes the IR more readable, especially for the cases where
one mix array types and struct types.

[denis13]

LGTM! Looks more pretty than it was, thanks!

mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp

548–549

What do you think about

if (!type.getArrayStride())

555

What do you think about

if (!type.getArrayStride())

[denis13]

mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp

548–549

Sorry, I meant

if (type.getArrayStride())

555

Sorry I meant

if (type.getArrayStide())

[antiagainst]

mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp

555

Oops, good catch! I meant to reuse the stride value in the following line. Fixed. :)

[mravishankar]

Sorry, just getting to this for review.

See comments below.

mlir/include/mlir/Dialect/SPIRV/SPIRVTypes.h

157

Seems to be a mismatch here. elementCount and stride should be size_t. Of course on many devices unsigned is enough, but when possible we might need to support objects of large sizes. Maybe better to make both size_t?

mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp

231

This makes it too verbose. Why change the parsing/printing?

[mravishankar]

I know this is already committed, but I just wanted to reopen this (posted comments here). Didnt get around to reviewing it till now

[antiagainst]

mlir/include/mlir/Dialect/SPIRV/SPIRVTypes.h

157

SGTM to be future proof. I used unsigned here to be consistent; better to change them together. (The serialization/deserialization side we need to be careful; the array stride is directly embedded as a literal there. I wonder whether all the drivers support non-32-bit cases nicely though. ;-P)

mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp

231

It makes the IR more readable. Otherwise if you have structs containing arrays, everything (struct member offsets, array strides) are represented as [N]. It's really hard to differentiate. And I typically think [N] is more naturally for indexing; thus good for offsets.

[bondhugula]

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

61

You don't need the %<num> = part. In any case, please don't use numbers in such places.

[antiagainst]

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

61

Good point! Thanks for the catch! Fixed with https://github.com/llvm/llvm-project/commit/f83d502febb54b2cf3866ebea247380cd6263f13.

mlir/docs/Dialects/SPIR-V.md

 This corresponds to SPIR-V [array type][ArrayType]. Its syntax is
 
 ```
 element-type ::= integer-type
                | floating-point-type
                | vector-type
                | spirv-type
 
-array-type ::= `!spv.array<` integer-literal `x` element-type `>`
+array-type ::= `!spv.array` `<` integer-literal `x` element-type
+               (`,` `stride` `=` integer-literal)? `>`
 ```
 
 For example,
 
 ```mlir
 !spv.array<4 x i32>
+!spv.array<4 x i32, stride = 4>
 !spv.array<16 x vector<4 x f32>>
 ```
 
 ### Image type
 
 This corresponds to SPIR-V [image type][ImageType]. Its syntax is
 
 ```

 !spv.ptr<vector<4 x f32>, Uniform>
 ```
 
 ### Runtime array type
 
 This corresponds to SPIR-V [runtime array type][RuntimeArrayType]. Its syntax is
 
 ```
-runtime-array-type ::= `!spv.rtarray<` element-type `>`
+runtime-array-type ::= `!spv.rtarray` `<` element-type (`,` `stride` `=` integer-literal)? `>`
 ```
 
 For example,
 
 ```mlir
 !spv.rtarray<i32>
+!spv.rtarray<i32, stride=4>
 !spv.rtarray<vector<4 x f32>>
 ```
 
 ### Struct type
 
 This corresponds to SPIR-V [struct type][StructType]. Its syntax is
 
 ```

mlir/include/mlir/Dialect/SPIRV/LayoutUtils.h

 #ifndef MLIR_DIALECT_SPIRV_LAYOUTUTILS_H_
 #define MLIR_DIALECT_SPIRV_LAYOUTUTILS_H_
 
 #include <cstdint>
 
 namespace mlir {
 class Type;
 class VectorType;
+
 namespace spirv {
-class StructType;
 class ArrayType;
+class RuntimeArrayType;
+class StructType;
 } // namespace spirv
 
 /// According to the Vulkan spec "14.5.4. Offset and Stride Assignment":
 /// "There are different alignment requirements depending on the specific
 /// resources and on the features enabled on the device."
 ///
 /// There are 3 types of alignment: scalar, base, extended.
 /// See the spec for details.

 /// According to the SPIR-V spec:
 /// "The ArrayStride, MatrixStride, and Offset decorations must be large
 /// enough to hold the size of the objects they affect (that is, specifying
 /// overlap is invalid)."
 class VulkanLayoutUtils {
 public:
   using Size = uint64_t;
 
-  /// Returns a new StructType with layout info. Assigns the type size in bytes
-  /// to the `size`. Assigns the type alignment in bytes to the `alignment`.
-  static spirv::StructType decorateType(spirv::StructType structType,
-                                        Size &size, Size &alignment);
+  /// Returns a new StructType with layout decoration.
+  static spirv::StructType decorateType(spirv::StructType structType);
+
   /// Checks whether a type is legal in terms of Vulkan layout info
   /// decoration. A type is dynamically illegal if it's a composite type in the
   /// StorageBuffer, PhysicalStorageBuffer, Uniform, and PushConstant Storage
   /// Classes without layout information.
   static bool isLegalType(Type type);
 
 private:
+  /// Returns a new type with layout decoration. Assigns the type size in bytes
+  /// to the `size`. Assigns the type alignment in bytes to the `alignment`.
   static Type decorateType(Type type, Size &size, Size &alignment);
+
   static Type decorateType(VectorType vectorType, Size &size, Size &alignment);
   static Type decorateType(spirv::ArrayType arrayType, Size &size,
                            Size &alignment);
+  static Type decorateType(spirv::RuntimeArrayType arrayType, Size &alignment);
+  static spirv::StructType decorateType(spirv::StructType structType,
+                                        Size &size, Size &alignment);
+
   /// Calculates the alignment for the given scalar type.
   static Size getScalarTypeAlignment(Type scalarType);
 };
 
 } // namespace mlir
 
 #endif // MLIR_DIALECT_SPIRV_LAYOUTUTILS_H_

mlir/include/mlir/Dialect/SPIRV/SPIRVTypes.h

void getCapabilities(SPIRVType::CapabilityArrayRefVector &capabilities,
                        Optional<spirv::StorageClass> storage = llvm::None);
 };
 
 // SPIR-V array type
 class ArrayType : public Type::TypeBase<ArrayType, CompositeType,
                                         detail::ArrayTypeStorage> {
 public:
   using Base::Base;
-  // Zero layout specifies that is no layout
-  using LayoutInfo = uint64_t;
 
   static bool kindof(unsigned kind) { return kind == TypeKind::Array; }
 
   static ArrayType get(Type elementType, unsigned elementCount);
 
+  /// Returns an array type with the given stride in bytes.
   static ArrayType get(Type elementType, unsigned elementCount,
-                       LayoutInfo layoutInfo);
+                       unsigned stride);
 
   unsigned getNumElements() const;
 
   Type getElementType() const;
 
-  bool hasLayout() const;
-
-  uint64_t getArrayStride() const;
+  /// Returns the array stride in bytes. 0 means no stride decorated on this
+  /// type.
+  unsigned getArrayStride() const;
 
   void getExtensions(SPIRVType::ExtensionArrayRefVector &extensions,
                      Optional<spirv::StorageClass> storage = llvm::None);
   void getCapabilities(SPIRVType::CapabilityArrayRefVector &capabilities,
                        Optional<spirv::StorageClass> storage = llvm::None);
 };
 
 // SPIR-V image type
: public Type::TypeBase<RuntimeArrayType, SPIRVType,
                             detail::RuntimeArrayTypeStorage> {
 public:
   using Base::Base;
 
   static bool kindof(unsigned kind) { return kind == TypeKind::RuntimeArray; }
 
   static RuntimeArrayType get(Type elementType);
 
+  /// Returns a runtime array type with the given stride in bytes.
+  static RuntimeArrayType get(Type elementType, unsigned stride);
+
   Type getElementType() const;
 
+  /// Returns the array stride in bytes. 0 means no stride decorated on this
+  /// type.
+  unsigned getArrayStride() const;
+
   void getExtensions(SPIRVType::ExtensionArrayRefVector &extensions,
                      Optional<spirv::StorageClass> storage = llvm::None);
   void getCapabilities(SPIRVType::CapabilityArrayRefVector &capabilities,
                        Optional<spirv::StorageClass> storage = llvm::None);
 };
 
 // SPIR-V struct type
 class StructType : public Type::TypeBase<StructType, CompositeType,

mlir/lib/Dialect/SPIRV/LayoutUtils.cpp

 //===----------------------------------------------------------------------===//
 #include "mlir/Dialect/SPIRV/LayoutUtils.h"
 #include "mlir/Dialect/SPIRV/SPIRVDialect.h"
 #include "mlir/Dialect/SPIRV/SPIRVTypes.h"
 
 using namespace mlir;
 
 spirv::StructType
+VulkanLayoutUtils::decorateType(spirv::StructType structType) {
+  Size size = 0;
+  Size alignment = 1;
+  return decorateType(structType, size, alignment);
+}
+
+spirv::StructType
 VulkanLayoutUtils::decorateType(spirv::StructType structType,
                                 VulkanLayoutUtils::Size &size,
                                 VulkanLayoutUtils::Size &alignment) {
   if (structType.getNumElements() == 0) {
     return structType;
   }
 
   SmallVector<Type, 4> memberTypes;
-  SmallVector<VulkanLayoutUtils::Size, 4> layoutInfo;
+  SmallVector<Size, 4> layoutInfo;
   SmallVector<spirv::StructType::MemberDecorationInfo, 4> memberDecorations;
 
-  VulkanLayoutUtils::Size structMemberOffset = 0;
-  VulkanLayoutUtils::Size maxMemberAlignment = 1;
+  Size structMemberOffset = 0;
+  Size maxMemberAlignment = 1;
 
   for (uint32_t i = 0, e = structType.getNumElements(); i < e; ++i) {
-    VulkanLayoutUtils::Size memberSize = 0;
-    VulkanLayoutUtils::Size memberAlignment = 1;
+    Size memberSize = 0;
+    Size memberAlignment = 1;
 
-    auto memberType = VulkanLayoutUtils::decorateType(
-        structType.getElementType(i), memberSize, memberAlignment);
+    auto memberType =
+        decorateType(structType.getElementType(i), memberSize, memberAlignment);
     structMemberOffset = llvm::alignTo(structMemberOffset, memberAlignment);
     memberTypes.push_back(memberType);
     layoutInfo.push_back(structMemberOffset);
+    // If the member's size is the max value, it must be the last member and it
+    // must be a runtime array.
+    assert(memberSize != std::numeric_limits<Size>().max() ||
+           (i + 1 == e &&
+            structType.getElementType(i).isa<spirv::RuntimeArrayType>()));
     // According to the Vulkan spec:
     // "A structure has a base alignment equal to the largest base alignment of
     // any of its members."
     structMemberOffset += memberSize;
     maxMemberAlignment = std::max(maxMemberAlignment, memberAlignment);
   }
 
   // According to the Vulkan spec:
   // "The Offset decoration of a member must not place it between the end of a
   // structure or an array and the next multiple of the alignment of that
   // structure or array."
   size = llvm::alignTo(structMemberOffset, maxMemberAlignment);
   alignment = maxMemberAlignment;
   structType.getMemberDecorations(memberDecorations);
   return spirv::StructType::get(memberTypes, layoutInfo, memberDecorations);
 }
 
 Type VulkanLayoutUtils::decorateType(Type type, VulkanLayoutUtils::Size &size,
                                      VulkanLayoutUtils::Size &alignment) {
   if (type.isa<spirv::ScalarType>()) {
-    alignment = VulkanLayoutUtils::getScalarTypeAlignment(type);
+    alignment = getScalarTypeAlignment(type);
     // Vulkan spec does not specify any padding for a scalar type.
     size = alignment;
     return type;
   }
 
   switch (type.getKind()) {
   case spirv::TypeKind::Struct:
-    return VulkanLayoutUtils::decorateType(type.cast<spirv::StructType>(), size,
-                                           alignment);
+    return decorateType(type.cast<spirv::StructType>(), size, alignment);
   case spirv::TypeKind::Array:
-    return VulkanLayoutUtils::decorateType(type.cast<spirv::ArrayType>(), size,
-                                           alignment);
+    return decorateType(type.cast<spirv::ArrayType>(), size, alignment);
   case StandardTypes::Vector:
-    return VulkanLayoutUtils::decorateType(type.cast<VectorType>(), size,
-                                           alignment);
+    return decorateType(type.cast<VectorType>(), size, alignment);
+  case spirv::TypeKind::RuntimeArray:
+    size = std::numeric_limits<Size>().max();
+    return decorateType(type.cast<spirv::RuntimeArrayType>(), alignment);
   default:
     llvm_unreachable("unhandled SPIR-V type");
   }
 }
 
 Type VulkanLayoutUtils::decorateType(VectorType vectorType,
                                      VulkanLayoutUtils::Size &size,
                                      VulkanLayoutUtils::Size &alignment) {
   const auto numElements = vectorType.getNumElements();
   auto elementType = vectorType.getElementType();
-  VulkanLayoutUtils::Size elementSize = 0;
-  VulkanLayoutUtils::Size elementAlignment = 1;
+  Size elementSize = 0;
+  Size elementAlignment = 1;
 
-  auto memberType = VulkanLayoutUtils::decorateType(elementType, elementSize,
-                                                    elementAlignment);
+  auto memberType = decorateType(elementType, elementSize, elementAlignment);
   // According to the Vulkan spec:
   // 1. "A two-component vector has a base alignment equal to twice its scalar
   // alignment."
   // 2. "A three- or four-component vector has a base alignment equal to four
   // times its scalar alignment."
   size = elementSize * numElements;
   alignment = numElements == 2 ? elementAlignment * 2 : elementAlignment * 4;
   return VectorType::get(numElements, memberType);
 }
 
 Type VulkanLayoutUtils::decorateType(spirv::ArrayType arrayType,
                                      VulkanLayoutUtils::Size &size,
                                      VulkanLayoutUtils::Size &alignment) {
   const auto numElements = arrayType.getNumElements();
   auto elementType = arrayType.getElementType();
-  spirv::ArrayType::LayoutInfo elementSize = 0;
-  VulkanLayoutUtils::Size elementAlignment = 1;
+  Size elementSize = 0;
+  Size elementAlignment = 1;
 
-  auto memberType = VulkanLayoutUtils::decorateType(elementType, elementSize,
-                                                    elementAlignment);
+  auto memberType = decorateType(elementType, elementSize, elementAlignment);
   // According to the Vulkan spec:
   // "An array has a base alignment equal to the base alignment of its element
   // type."
   size = elementSize * numElements;
   alignment = elementAlignment;
   return spirv::ArrayType::get(memberType, numElements, elementSize);
 }
 
+Type VulkanLayoutUtils::decorateType(spirv::RuntimeArrayType arrayType,
+                                     VulkanLayoutUtils::Size &alignment) {
+  auto elementType = arrayType.getElementType();
+  Size elementSize = 0;
+
+  auto memberType = decorateType(elementType, elementSize, alignment);
+  return spirv::RuntimeArrayType::get(memberType, elementSize);
+}
+
 VulkanLayoutUtils::Size
 VulkanLayoutUtils::getScalarTypeAlignment(Type scalarType) {
   // According to the Vulkan spec:
   // 1. "A scalar of size N has a scalar alignment of N."
   // 2. "A scalar has a base alignment equal to its scalar alignment."
   // 3. "A scalar, vector or matrix type has an extended alignment equal to its
   // base alignment."
   auto bitWidth = scalarType.getIntOrFloatBitWidth();

mlir/lib/Dialect/SPIRV/SPIRVDialect.cpp

 template <typename ValTy>
 static Optional<ValTy> parseAndVerify(SPIRVDialect const &dialect,
                                       DialectAsmParser &parser);
 template <>
 Optional<Type> parseAndVerify<Type>(SPIRVDialect const &dialect,
                                     DialectAsmParser &parser);
 
 template <>
-Optional<uint64_t> parseAndVerify<uint64_t>(SPIRVDialect const &dialect,
+Optional<unsigned> parseAndVerify<unsigned>(SPIRVDialect const &dialect,
                                             DialectAsmParser &parser);
 
 static Type parseAndVerifyType(SPIRVDialect const &dialect,
                                DialectAsmParser &parser) {
   Type type;
   llvm::SMLoc typeLoc = parser.getCurrentLocation();
   if (parser.parseType(type))
     return Type();
if (auto t = type.dyn_cast<FloatType>()) {
     parser.emitError(typeLoc, "cannot use ")
         << type << " to compose SPIR-V types";
     return Type();
   }
 
   return type;
 }
 
+/// Parses an optional `, stride = N` assembly segment. If no parsing failure
+/// occurs, writes `N` to `stride` if existing and writes 0 to `stride` if
+/// missing.
+static LogicalResult parseOptionalArrayStride(const SPIRVDialect &dialect,
+                                              DialectAsmParser &parser,
+                                              unsigned &stride) {
+  if (failed(parser.parseOptionalComma())) {
+    stride = 0;
+    return success();
+  }
+
+  if (parser.parseKeyword("stride") || parser.parseEqual())
+    return failure();
+
+  llvm::SMLoc strideLoc = parser.getCurrentLocation();
+  Optional<unsigned> optStride = parseAndVerify<unsigned>(dialect, parser);
+  if (!optStride)
+    return failure();
+
+  if (!(stride = optStride.getValue())) {
+    parser.emitError(strideLoc, "ArrayStride must be greater than zero");
+    return failure();
+  }
+  return success();
+}
+
 // element-type ::= integer-type
 //                | floating-point-type
 //                | vector-type
 //                | spirv-type
 //
-// array-type ::= `!spv.array<` integer-literal `x` element-type
-//                (`[` integer-literal `]`)? `>`
+// array-type ::= `!spv.array` `<` integer-literal `x` element-type
+//                (`,` `stride` `=` integer-literal)? `>`
 static Type parseArrayType(SPIRVDialect const &dialect,
                            DialectAsmParser &parser) {
   if (parser.parseLess())
     return Type();
 
   SmallVector<int64_t, 1> countDims;
   llvm::SMLoc countLoc = parser.getCurrentLocation();
   if (parser.parseDimensionList(countDims, /*allowDynamic=*/false))
if (count == 0) {
     parser.emitError(countLoc, "expected array length greater than 0");
     return Type();
   }
 
   Type elementType = parseAndVerifyType(dialect, parser);
   if (!elementType)
     return Type();
 
-  ArrayType::LayoutInfo layoutInfo = 0;
-  if (succeeded(parser.parseOptionalLSquare())) {
-    llvm::SMLoc layoutLoc = parser.getCurrentLocation();
-    auto layout = parseAndVerify<ArrayType::LayoutInfo>(dialect, parser);
-    if (!layout)
-      return Type();
-
-    if (!(layoutInfo = layout.getValue())) {
-      parser.emitError(layoutLoc, "ArrayStride must be greater than zero");
-      return Type();
-    }
-
-    if (parser.parseRSquare())
-      return Type();
-  }
+  unsigned stride = 0;
+  if (failed(parseOptionalArrayStride(dialect, parser, stride)))
 
   if (parser.parseGreater())
     return Type();
-  return ArrayType::get(elementType, count, layoutInfo);
+  return ArrayType::get(elementType, count, stride);
 }
 
 // TODO(ravishankarm) : Reorder methods to be utilities first and parse*Type
 // methods in alphabetical order
 //
 // storage-class ::= `UniformConstant`
 //                 | `Uniform`
 //                 | `Workgroup`
parser.emitError(storageClassLoc, "unknown storage class: ")
         << storageClassSpec;
     return Type();
   }
   if (parser.parseGreater())
     return Type();
   return PointerType::get(pointeeType, *storageClass);
 }
 
-// runtime-array-type ::= `!spv.rtarray<` element-type `>`
+// runtime-array-type ::= `!spv.rtarray` `<` element-type
+//                        (`,` `stride` `=` integer-literal)? `>`
 static Type parseRuntimeArrayType(SPIRVDialect const &dialect,
                                   DialectAsmParser &parser) {
   if (parser.parseLess())
     return Type();
 
   Type elementType = parseAndVerifyType(dialect, parser);
   if (!elementType)
     return Type();
 
+  unsigned stride = 0;
+  if (failed(parseOptionalArrayStride(dialect, parser, stride)))
+    return Type();
+
   if (parser.parseGreater())
     return Type();
-  return RuntimeArrayType::get(elementType);
+  return RuntimeArrayType::get(elementType, stride);
 }
 
 // Specialize this function to parse each of the parameters that define an
 // ImageType. By default it assumes this is an enum type.
 template <typename ValTy>
 static Optional<ValTy> parseAndVerify(SPIRVDialect const &dialect,
                                       DialectAsmParser &parser) {
   StringRef enumSpec;
static Optional<IntTy> parseAndVerifyInteger(SPIRVDialect const &dialect,
                                              DialectAsmParser &parser) {
   IntTy offsetVal = std::numeric_limits<IntTy>::max();
   if (parser.parseInteger(offsetVal))
     return llvm::None;
   return offsetVal;
 }
 
 template <>
-Optional<uint64_t> parseAndVerify<uint64_t>(SPIRVDialect const &dialect,
+Optional<unsigned> parseAndVerify<unsigned>(SPIRVDialect const &dialect,
                                             DialectAsmParser &parser) {
-  return parseAndVerifyInteger<uint64_t>(dialect, parser);
+  return parseAndVerifyInteger<unsigned>(dialect, parser);
 }
 
 namespace {
 // Functor object to parse a comma separated list of specs. The function
 // parseAndVerify does the actual parsing and verification of individual
 // elements. This is a functor since parsing the last element of the list
 // (termination condition) needs partial specialization.
 template <typename ParseType, typename... Args>

 }
 
 //===----------------------------------------------------------------------===//
 // Type Printing
 //===----------------------------------------------------------------------===//
 
 static void print(ArrayType type, DialectAsmPrinter &os) {
   os << "array<" << type.getNumElements() << " x " << type.getElementType();
-  if (type.hasLayout()) {
-    os << " [" << type.getArrayStride() << "]";
-  }
+  if (unsigned stride = type.getArrayStride())
+    os << ", stride=" << stride;
   os << ">";
 }
 
 static void print(RuntimeArrayType type, DialectAsmPrinter &os) {
-  os << "rtarray<" << type.getElementType() << ">";
+  os << "rtarray<" << type.getElementType();
+  if (unsigned stride = type.getArrayStride())
+    os << ", stride=" << stride;
+  os << ">";
 }
 
 static void print(PointerType type, DialectAsmPrinter &os) {
   os << "ptr<" << type.getPointeeType() << ", "
      << stringifyStorageClass(type.getStorageClass()) << ">";
 }
 
 static void print(ImageType type, DialectAsmPrinter &os) {

mlir/lib/Dialect/SPIRV/SPIRVTypes.cpp

spirv::getRecursiveImpliedCapabilities(Capability cap) {
   return allCaps.takeVector();
 }
 
 //===----------------------------------------------------------------------===//
 // ArrayType
 //===----------------------------------------------------------------------===//
 
 struct spirv::detail::ArrayTypeStorage : public TypeStorage {
-  using KeyTy = std::tuple<Type, unsigned, ArrayType::LayoutInfo>;
+  using KeyTy = std::tuple<Type, unsigned, unsigned>;
 
   static ArrayTypeStorage *construct(TypeStorageAllocator &allocator,
                                      const KeyTy &key) {
     return new (allocator.allocate<ArrayTypeStorage>()) ArrayTypeStorage(key);
   }
 
   bool operator==(const KeyTy &key) const {
-    return key == KeyTy(elementType, getSubclassData(), layoutInfo);
+    return key == KeyTy(elementType, getSubclassData(), stride);
   }
 
   ArrayTypeStorage(const KeyTy &key)
       : TypeStorage(std::get<1>(key)), elementType(std::get<0>(key)),
-        layoutInfo(std::get<2>(key)) {}
+        stride(std::get<2>(key)) {}
 
   Type elementType;
-  ArrayType::LayoutInfo layoutInfo;
+  unsigned stride;
 };
 
 ArrayType ArrayType::get(Type elementType, unsigned elementCount) {
   assert(elementCount && "ArrayType needs at least one element");
   return Base::get(elementType.getContext(), TypeKind::Array, elementType,
-                   elementCount, 0);
+                   elementCount, /*stride=*/0);
 }
 
 ArrayType ArrayType::get(Type elementType, unsigned elementCount,
-                         ArrayType::LayoutInfo layoutInfo) {
+                         unsigned stride) {
   assert(elementCount && "ArrayType needs at least one element");
   return Base::get(elementType.getContext(), TypeKind::Array, elementType,
-                   elementCount, layoutInfo);
+                   elementCount, stride);
 }
 
 unsigned ArrayType::getNumElements() const {
   return getImpl()->getSubclassData();
 }
 
 Type ArrayType::getElementType() const { return getImpl()->elementType; }
 
-// ArrayStride must be greater than zero
-bool ArrayType::hasLayout() const { return getImpl()->layoutInfo; }
-
-uint64_t ArrayType::getArrayStride() const { return getImpl()->layoutInfo; }
+unsigned ArrayType::getArrayStride() const { return getImpl()->stride; }
 
 void ArrayType::getExtensions(SPIRVType::ExtensionArrayRefVector &extensions,
                               Optional<StorageClass> storage) {
   getElementType().cast<SPIRVType>().getExtensions(extensions, storage);
 }
 
 void ArrayType::getCapabilities(
     SPIRVType::CapabilityArrayRefVector &capabilities,

 void CompositeType::getExtensions(
     SPIRVType::ExtensionArrayRefVector &extensions,
     Optional<StorageClass> storage) {
   switch (getKind()) {
   case spirv::TypeKind::Array:
     cast<ArrayType>().getExtensions(extensions, storage);
     break;
   case spirv::TypeKind::RuntimeArray:
-    cast<ArrayType>().getExtensions(extensions, storage);
+    cast<RuntimeArrayType>().getExtensions(extensions, storage);
     break;
   case spirv::TypeKind::Struct:
     cast<StructType>().getExtensions(extensions, storage);
     break;
   case StandardTypes::Vector:
     cast<VectorType>().getElementType().cast<ScalarType>().getExtensions(
         extensions, storage);
     break;
   default:
     llvm_unreachable("invalid composite type");
   }
 }
 
 void CompositeType::getCapabilities(
     SPIRVType::CapabilityArrayRefVector &capabilities,
     Optional<StorageClass> storage) {
   switch (getKind()) {
   case spirv::TypeKind::Array:
     cast<ArrayType>().getCapabilities(capabilities, storage);
     break;
   case spirv::TypeKind::RuntimeArray:
-    cast<ArrayType>().getCapabilities(capabilities, storage);
+    cast<RuntimeArrayType>().getCapabilities(capabilities, storage);
     break;
   case spirv::TypeKind::Struct:
     cast<StructType>().getCapabilities(capabilities, storage);
     break;
   case StandardTypes::Vector:
     cast<VectorType>().getElementType().cast<ScalarType>().getCapabilities(
         capabilities, storage);
     break;
if (auto scCaps = spirv::getCapabilities(getStorageClass()))
     capabilities.push_back(*scCaps);
 }
 
 //===----------------------------------------------------------------------===//
 // RuntimeArrayType
 //===----------------------------------------------------------------------===//
 
 struct spirv::detail::RuntimeArrayTypeStorage : public TypeStorage {
-  using KeyTy = Type;
+  using KeyTy = std::pair<Type, unsigned>;
 
   static RuntimeArrayTypeStorage *construct(TypeStorageAllocator &allocator,
                                             const KeyTy &key) {
     return new (allocator.allocate<RuntimeArrayTypeStorage>())
         RuntimeArrayTypeStorage(key);
   }
 
-  bool operator==(const KeyTy &key) const { return elementType == key; }
+  bool operator==(const KeyTy &key) const {
+    return key == KeyTy(elementType, getSubclassData());
+  }
 
-  RuntimeArrayTypeStorage(const KeyTy &key) : elementType(key) {}
+  RuntimeArrayTypeStorage(const KeyTy &key)
+      : TypeStorage(key.second), elementType(key.first) {}
 
   Type elementType;
 };
 
 RuntimeArrayType RuntimeArrayType::get(Type elementType) {
   return Base::get(elementType.getContext(), TypeKind::RuntimeArray,
-                   elementType);
+                   elementType, /*stride=*/0);
+}
+
+RuntimeArrayType RuntimeArrayType::get(Type elementType, unsigned stride) {
+  return Base::get(elementType.getContext(), TypeKind::RuntimeArray,
+                   elementType, stride);
 }
 
 Type RuntimeArrayType::getElementType() const { return getImpl()->elementType; }
 
+unsigned RuntimeArrayType::getArrayStride() const {
+  return getImpl()->getSubclassData();
+}
+
 void RuntimeArrayType::getExtensions(
     SPIRVType::ExtensionArrayRefVector &extensions,
     Optional<StorageClass> storage) {
   getElementType().cast<SPIRVType>().getExtensions(extensions, storage);
 }
 
 void RuntimeArrayType::getCapabilities(
     SPIRVType::CapabilityArrayRefVector &capabilities,

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

if (operands.size() != 2) {
     return emitError(unknownLoc, "OpTypeRuntimeArray must have two operands");
   }
   Type memberType = getType(operands[1]);
   if (!memberType) {
     return emitError(unknownLoc,
                      "OpTypeRuntimeArray references undefined <id> ")
            << operands[1];
   }
-  typeMap[operands[0]] = spirv::RuntimeArrayType::get(memberType);
+  typeMap[operands[0]] = spirv::RuntimeArrayType::get(
+      memberType, typeDecorations.lookup(operands[0]));
   return success();
 }
 
 LogicalResult Deserializer::processStructType(ArrayRef<uint32_t> operands) {
   if (operands.empty()) {
     return emitError(unknownLoc, "OpTypeStruct must have at least result <id>");
   }
   if (operands.size() == 1) {

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

LogicalResult Serializer::processName(uint32_t resultID, StringRef name) {
   }
   return encodeInstructionInto(names, spirv::Opcode::OpName, nameOperands);
 }
 
 namespace {
 template <>
 LogicalResult Serializer::processTypeDecoration<spirv::ArrayType>(
     Location loc, spirv::ArrayType type, uint32_t resultID) {
-  if (type.hasLayout()) {
+  if (unsigned stride = type.getArrayStride()) {
     // OpDecorate %arrayTypeSSA ArrayStride strideLiteral
-    return emitDecoration(resultID, spirv::Decoration::ArrayStride,
-                          {static_cast<uint32_t>(type.getArrayStride())});
+    return emitDecoration(resultID, spirv::Decoration::ArrayStride, {stride});
+  }
+  return success();
+}
+
+template <>
+LogicalResult Serializer::processTypeDecoration<spirv::RuntimeArrayType>(
+    Location Loc, spirv::RuntimeArrayType type, uint32_t resultID) {
+  if (unsigned stride = type.getArrayStride()) {
+    // OpDecorate %arrayTypeSSA ArrayStride strideLiteral
+    return emitDecoration(resultID, spirv::Decoration::ArrayStride, {stride});
   }
   return success();
 }
 
 LogicalResult
 Serializer::processMemberDecoration(uint32_t structID, uint32_t memberIndex,
                                     spirv::Decoration decorationType,
                                     ArrayRef<uint32_t> values) {
Serializer::prepareBasicType(Location loc, Type type, uint32_t resultID,
   }
 
   if (auto runtimeArrayType = type.dyn_cast<spirv::RuntimeArrayType>()) {
     uint32_t elementTypeID = 0;
     if (failed(processType(loc, runtimeArrayType.getElementType(),
                            elementTypeID))) {
       return failure();
     }
-    operands.push_back(elementTypeID);
     typeEnum = spirv::Opcode::OpTypeRuntimeArray;
-    return success();
+    operands.push_back(elementTypeID);
+    return processTypeDecoration(loc, runtimeArrayType, resultID);
   }
 
   if (auto structType = type.dyn_cast<spirv::StructType>()) {
     bool hasLayout = structType.hasLayout();
     for (auto elementIndex :
          llvm::seq<uint32_t>(0, structType.getNumElements())) {
       uint32_t elementTypeID = 0;
       if (failed(processType(loc, structType.getElementType(elementIndex),

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

 namespace {
 class SPIRVGlobalVariableOpLayoutInfoDecoration
     : public OpRewritePattern<spirv::GlobalVariableOp> {
 public:
   using OpRewritePattern<spirv::GlobalVariableOp>::OpRewritePattern;
 
   LogicalResult matchAndRewrite(spirv::GlobalVariableOp op,
                                 PatternRewriter &rewriter) const override {
-    spirv::StructType::LayoutInfo structSize = 0;
-    VulkanLayoutUtils::Size structAlignment = 1;
     SmallVector<NamedAttribute, 4> globalVarAttrs;
 
     auto ptrType = op.type().cast<spirv::PointerType>();
     auto structType = VulkanLayoutUtils::decorateType(
-        ptrType.getPointeeType().cast<spirv::StructType>(), structSize,
-        structAlignment);
+        ptrType.getPointeeType().cast<spirv::StructType>());
     auto decoratedType =
         spirv::PointerType::get(structType, ptrType.getStorageClass());
 
     // Save all named attributes except "type" attribute.
     for (const auto &attr : op.getAttrs()) {
       if (attr.first == "type") {
         continue;
       }

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

auto storageClass =
         static_cast<spirv::StorageClass>(abiInfo.storage_class().getInt());
     varType =
         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.
-  VulkanLayoutUtils::Size size = 0, alignment = 0;
   varPointeeType =
-      VulkanLayoutUtils::decorateType(varPointeeType, size, alignment)
-          .cast<spirv::StructType>();
+      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());
 }
 

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 [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 [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 [4]> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 2 : i32, descriptor_set = 0 : i32{{[}][}]}}
+    // 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{{[}][}]}}
     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]]

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 [4]> [0]>, StorageBuffer> {spv.interface_var_abi = {binding = 1 : i32, descriptor_set = 0 : 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: 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/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 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.globalVariable @kernel_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<12 x f32, stride=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>
+      %0 = spv._address_of @kernel_arg_0 : !spv.ptr<!spv.struct<!spv.array<12 x f32, stride=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>
+      %3 = spv._address_of @kernel_arg_0 : !spv.ptr<!spv.struct<!spv.array<12 x f32, stride=4> [0]>, StorageBuffer>
+      %4 = spv.AccessChain %0[%2, %2] : !spv.ptr<!spv.struct<!spv.array<12 x f32, stride=4> [0]>, StorageBuffer>
       %5 = spv.Load "StorageBuffer" %4 : f32
       spv.Return
     }
     spv.EntryPoint "GLCompute" @kernel
     spv.ExecutionMode @kernel "LocalSize", 1, 1, 1
   }
   gpu.module @kernels {
     gpu.func @kernel(%arg0: memref<12xf32>) kernel {

mlir/test/Conversion/StandardToSPIRV/std-to-spirv.mlir

func @constant() {
   // CHECK: spv.constant 42 : i32
   %1 = constant 42 : i32
   // CHECK: spv.constant 5.000000e-01 : f32
   %2 = constant 0.5 : f32
   // CHECK: spv.constant dense<[2, 3]> : vector<2xi32>
   %3 = constant dense<[2, 3]> : vector<2xi32>
   // CHECK: spv.constant 1 : i32
   %4 = constant 1 : index
-  // CHECK: spv.constant dense<1> : tensor<6xi32> : !spv.array<6 x i32 [4]>
+  // CHECK: spv.constant dense<1> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
   %5 = constant dense<1> : tensor<2x3xi32>
-  // CHECK: spv.constant dense<1.000000e+00> : tensor<6xf32> : !spv.array<6 x f32 [4]>
+  // CHECK: spv.constant dense<1.000000e+00> : tensor<6xf32> : !spv.array<6 x f32, stride=4>
   %6 = constant dense<1.0> : tensor<2x3xf32>
-  // CHECK: spv.constant dense<{{\[}}1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf32> : !spv.array<6 x f32 [4]>
+  // CHECK: spv.constant dense<{{\[}}1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf32> : !spv.array<6 x f32, stride=4>
   %7 = constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32>
-  // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32 [4]>
+  // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
   %8 = constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32>
-  // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32 [4]>
+  // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
   %9 =  constant dense<[[1, 2], [3, 4], [5, 6]]> : tensor<3x2xi32>
-  // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32 [4]>
+  // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
   %10 =  constant dense<[1, 2, 3, 4, 5, 6]> : tensor<6xi32>
   return
 }
 
 // CHECK-LABEL: @constant_16bit
 func @constant_16bit() {
   // CHECK: spv.constant 4 : i16
   %0 = constant 4 : i16
   // CHECK: spv.constant 5.000000e+00 : f16
   %1 = constant 5.0 : f16
   // CHECK: spv.constant dense<[2, 3]> : vector<2xi16>
   %2 = constant dense<[2, 3]> : vector<2xi16>
-  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf16> : !spv.array<5 x f16 [2]>
+  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf16> : !spv.array<5 x f16, stride=2>
   %3 = constant dense<4.0> : tensor<5xf16>
   return
 }
 
 // CHECK-LABEL: @constant_64bit
 func @constant_64bit() {
   // CHECK: spv.constant 4 : i64
   %0 = constant 4 : i64
   // CHECK: spv.constant 5.000000e+00 : f64
   %1 = constant 5.0 : f64
   // CHECK: spv.constant dense<[2, 3]> : vector<2xi64>
   %2 = constant dense<[2, 3]> : vector<2xi64>
-  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf64> : !spv.array<5 x f64 [8]>
+  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf64> : !spv.array<5 x f64, stride=8>
   %3 = constant dense<4.0> : tensor<5xf64>
   return
 }
 
 } // end module
 
 // -----
 
 // Check that constants are converted to 32-bit when no special capability.
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
 // CHECK-LABEL: @constant_16bit
 func @constant_16bit() {
   // CHECK: spv.constant 4 : i32
   %0 = constant 4 : i16
   // CHECK: spv.constant 5.000000e+00 : f32
   %1 = constant 5.0 : f16
   // CHECK: spv.constant dense<[2, 3]> : vector<2xi32>
   %2 = constant dense<[2, 3]> : vector<2xi16>
-  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32 [4]>
+  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4>
   %3 = constant dense<4.0> : tensor<5xf16>
-  // CHECK: spv.constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32 [4]>
+  // CHECK: spv.constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4>
   %4 = constant dense<[[1.0, 2.0], [3.0, 4.0]]> : tensor<2x2xf16>
   return
 }
 
 // CHECK-LABEL: @constant_64bit
 func @constant_64bit() {
   // CHECK: spv.constant 4 : i32
   %0 = constant 4 : i64
   // CHECK: spv.constant 5.000000e+00 : f32
   %1 = constant 5.0 : f64
   // CHECK: spv.constant dense<[2, 3]> : vector<2xi32>
   %2 = constant dense<[2, 3]> : vector<2xi64>
-  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32 [4]>
+  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4>
   %3 = constant dense<4.0> : tensor<5xf64>
-  // CHECK: spv.constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32 [4]>
+  // CHECK: spv.constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4>
   %4 = constant dense<[[1.0, 2.0], [3.0, 4.0]]> : tensor<2x2xf16>
   return
 }
 
 // CHECK-LABEL: @corner_cases
 func @corner_cases() {
   // CHECK: %{{.*}} = spv.constant -1 : i32
   %0 = constant 4294967295  : i64 // 2^32 - 1
func @select(%arg0 : i32, %arg1 : i32) {
   return
 }
 
 //===----------------------------------------------------------------------===//
 // std load/store ops
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @load_store_zero_rank_float
-// CHECK: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<1 x f32 [4]> [0]>, StorageBuffer>,
-// CHECK: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<1 x f32 [4]> [0]>, StorageBuffer>)
+// CHECK: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<1 x f32, stride=4> [0]>, StorageBuffer>,
+// CHECK: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<1 x f32, stride=4> [0]>, StorageBuffer>)
 func @load_store_zero_rank_float(%arg0: memref<f32>, %arg1: memref<f32>) {
   //      CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32
   //      CHECK: spv.AccessChain [[ARG0]][
   // CHECK-SAME: [[ZERO1]], [[ZERO1]]
   // CHECK-SAME: ] :
   //      CHECK: spv.Load "StorageBuffer" %{{.*}} : f32
   %0 = load %arg0[] : memref<f32>
   //      CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32
   //      CHECK: spv.AccessChain [[ARG1]][
   // CHECK-SAME: [[ZERO2]], [[ZERO2]]
   // CHECK-SAME: ] :
   //      CHECK: spv.Store "StorageBuffer" %{{.*}} : f32
   store %0, %arg1[] : memref<f32>
   return
 }
 
 // CHECK-LABEL: @load_store_zero_rank_int
-// CHECK: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<1 x i32 [4]> [0]>, StorageBuffer>,
-// CHECK: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<1 x i32 [4]> [0]>, StorageBuffer>)
+// CHECK: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>,
+// CHECK: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>)
 func @load_store_zero_rank_int(%arg0: memref<i32>, %arg1: memref<i32>) {
   //      CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32
   //      CHECK: spv.AccessChain [[ARG0]][
   // CHECK-SAME: [[ZERO1]], [[ZERO1]]
   // CHECK-SAME: ] :
   //      CHECK: spv.Load "StorageBuffer" %{{.*}} : i32
   %0 = load %arg0[] : memref<i32>
   //      CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32

mlir/test/Conversion/StandardToSPIRV/std-types-to-spirv.mlir

 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
 // CHECK-LABEL: spv.func @memref_8bit_StorageBuffer
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i32 [4]> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i32, stride=4> [0]>, StorageBuffer>
 func @memref_8bit_StorageBuffer(%arg0: memref<16xi8, 0>) { return }
 
 // CHECK-LABEL: spv.func @memref_8bit_Uniform
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x si32 [4]> [0]>, Uniform>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x si32, stride=4> [0]>, Uniform>
 func @memref_8bit_Uniform(%arg0: memref<16xsi8, 4>) { return }
 
 // CHECK-LABEL: spv.func @memref_8bit_PushConstant
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x ui32 [4]> [0]>, PushConstant>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x ui32, stride=4> [0]>, PushConstant>
 func @memref_8bit_PushConstant(%arg0: memref<16xui8, 7>) { return }
 
 // CHECK-LABEL: spv.func @memref_16bit_StorageBuffer
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i32 [4]> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i32, stride=4> [0]>, StorageBuffer>
 func @memref_16bit_StorageBuffer(%arg0: memref<16xi16, 0>) { return }
 
 // CHECK-LABEL: spv.func @memref_16bit_Uniform
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x si32 [4]> [0]>, Uniform>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x si32, stride=4> [0]>, Uniform>
 func @memref_16bit_Uniform(%arg0: memref<16xsi16, 4>) { return }
 
 // CHECK-LABEL: spv.func @memref_16bit_PushConstant
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x ui32 [4]> [0]>, PushConstant>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x ui32, stride=4> [0]>, PushConstant>
 func @memref_16bit_PushConstant(%arg0: memref<16xui16, 7>) { return }
 
 // CHECK-LABEL: spv.func @memref_16bit_Input
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f32 [4]> [0]>, Input>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f32, stride=4> [0]>, Input>
 func @memref_16bit_Input(%arg3: memref<16xf16, 9>) { return }
 
 // CHECK-LABEL: spv.func @memref_16bit_Output
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f32 [4]> [0]>, Output>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f32, stride=4> [0]>, Output>
 func @memref_16bit_Output(%arg4: memref<16xf16, 10>) { return }
 
 } // end module
 
 // -----
 
 // Check that using non-32-bit scalar types in interface storage classes
 // requires special capability and extension: keep as-is when the capability
 // and extension is available.
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [StoragePushConstant8, StoragePushConstant16],
              [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
 // CHECK-LABEL: spv.func @memref_8bit_PushConstant
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i8 [1]> [0]>, PushConstant>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i8, stride=1> [0]>, PushConstant>
 func @memref_8bit_PushConstant(%arg0: memref<16xi8, 7>) { return }
 
 // CHECK-LABEL: spv.func @memref_16bit_PushConstant
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i16 [2]> [0]>, PushConstant>
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f16 [2]> [0]>, PushConstant>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i16, stride=2> [0]>, PushConstant>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f16, stride=2> [0]>, PushConstant>
 func @memref_16bit_PushConstant(
   %arg0: memref<16xi16, 7>,
   %arg1: memref<16xf16, 7>
 ) { return }
 
 } // end module
 
 // -----
 
 // Check that using non-32-bit scalar types in interface storage classes
 // requires special capability and extension: keep as-is when the capability
 // and extension is available.
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [StorageBuffer8BitAccess, StorageBuffer16BitAccess],
              [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
 // CHECK-LABEL: spv.func @memref_8bit_StorageBuffer
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i8 [1]> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i8, stride=1> [0]>, StorageBuffer>
 func @memref_8bit_StorageBuffer(%arg0: memref<16xi8, 0>) { return }
 
 // CHECK-LABEL: spv.func @memref_16bit_StorageBuffer
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i16 [2]> [0]>, StorageBuffer>
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f16 [2]> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i16, stride=2> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f16, stride=2> [0]>, StorageBuffer>
 func @memref_16bit_StorageBuffer(
   %arg0: memref<16xi16, 0>,
   %arg1: memref<16xf16, 0>
 ) { return }
 
 } // end module
 
 // -----
 
 // Check that using non-32-bit scalar types in interface storage classes
 // requires special capability and extension: keep as-is when the capability
 // and extension is available.
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [UniformAndStorageBuffer8BitAccess, StorageUniform16],
              [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
 // CHECK-LABEL: spv.func @memref_8bit_Uniform
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i8 [1]> [0]>, Uniform>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i8, stride=1> [0]>, Uniform>
 func @memref_8bit_Uniform(%arg0: memref<16xi8, 4>) { return }
 
 // CHECK-LABEL: spv.func @memref_16bit_Uniform
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i16 [2]> [0]>, Uniform>
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f16 [2]> [0]>, Uniform>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i16, stride=2> [0]>, Uniform>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f16, stride=2> [0]>, Uniform>
 func @memref_16bit_Uniform(
   %arg0: memref<16xi16, 4>,
   %arg1: memref<16xf16, 4>
 ) { return }
 
 } // end module
 
 // -----
 
 // Check that using non-32-bit scalar types in interface storage classes
 // requires special capability and extension: keep as-is when the capability
 // and extension is available.
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [StorageInputOutput16], [SPV_KHR_16bit_storage]>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
 // CHECK-LABEL: spv.func @memref_16bit_Input
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f16 [2]> [0]>, Input>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x f16, stride=2> [0]>, Input>
 func @memref_16bit_Input(%arg3: memref<16xf16, 9>) { return }
 
 // CHECK-LABEL: spv.func @memref_16bit_Output
-// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i16 [2]> [0]>, Output>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<16 x i16, stride=2> [0]>, Output>
 func @memref_16bit_Output(%arg4: memref<16xi16, 10>) { return }
 
 } // end module
 
 // -----
 
 // Check that memref offset and strides affect the array size.
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [StorageBuffer16BitAccess], [SPV_KHR_16bit_storage]>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
 // CHECK-LABEL: spv.func @memref_offset_strides
 func @memref_offset_strides(
-// CHECK-SAME: !spv.array<64 x f32 [4]> [0]>, StorageBuffer>
-// CHECK-SAME: !spv.array<72 x f32 [4]> [0]>, StorageBuffer>
-// CHECK-SAME: !spv.array<256 x f32 [4]> [0]>, StorageBuffer>
-// CHECK-SAME: !spv.array<64 x f32 [4]> [0]>, StorageBuffer>
-// CHECK-SAME: !spv.array<88 x f32 [4]> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.array<64 x f32, stride=4> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.array<72 x f32, stride=4> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.array<256 x f32, stride=4> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.array<64 x f32, stride=4> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.array<88 x f32, stride=4> [0]>, StorageBuffer>
   %arg0: memref<16x4xf32, offset: 0, strides: [4, 1]>,  // tightly packed; row major
   %arg1: memref<16x4xf32, offset: 8, strides: [4, 1]>,  // offset 8
   %arg2: memref<16x4xf32, offset: 0, strides: [16, 1]>, // pad 12 after each row
   %arg3: memref<16x4xf32, offset: 0, strides: [1, 16]>, // tightly packed; col major
   %arg4: memref<16x4xf32, offset: 0, strides: [1, 22]>, // pad 4 after each col
 
-// CHECK-SAME: !spv.array<64 x f16 [2]> [0]>, StorageBuffer>
-// CHECK-SAME: !spv.array<72 x f16 [2]> [0]>, StorageBuffer>
-// CHECK-SAME: !spv.array<256 x f16 [2]> [0]>, StorageBuffer>
-// CHECK-SAME: !spv.array<64 x f16 [2]> [0]>, StorageBuffer>
-// CHECK-SAME: !spv.array<88 x f16 [2]> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.array<64 x f16, stride=2> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.array<72 x f16, stride=2> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.array<256 x f16, stride=2> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.array<64 x f16, stride=2> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.array<88 x f16, stride=2> [0]>, StorageBuffer>
   %arg5: memref<16x4xf16, offset: 0, strides: [4, 1]>,
   %arg6: memref<16x4xf16, offset: 8, strides: [4, 1]>,
   %arg7: memref<16x4xf16, offset: 0, strides: [16, 1]>,
   %arg8: memref<16x4xf16, offset: 0, strides: [1, 16]>,
   %arg9: memref<16x4xf16, offset: 0, strides: [1, 22]>
 ) { return }
 
 } // end module

 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
 // CHECK-LABEL: spv.func @int_tensor_types
-// CHECK-SAME: !spv.array<32 x i64 [8]>
-// CHECK-SAME: !spv.array<32 x i32 [4]>
-// CHECK-SAME: !spv.array<32 x i16 [2]>
-// CHECK-SAME: !spv.array<32 x i8 [1]>
+// CHECK-SAME: !spv.array<32 x i64, stride=8>
+// CHECK-SAME: !spv.array<32 x i32, stride=4>
+// CHECK-SAME: !spv.array<32 x i16, stride=2>
+// CHECK-SAME: !spv.array<32 x i8, stride=1>
 func @int_tensor_types(
   %arg0: tensor<8x4xi64>,
   %arg1: tensor<8x4xi32>,
   %arg2: tensor<8x4xi16>,
   %arg3: tensor<8x4xi8>
 ) { return }
 
 // CHECK-LABEL: spv.func @float_tensor_types
-// CHECK-SAME: !spv.array<32 x f64 [8]>
-// CHECK-SAME: !spv.array<32 x f32 [4]>
-// CHECK-SAME: !spv.array<32 x f16 [2]>
+// CHECK-SAME: !spv.array<32 x f64, stride=8>
+// CHECK-SAME: !spv.array<32 x f32, stride=4>
+// CHECK-SAME: !spv.array<32 x f16, stride=2>
 func @float_tensor_types(
   %arg0: tensor<8x4xf64>,
   %arg1: tensor<8x4xf32>,
   %arg2: tensor<8x4xf16>
 ) { return }
 
 } // end module
 
 // -----
 
 // Check that tensor element types are changed to 32-bit without capabilities.
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
 // CHECK-LABEL: spv.func @int_tensor_types
-// CHECK-SAME: !spv.array<32 x i32 [4]>
-// CHECK-SAME: !spv.array<32 x i32 [4]>
-// CHECK-SAME: !spv.array<32 x i32 [4]>
-// CHECK-SAME: !spv.array<32 x i32 [4]>
+// CHECK-SAME: !spv.array<32 x i32, stride=4>
+// CHECK-SAME: !spv.array<32 x i32, stride=4>
+// CHECK-SAME: !spv.array<32 x i32, stride=4>
+// CHECK-SAME: !spv.array<32 x i32, stride=4>
 func @int_tensor_types(
   %arg0: tensor<8x4xi64>,
   %arg1: tensor<8x4xi32>,
   %arg2: tensor<8x4xi16>,
   %arg3: tensor<8x4xi8>
 ) { return }
 
 // CHECK-LABEL: spv.func @float_tensor_types
-// CHECK-SAME: !spv.array<32 x f32 [4]>
-// CHECK-SAME: !spv.array<32 x f32 [4]>
-// CHECK-SAME: !spv.array<32 x f32 [4]>
+// CHECK-SAME: !spv.array<32 x f32, stride=4>
+// CHECK-SAME: !spv.array<32 x f32, stride=4>
+// CHECK-SAME: !spv.array<32 x f32, stride=4>
 func @float_tensor_types(
   %arg0: tensor<8x4xf64>,
   %arg1: tensor<8x4xf32>,
   %arg2: tensor<8x4xf16>
 ) { return }
 
 } // end module
 

mlir/test/Conversion/StandardToSPIRV/subview-to-spirv.mlir

spv.target_env = #spv.target_env<
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
 //===----------------------------------------------------------------------===//
 // std.subview
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @fold_static_stride_subview_with_load
-// CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<384 x f32 [4]> [0]>, StorageBuffer>, [[ARG1:%.*]]: i32, [[ARG2:%.*]]: i32, [[ARG3:%.*]]: i32, [[ARG4:%.*]]: i32
+// CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<384 x f32, stride=4> [0]>, StorageBuffer>, [[ARG1:%.*]]: i32, [[ARG2:%.*]]: i32, [[ARG3:%.*]]: i32, [[ARG4:%.*]]: i32
 func @fold_static_stride_subview_with_load(%arg0 : memref<12x32xf32>, %arg1 : index, %arg2 : index, %arg3 : index, %arg4 : index) {
   // CHECK: [[C2:%.*]] = spv.constant 2
   // CHECK: [[C3:%.*]] = spv.constant 3
   // CHECK: [[T2:%.*]] = spv.IMul [[ARG3]], [[C2]]
   // CHECK: [[T3:%.*]] = spv.IAdd [[ARG1]], [[T2]]
   // CHECK: [[T4:%.*]] = spv.IMul [[ARG4]], [[C3]]
   // CHECK: [[T5:%.*]] = spv.IAdd [[ARG2]], [[T4]]
   // CHECK: [[C32:%.*]] = spv.constant 32
   // CHECK: [[T7:%.*]] = spv.IMul [[C32]], [[T3]]
   // CHECK: [[C1:%.*]] = spv.constant 1
   // CHECK: [[T9:%.*]] = spv.IMul [[C1]], [[T5]]
   // CHECK: [[T10:%.*]] = spv.IAdd [[T7]], [[T9]]
   // CHECK: [[C0:%.*]] = spv.constant 0
   // CHECK: [[T12:%.*]] = spv.AccessChain [[ARG0]]{{\[}}[[C0]], [[T10]]
   // CHECK: spv.Load "StorageBuffer" [[T12]] : f32
   %0 = subview %arg0[%arg1, %arg2][][] : memref<12x32xf32> to memref<4x4xf32, offset:?, strides: [64, 3]>
   %1 = load %0[%arg3, %arg4] : memref<4x4xf32, offset:?, strides: [64, 3]>
   return
 }
 
 // CHECK-LABEL: @fold_static_stride_subview_with_store
-// CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<384 x f32 [4]> [0]>, StorageBuffer>, [[ARG1:%.*]]: i32, [[ARG2:%.*]]: i32, [[ARG3:%.*]]: i32, [[ARG4:%.*]]: i32, [[ARG5:%.*]]: f32
+// CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<384 x f32, stride=4> [0]>, StorageBuffer>, [[ARG1:%.*]]: i32, [[ARG2:%.*]]: i32, [[ARG3:%.*]]: i32, [[ARG4:%.*]]: i32, [[ARG5:%.*]]: f32
 func @fold_static_stride_subview_with_store(%arg0 : memref<12x32xf32>, %arg1 : index, %arg2 : index, %arg3 : index, %arg4 : index, %arg5 : f32) {
   // CHECK: [[C2:%.*]] = spv.constant 2
   // CHECK: [[C3:%.*]] = spv.constant 3
   // CHECK: [[T2:%.*]] = spv.IMul [[ARG3]], [[C2]]
   // CHECK: [[T3:%.*]] = spv.IAdd [[ARG1]], [[T2]]
   // CHECK: [[T4:%.*]] = spv.IMul [[ARG4]], [[C3]]
   // CHECK: [[T5:%.*]] = spv.IAdd [[ARG2]], [[T4]]
   // CHECK: [[C32:%.*]] = spv.constant 32

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

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 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.func @array_stride(%arg0 : !spv.ptr<!spv.array<4x!spv.array<4xf32, stride=4>, stride=128>, StorageBuffer>, %arg1 : i32, %arg2 : i32) "None" {
+    // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.array<4 x !spv.array<4 x f32, stride=4>, stride=128>, StorageBuffer>
+    %2 = spv.AccessChain %arg0[%arg1, %arg2] : !spv.ptr<!spv.array<4x!spv.array<4xf32, stride=4>, stride=128>, StorageBuffer>
     spv.Return
   }
 }
 
 // -----
 
 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<f32, stride=4>, StorageBuffer>
+  spv.globalVariable @var0 : !spv.ptr<!spv.rtarray<f32, stride=4>, 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/constant.mlir

spv.func @const_variable(%arg0 : i32, %arg1 : i32) -> () "None" {
     %1 = spv.constant 5 : i32
     %2 = spv.Variable init(%1) : !spv.ptr<i32, Function>
     %3 = spv.Load "Function" %2 : i32
     %4 = spv.IAdd %0, %3 : i32
     spv.Return
   }
 
   // CHECK-LABEL: @multi_dimensions_const
-  spv.func @multi_dimensions_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>) "None" {
-    // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 2 : i32, 3 : i32], [4 : i32, 5 : i32, 6 : i32]], {{\[}}[7 : i32, 8 : i32, 9 : i32], [10 : i32, 11 : i32, 12 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
-    %0 = spv.constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
-    spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
+  spv.func @multi_dimensions_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>) "None" {
+    // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 2 : i32, 3 : i32], [4 : i32, 5 : i32, 6 : i32]], {{\[}}[7 : i32, 8 : i32, 9 : i32], [10 : i32, 11 : i32, 12 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+    %0 = spv.constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+    spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
   }
 
   // CHECK-LABEL: @multi_dimensions_splat_const
-  spv.func @multi_dimensions_splat_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>) "None" {
-    // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]], {{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
-    %0 = spv.constant dense<1> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
-    spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32 [4]> [12]> [24]>
+  spv.func @multi_dimensions_splat_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>) "None" {
+    // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]], {{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+    %0 = spv.constant dense<1> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+    spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
   }
 }

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

spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 }
 
 // -----
 
 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.globalVariable @GV1 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<10 x f32, stride=4> [0]>, StorageBuffer>
+  spv.globalVariable @GV2 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<10 x f32, stride=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>
+    %0 = spv._address_of @GV1 : !spv.ptr<!spv.struct<!spv.array<10 x f32, stride=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>
+    %2 = spv.AccessChain %0[%1] : !spv.ptr<!spv.struct<!spv.array<10 x f32, stride=4> [0]>, StorageBuffer>
+    %3 = spv._address_of @GV2 : !spv.ptr<!spv.struct<!spv.array<10 x f32, stride=4> [0]>, StorageBuffer>
+    %5 = spv.AccessChain %3[%1] : !spv.ptr<!spv.struct<!spv.array<10 x f32, stride=4> [0]>, StorageBuffer>
     %6 = spv.constant 4 : i32
     %7 = spv.constant 42 : i32
     %8 = spv.constant 2 : i32
 // CHECK:        spv.Branch ^bb1(%{{.*}} : i32)
 // CHECK-NEXT: ^bb1(%[[OUTARG:.*]]: i32):
 // CHECK-NEXT:   spv.loop {
     spv.loop {
 // CHECK-NEXT:     spv.Branch ^bb1(%[[OUTARG]] : i32)
       spv.Branch ^header(%6 : i32)
 // CHECK-NEXT:   ^bb1(%[[HEADARG:.*]]: i32):
     ^header(%9: i32):
       %10 = spv.SLessThan %9, %7 : i32
 // CHECK:          spv.BranchConditional %{{.*}}, ^bb2, ^bb3
       spv.BranchConditional %10, ^body, ^merge
 // CHECK-NEXT:   ^bb2:     // pred: ^bb1
     ^body:
-      %11 = spv.AccessChain %2[%9] : !spv.ptr<!spv.array<10 x f32 [4]>, StorageBuffer>
+      %11 = spv.AccessChain %2[%9] : !spv.ptr<!spv.array<10 x f32, stride=4>, StorageBuffer>
       %12 = spv.Load "StorageBuffer" %11 : f32
-      %13 = spv.AccessChain %5[%9] : !spv.ptr<!spv.array<10 x f32 [4]>, StorageBuffer>
+      %13 = spv.AccessChain %5[%9] : !spv.ptr<!spv.array<10 x f32, stride=4>, StorageBuffer>
       spv.Store "StorageBuffer" %13, %12 : f32
 // CHECK:          %[[ADD:.*]] = spv.IAdd
       %14 = spv.IAdd %9, %8 : i32
 // CHECK-NEXT:     spv.Branch ^bb1(%[[ADD]] : i32)
       spv.Branch ^header(%14 : i32)
 // CHECK-NEXT:   ^bb3:
     ^merge:
 // CHECK-NEXT:     spv._merge

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

spv.func @access_chain(%arg0 : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>, %arg1 : i32, %arg2 : i32) "None" {
     %2 = spv.AccessChain %arg0[%arg1, %arg2] : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>
     spv.Return
   }
 }
 
 // -----
 
 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]>
+  spv.func @load_store_zero_rank_float(%arg0: !spv.ptr<!spv.struct<!spv.array<1 x f32, stride=4> [0]>, StorageBuffer>, %arg1: !spv.ptr<!spv.struct<!spv.array<1 x f32, stride=4> [0]>, StorageBuffer>) "None" {
+    // CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<!spv.array<1 x f32, stride=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>
+    %1 = spv.AccessChain %arg0[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x f32, stride=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]>
+    // CHECK: [[STORE_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<!spv.array<1 x f32, stride=4> [0]>
     // CHECK-NEXT: spv.Store "StorageBuffer" [[STORE_PTR]], [[VAL]] : f32
     %3 = spv.constant 0 : i32
-    %4 = spv.AccessChain %arg1[%3, %3] : !spv.ptr<!spv.struct<!spv.array<1 x f32 [4]> [0]>, StorageBuffer>
+    %4 = spv.AccessChain %arg1[%3, %3] : !spv.ptr<!spv.struct<!spv.array<1 x f32, stride=4> [0]>, StorageBuffer>
     spv.Store "StorageBuffer" %4, %2 : f32
     spv.Return
   }
 
-  spv.func @load_store_zero_rank_int(%arg0: !spv.ptr<!spv.struct<!spv.array<1 x i32 [4]> [0]>, StorageBuffer>, %arg1: !spv.ptr<!spv.struct<!spv.array<1 x i32 [4]> [0]>, StorageBuffer>) "None" {
-    // CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<!spv.array<1 x i32 [4]> [0]>
+  spv.func @load_store_zero_rank_int(%arg0: !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, %arg1: !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>) "None" {
+    // CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>
     // CHECK-NEXT: [[VAL:%.*]] = spv.Load "StorageBuffer" [[LOAD_PTR]] : i32
     %0 = spv.constant 0 : i32
-    %1 = spv.AccessChain %arg0[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32 [4]> [0]>, StorageBuffer>
+    %1 = spv.AccessChain %arg0[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
     %2 = spv.Load "StorageBuffer" %1 : i32
 
-    // CHECK: [[STORE_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<!spv.array<1 x i32 [4]> [0]>
+    // CHECK: [[STORE_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>
     // CHECK-NEXT: spv.Store "StorageBuffer" [[STORE_PTR]], [[VAL]] : i32
     %3 = spv.constant 0 : i32
-    %4 = spv.AccessChain %arg1[%3, %3] : !spv.ptr<!spv.struct<!spv.array<1 x i32 [4]> [0]>, StorageBuffer>
+    %4 = spv.AccessChain %arg1[%3, %3] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
     spv.Store "StorageBuffer" %4, %2 : i32
     spv.Return
   }
 }

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

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 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<!spv.array<128 x f32, stride=4> [0]>, Input>
+  spv.globalVariable @var0 bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<128 x f32, stride=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], !spv.struct<f32 [0], !spv.array<16 x f32, stride=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, stride=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>
 
-  // CHECK: !spv.ptr<!spv.struct<!spv.array<128 x !spv.struct<!spv.array<128 x f32 [4]> [0]> [4]> [0]>, StorageBuffer>
-  spv.globalVariable @var3 : !spv.ptr<!spv.struct<!spv.array<128 x !spv.struct<!spv.array<128 x f32 [4]> [0]> [4]> [0]>, StorageBuffer>
+  // CHECK: !spv.ptr<!spv.struct<!spv.array<128 x !spv.struct<!spv.array<128 x f32, stride=4> [0]>, stride=512> [0]>, StorageBuffer>
+  spv.globalVariable @var3 : !spv.ptr<!spv.struct<!spv.array<128 x !spv.struct<!spv.array<128 x f32, stride=4> [0]>, stride=512> [0]>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<f32 [0, NonWritable], i32 [4]>, StorageBuffer>
   spv.globalVariable @var4 : !spv.ptr<!spv.struct<f32 [0, NonWritable], i32 [4]>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<f32 [NonWritable], i32 [NonWritable, NonReadable]>, StorageBuffer>
   spv.globalVariable @var5 : !spv.ptr<!spv.struct<f32 [NonWritable], i32 [NonWritable, NonReadable]>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<f32 [0, NonWritable], i32 [4, NonWritable, NonReadable]>, StorageBuffer>
   spv.globalVariable @var6 : !spv.ptr<!spv.struct<f32 [0, NonWritable], i32 [4, NonWritable, NonReadable]>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<>, StorageBuffer>
   spv.globalVariable @empty : !spv.ptr<!spv.struct<>, StorageBuffer>
 
-  // CHECK: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Input>,
-  // CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Output>
-  spv.func @kernel_1(%arg0: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Input>, %arg1: !spv.ptr<!spv.struct<!spv.array<128 x f32 [4]> [0]>, Output>) -> () "None" {
+  // CHECK: !spv.ptr<!spv.struct<!spv.array<128 x f32, stride=4> [0]>, Input>,
+  // CHECK-SAME: !spv.ptr<!spv.struct<!spv.array<128 x f32, stride=4> [0]>, Output>
+  spv.func @kernel(%arg0: !spv.ptr<!spv.struct<!spv.array<128 x f32, stride=4> [0]>, Input>, %arg1: !spv.ptr<!spv.struct<!spv.array<128 x f32, stride=4> [0]>, Output>) -> () "None" {
     spv.Return
   }
 }

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

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
     {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 [4]> [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}},
     %arg1: !spv.ptr<!spv.struct<!spv.array<12 x f32>>, StorageBuffer>
            {spv.interface_var_abi = {binding = 1 : i32,

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

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>
-  // CHECK-DAG: spv.globalVariable [[VAR0:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32 [4]> [16]> [0]>, StorageBuffer>
-  // CHECK-DAG: spv.globalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32 [4]> [16]> [0]>, StorageBuffer>
-  // CHECK-DAG: spv.globalVariable [[VAR2:@.*]] bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32 [4]> [16]> [0]>, StorageBuffer>
+  // CHECK-DAG: spv.globalVariable [[VAR0:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32, stride=4>, stride=16> [0]>, StorageBuffer>
+  // CHECK-DAG: spv.globalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<12 x !spv.array<4 x f32, stride=4>, stride=16> [0]>, StorageBuffer>
+  // 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,

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 {
   // 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>
+  // CHECK: spv.globalVariable @var1 bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<64 x i32, stride=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>
+  // CHECK: spv.globalVariable @var2 bind(1, 0) : !spv.ptr<!spv.struct<!spv.struct<!spv.array<64 x i32, stride=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>
 
-  // CHECK: spv.globalVariable @var3 : !spv.ptr<!spv.struct<!spv.array<16 x !spv.struct<f32 [0], f32 [4], !spv.array<16 x f32 [4]> [8]> [72]> [0], f32 [1152]>, StorageBuffer>
+  // CHECK: spv.globalVariable @var3 : !spv.ptr<!spv.struct<!spv.array<16 x !spv.struct<f32 [0], f32 [4], !spv.array<16 x f32, stride=4> [8]>, stride=72> [0], f32 [1152]>, StorageBuffer>
   spv.globalVariable @var3 : !spv.ptr<!spv.struct<!spv.array<16x!spv.struct<f32, f32, !spv.array<16xf32>>>, f32>, StorageBuffer>
 
   // CHECK: spv.globalVariable @var4 bind(1, 2) : !spv.ptr<!spv.struct<!spv.struct<!spv.struct<i1 [0], i8 [1], i16 [2], i32 [4], i64 [8]> [0], f32 [16], i1 [20]> [0], i1 [24]>, StorageBuffer>
   spv.globalVariable @var4 bind(1,2) : !spv.ptr<!spv.struct<!spv.struct<!spv.struct<i1, i8, i16, i32, i64>, f32, i1>, i1>, StorageBuffer>
 
-  // CHECK: spv.globalVariable @var5 bind(1, 3) : !spv.ptr<!spv.struct<!spv.array<256 x f32 [4]> [0]>, StorageBuffer>
+  // CHECK: spv.globalVariable @var5 bind(1, 3) : !spv.ptr<!spv.struct<!spv.array<256 x f32, stride=4> [0]>, StorageBuffer>
   spv.globalVariable @var5 bind(1,3) : !spv.ptr<!spv.struct<!spv.array<256xf32>>, StorageBuffer>
 
   spv.func @kernel() -> () "None" {
     %c0 = spv.constant 0 : i32
     // CHECK: {{%.*}} = spv._address_of @var0 : !spv.ptr<!spv.struct<i32 [0], !spv.struct<f32 [0], i32 [4]> [4], f32 [12]>, Uniform>
     %0 = spv._address_of @var0 : !spv.ptr<!spv.struct<i32, !spv.struct<f32, i32>, f32>, Uniform>
     // 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 {
   // 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>
+  // 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, stride=8>, stride=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>
 
-  // CHECK: spv.globalVariable @var3 : !spv.ptr<!spv.struct<!spv.struct<!spv.array<64 x i64 [8]> [0], i1 [512]> [0], i1 [520]>, Uniform>
+  // CHECK: spv.globalVariable @var3 : !spv.ptr<!spv.struct<!spv.struct<!spv.array<64 x i64, stride=8> [0], i1 [512]> [0], i1 [520]>, Uniform>
   spv.globalVariable @var3 : !spv.ptr<!spv.struct<!spv.struct<!spv.array<64xi64>, i1>, i1>, Uniform>
 
   // CHECK: spv.globalVariable @var4 : !spv.ptr<!spv.struct<i1 [0], !spv.struct<i64 [0], i1 [8], i1 [9], i1 [10], i1 [11]> [8], i1 [24]>, Uniform>
   spv.globalVariable @var4 : !spv.ptr<!spv.struct<i1, !spv.struct<i64, i1, i1, i1, i1>, i1>, Uniform>
 
   // CHECK: spv.globalVariable @var5 : !spv.ptr<!spv.struct<i1 [0], !spv.struct<i1 [0], i1 [1], i1 [2], i1 [3], i64 [8]> [8], i1 [24]>, Uniform>
   spv.globalVariable @var5 : !spv.ptr<!spv.struct<i1, !spv.struct<i1, i1, i1, i1, i64>, i1>, Uniform>
 

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

 //===----------------------------------------------------------------------===//
 
 func @const() -> () {
   // CHECK: %0 = spv.constant true
   // CHECK: %1 = spv.constant 42 : i32
   // CHECK: %2 = spv.constant 5.000000e-01 : f32
   // CHECK: %3 = spv.constant dense<[2, 3]> : vector<2xi32>
   // CHECK: %4 = spv.constant [dense<3.000000e+00> : vector<2xf32>] : !spv.array<1 x vector<2xf32>>
-  // CHECK: %5 = spv.constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32 [4]> [12]>
-  // CHECK: %6 = spv.constant dense<1.000000e+00> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32 [4]> [12]>
-  // CHECK: %7 = spv.constant dense<{{\[}}[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32 [4]> [12]>
-  // CHECK: %8 = spv.constant dense<{{\[}}[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32 [4]> [12]>
+  // CHECK: %5 = spv.constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+  // CHECK: %6 = spv.constant dense<1.000000e+00> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
+  // CHECK: %7 = spv.constant dense<{{\[}}[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+  // CHECK: %8 = spv.constant dense<{{\[}}[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
 
   %0 = spv.constant true
   %1 = spv.constant 42 : i32
   %2 = spv.constant 0.5 : f32
   %3 = spv.constant dense<[2, 3]> : vector<2xi32>
   %4 = spv.constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>>
-  %5 = spv.constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32 [4]> [12]>
-  %6 = spv.constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32 [4]> [12]>
-  %7 = spv.constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32 [4]> [12]>
-  %8 = spv.constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32 [4]> [12]>
+  %5 = spv.constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+  %6 = spv.constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
+  %7 = spv.constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+  %8 = spv.constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
   return
 }
 
 // -----
 
 func @unaccepted_std_attr() -> () {
   // expected-error @+1 {{cannot have value of type 'none'}}
   %0 = spv.constant unit : none
func @value_result_type_mismatch() -> () {
   // expected-error @+1 {{must have spv.array result type for array value}}
   %0 = "spv.constant"() {value = dense<0> : tensor<4xi32>} : () -> (vector<4xi32>)
 }
 
 // -----
 
 func @value_result_type_mismatch() -> () {
   // expected-error @+1 {{result element type ('i32') does not match value element type ('f32')}}
-  %0 = spv.constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x i32 [4]> [12]>
+  %0 = spv.constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x i32>>
 }
 
 // -----
 
 func @value_result_num_elements_mismatch() -> () {
   // expected-error @+1 {{result number of elements (6) does not match value number of elements (4)}}
-  %0 = spv.constant dense<1.0> : tensor<2x2xf32> : !spv.array<2 x !spv.array<3 x f32 [4]> [12]>
+  %0 = spv.constant dense<1.0> : tensor<2x2xf32> : !spv.array<2 x !spv.array<3 x f32>>
   return
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.EntryPoint
 //===----------------------------------------------------------------------===//

mlir/test/Dialect/SPIRV/types.mlir

 // RUN: mlir-opt -split-input-file -verify-diagnostics %s | FileCheck %s
 
 // TODO(b/133530217): Add more tests after switching to the generic parser.
 
 //===----------------------------------------------------------------------===//
 // ArrayType
 //===----------------------------------------------------------------------===//
 
 // CHECK: func @scalar_array_type(!spv.array<16 x f32>, !spv.array<8 x i32>)
 func @scalar_array_type(!spv.array<16xf32>, !spv.array<8 x i32>) -> ()
 
 // CHECK: func @vector_array_type(!spv.array<32 x vector<4xf32>>)
 func @vector_array_type(!spv.array< 32 x vector<4xf32> >) -> ()
 
-// CHECK: func @array_type_stride(!spv.array<4 x !spv.array<4 x f32 [4]> [128]>)
-func @array_type_stride(!spv.array< 4 x !spv.array<4 x f32 [4]> [128]>) -> ()
+// CHECK: func @array_type_stride(!spv.array<4 x !spv.array<4 x f32, stride=4>, stride=128>)
+func @array_type_stride(!spv.array< 4 x !spv.array<4 x f32, stride=4>, stride = 128>) -> ()
 
 // -----
 
 // expected-error @+1 {{expected '<'}}
 func @missing_left_angle_bracket(!spv.array 4xf32>) -> ()
 
 // -----
 

 // -----
 
 // expected-error @+1 {{cannot use '!llvm.i32' to compose SPIR-V types}}
 func @llvm_type(!spv.array<4x!llvm.i32>) -> ()
 
 // -----
 
 // expected-error @+1 {{ArrayStride must be greater than zero}}
-func @array_type_zero_stride(!spv.array<4xi32 [0]>) -> ()
+func @array_type_zero_stride(!spv.array<4xi32, stride=0>) -> ()
 
 // -----
 
 // expected-error @+1 {{expected array length greater than 0}}
 func @array_type_zero_length(!spv.array<0xf32>) -> ()
 
 // -----
 

 //===----------------------------------------------------------------------===//
 
 // CHECK: func @scalar_runtime_array_type(!spv.rtarray<f32>, !spv.rtarray<i32>)
 func @scalar_runtime_array_type(!spv.rtarray<f32>, !spv.rtarray<i32>) -> ()
 
 // CHECK: func @vector_runtime_array_type(!spv.rtarray<vector<4xf32>>)
 func @vector_runtime_array_type(!spv.rtarray< vector<4xf32> >) -> ()
 
+// CHECK: func @runtime_array_type_stride(!spv.rtarray<f32, stride=4>)
+func @runtime_array_type_stride(!spv.rtarray<f32, stride=4>) -> ()
+
 // -----
 
 // expected-error @+1 {{expected '<'}}
 func @missing_left_angle_bracket(!spv.rtarray f32>) -> ()
 
 // -----
 
 // expected-error @+1 {{expected non-function type}}
 func @missing_element_type(!spv.rtarray<>) -> ()
 
 // -----
 
 // expected-error @+1 {{expected non-function type}}
 func @redundant_count(!spv.rtarray<4xf32>) -> ()
 
 // -----
 
+// expected-error @+1 {{ArrayStride must be greater than zero}}
+func @runtime_array_type_zero_stride(!spv.rtarray<i32, stride=0>) -> ()
+
+// -----
+
 //===----------------------------------------------------------------------===//
 // ImageType
 //===----------------------------------------------------------------------===//
 
 // CHECK: func @image_parameters_1D(!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Unknown>)
 func @image_parameters_1D(!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Unknown>) -> ()
 
 // -----