[WIP]

mlir/include/mlir/Dialect/Vector/IR/VectorOps.td

 }
 
 def Vector_TransferReadOp :
   Vector_Op<"transfer_read", [
       DeclareOpInterfaceMethods<VectorTransferOpInterface>,
       DeclareOpInterfaceMethods<VectorUnrollOpInterface, ["getShapeForUnroll"]>,
       DeclareOpInterfaceMethods<MaskableOpInterface>,
       DeclareOpInterfaceMethods<MemoryEffectsOpInterface>,
+      TypesMatchWith<"type of 'padding' matches element type of 'source'",
+                     "source", "padding",
+                     "$_self.cast<ShapedType>().getElementType()">,
+      // TODO: This can be only used once "broadcast"-behavior is retired and the
+      // mask type becomes unambiguous,
+      // TypesMatchWith<
+      //   "type of 'mask' matches type of 'vector' with an i1 elt type",
+      //   "vector", "mask",
+      //   [{
+      //     VectorType::get($_self.cast<VectorType>().getShape(),
+      //                     IntegerType::get($_self.getContext(), 1))
+      //   }]>,
       AttrSizedOperandSegments
     ]>,
     Arguments<(ins AnyShaped:$source,
                    Variadic<Index>:$indices,
-                   AffineMapAttr:$permutation_map,
+                   OptionalAttr<AffineMapAttr>:$permutation_map,
                    AnyType:$padding,
                    Optional<VectorOf<[I1]>>:$mask,
                    OptionalAttr<BoolArrayAttr>:$in_bounds)>,
     Results<(outs AnyVectorOfAnyRank:$vector)> {
 
-  let summary = "Reads a supervector from memory into an SSA vector value.";
+  let summary = "Reads an n-D from memory into an SSA vector value.";
 
   let description = [{
-    The `vector.transfer_read` op performs a read from a slice within a
+    The `vector.transfer_read` op performs a read from a slice of a
     [MemRef](../LangRef.md#memref-type) or a Ranked
     [Tensor](../LangRef.md#tensor-type) supplied as its first operand
-    into a [vector](../LangRef.md#vector-type) of the same base elemental type.
+    into an n-D [vector](../LangRef.md#vector-type) of the same base elemental
+    type.
 
     A memref/tensor operand with vector element type, must have its vector
     element type match a suffix (shape and element type) of the vector (e.g.
     memref<3x2x6x4x3xf32>, vector<1x1x4x3xf32>).
 
     The slice is further defined by a full-rank index within the MemRef/Tensor,
     supplied as the operands `[1 .. 1 + rank(memref/tensor))` that defines the
     starting point of the transfer (e.g. `%A[%i0, %i1, %i2]`).
 
-    The permutation_map [attribute](../LangRef.md#attributes) is an
+    The permutation_map [attribute](../LangRef.md#attributes) is an optional 
     [affine-map](Affine.md#affine-maps) which specifies the transposition on the
     slice to match the vector shape. The permutation map may be implicit and
     omitted from parsing and printing if it is the canonical minor identity map
     (i.e. if it does not permute or broadcast any dimension).
 
     The size of the slice is specified by the size of the vector, given as the
     return type.
 
let description = [{
     While the starting point of the transfer has to be in-bounds, accesses may
     run out-of-bounds as indices increase. Broadcast dimensions must always be
     in-bounds. If specified, the `in_bounds` array length has to be equal to the
     vector rank. In absence of the attribute, accesses along all dimensions
     (except for broadcasts) may run out-of-bounds. A `vector.transfer_read` can
     be lowered to a simple load if all dimensions are specified to be within
     bounds and no `mask` was specified.
 
-    This operation is called 'read' by opposition to 'load' because the
-    super-vector granularity is generally not representable with a single
-    hardware register. A `vector.transfer_read` is thus a mid-level abstraction
-    that supports super-vectorization with non-effecting padding for full-tile
-    only operations.
+    This operation is called 'read' by opposition to 'load' because the n-D
+    vector granularity is generally not representable with a single hardware
+    register.
+    A `vector.transfer_read` is thus a mid-level abstraction that supports 
+    super-vectorization with non-effecting padding for full-tile only operations.
 
     More precisely, let's dive deeper into the permutation_map for the following
-    MLIR:
+    MLIR snippet:
 
     ```mlir
     vector.transfer_read %A[%expr1, %expr2, %expr3, %expr4]
       { permutation_map : (d0,d1,d2,d3) -> (d2,0,d0) } :
       memref<?x?x?x?xf32>, vector<3x4x5xf32>
     ```
 
-    This operation always reads a slice starting at `%A[%expr1, %expr2, %expr3,
-    %expr4]`. The size of the slice is 3 along d2 and 5 along d0, so the slice
-    is: `%A[%expr1 : %expr1 + 5, %expr2, %expr3:%expr3 + 3, %expr4]`
-
-    That slice needs to be read into a `vector<3x4x5xf32>`. Since the
-    permutation map is not full rank, there must be a broadcast along vector
-    dimension `1`.
+    This operation always reads a slice starting at 
+      `%A[%expr1, %expr2, %expr3,  %expr4]`.
+    The size of the slice is 3 along d2 and 5 along d0, so the slice is: 
+      `%A[%expr1 : %expr1 + 5, %expr2, %expr3:%expr3 + 3, %expr4]`
+
+    That slice needs to be read into a `vector<3x4x5xf32>`. Since the permutation
+    map is not full rank, there must be a broadcast along vector dimension `1`.
 
     A notional lowering of vector.transfer_read could generate code resembling:
 
     ```mlir
     // %expr1, %expr2, %expr3, %expr4 defined before this point
-    %tmp = alloc() : vector<3x4x5xf32>
-    %view_in_tmp = "element_type_cast"(%tmp) : memref<1xvector<3x4x5xf32>>
+    %tmp = alloc() : memref<3x4x5xf32>
+    %view_in_tmp = "element_type_cast"(%tmp) : memref<vector<3x4x5xf32>>
     for %i = 0 to 3 {
       affine.for %j = 0 to 4 {
         affine.for %k = 0 to 5 {
           %a = load %A[%expr1 + %k, %expr2, %expr3 + %i, %expr4] :
             memref<?x?x?x?xf32>
-          store %tmp[%i, %j, %k] : vector<3x4x5xf32>
+          store %a, %tmp[%i, %j, %k] : memref<3x4x5xf32>
     }}}
     %c0 = arith.constant 0 : index
-    %vec = load %view_in_tmp[%c0] : vector<3x4x5xf32>
+    %vec = load %view_in_tmp[] : vector<3x4x5xf32>
     ```
+    Notice that the temporary storage footprint is `3 * 5` values but `3 * 4 * 5`
+    values are actually transferred between `%A` and `%tmp`.
     
-    On a GPU one could then map `i`, `j`, `k` to blocks and threads. Notice that
-    the temporary storage footprint is `3 * 5` values but `3 * 4 * 5` values are
-    actually transferred between `%A` and `%tmp`.
-
-    Alternatively, if a notional vector broadcast operation were available, the
-    lowered code would resemble:
+    Alternatively, the lowered code could use a vector.broadcast and resemble:
 
     ```mlir
     // %expr1, %expr2, %expr3, %expr4 defined before this point
-    %tmp = alloc() : vector<3x4x5xf32>
-    %view_in_tmp = "element_type_cast"(%tmp) : memref<1xvector<3x4x5xf32>>
+    %tmp = alloc() : memref<3x5xf32>
+    %view_in_tmp = "element_type_cast"(%tmp) : memref<vector<3x5xf32>>
     for %i = 0 to 3 {
       affine.for %k = 0 to 5 {
         %a = load %A[%expr1 + %k, %expr2, %expr3 + %i, %expr4] :
           memref<?x?x?x?xf32>
-        store %tmp[%i, 0, %k] : vector<3x4x5xf32>
+        store %tmp[%i, %k] : memref<3x5xf32>
     }}
     %c0 = arith.constant 0 : index
-    %tmpvec = load %view_in_tmp[%c0] : vector<3x4x5xf32>
-    %vec = broadcast %tmpvec, 1 : vector<3x4x5xf32>
+    %tmpvec = load %view_in_tmp[%c0] : vector<3x5xf32>
+    %tmpvec2 = vector.broadcast %tmpvec : vector<3x5xf32> to vector<4x3x5xf32>
+    %vec = vector.transpose %tmpvec2, [1, 0, 2] : vector<4x3x5xf32> to vector<3x4x5xf32>
     ```
 
-    where `broadcast` broadcasts from element 0 to all others along the
-    specified dimension. This time, the temporary storage footprint is `3 * 5`
-    values which is the same amount of data as the `3 * 5` values transferred.
-    An additional `1` broadcast is required. On a GPU this broadcast could be
-    implemented using a warp-shuffle if loop `j` were mapped to `threadIdx.x`.
+    This time, the temporary storage footprint is `3 * 5` values which is the
+    same amount of data as the `3 * 5` values transferred, at the expense of an
+    additional vector.broadcast and vector.transpose.
     
     Syntax
     ```
     operation ::= ssa-id `=` `vector.transfer_read` ssa-use-list
       `{` attribute-entry `} :` memref-type `,` vector-type
     ```
 
     Example:
 
let builders = [
     /// 'getMinorIdentityMap'.
     OpBuilder<(ins "VectorType":$vectorType,
                    "Value":$source,
                    "ValueRange":$indices,
                    CArg<"Optional<ArrayRef<bool>>", "::llvm::None">:$inBounds)>,
   ];
 
   let extraClassDeclaration = [{
+    AffineMap getPermutationMapOrMinorIdentity() {
+      if (getPermutationMap()) return *getPermutationMap();
+      return getTransferMinorIdentityMap(getSource().getType().cast<ShapedType>(),
+                                         getVector().getType().cast<VectorType>());
+    }
+
     // MaskableOpInterface methods.
     bool supportsPassthru() { return true; }
   }];
 
+  let assemblyFormat = [{
+    $source `[` $indices `]` 
+    `,` $padding 
+    (`,` $mask^ `:` type($mask))?
+    custom<CustomTransferAttrs>( $permutation_map, $in_bounds, attr-dict )
+    `:` type($source) `,` type($vector)
+  }];
   let hasCanonicalizer = 1;
-  let hasCustomAssemblyFormat = 1;
   let hasFolder = 1;
   let hasVerifier = 1;
 }
 
 def Vector_TransferWriteOp :
   Vector_Op<"transfer_write", [
       DeclareOpInterfaceMethods<VectorTransferOpInterface>,
       DeclareOpInterfaceMethods<VectorUnrollOpInterface, ["getShapeForUnroll"]>,
let builders = [
     /// map to 'getMinorIdentityMap'.
     OpBuilder<(ins "Value":$vector,
                    "Value":$dest,
                    "ValueRange":$indices,
                    CArg<"Optional<ArrayRef<bool>>", "::llvm::None">:$inBounds)>,
   ];
 
   let extraClassDeclaration = [{
+    AffineMap getPermutationMapOrMinorIdentity() {
+      return getPermutationMap();
+    }
+
     /// This method is added to maintain uniformity with load/store
     ///  ops of other dialects.
     Value getValue() { return getVector(); }
 
     std::pair<int64_t, int64_t> getDpsInitsPositionRange() {
       return {1, 2};  // `source` operand
     }
   }];

mlir/include/mlir/Interfaces/VectorInterfaces.td

InterfaceMethod<
       /*methodBody=*/"return $_op.getIndices();"
       /*defaultImplementation=*/
     >,
     InterfaceMethod<
       /*desc=*/"Return the permutation map.",
       /*retTy=*/"::mlir::AffineMap",
       /*methodName=*/"permutation_map",
       /*args=*/(ins),
-      /*methodBody=*/"return $_op.getPermutationMap();"
+      /*methodBody=*/"return $_op.getPermutationMapOrMinorIdentity();"
       /*defaultImplementation=*/
     >,
     InterfaceMethod<
       /*desc=*/[{ Returns true if the specified dimension is a broadcast. }],
       /*retTy=*/"bool",
       /*methodName=*/"isBroadcastDim",
       /*args=*/(ins "unsigned":$idx),
       /*methodBody=*/"",
       /*defaultImplementation=*/[{
-        auto expr = $_op.getPermutationMap().getResult(idx);
+        auto expr = $_op.getPermutationMapOrMinorIdentity().getResult(idx);
         return expr.template isa<::mlir::AffineConstantExpr>() &&
                expr.template dyn_cast<::mlir::AffineConstantExpr>().getValue() == 0;
       }]
     >,
     InterfaceMethod<
       /*desc=*/[{ Returns true if at least one of the dimensions in the
                   permutation map is a broadcast.}],
       /*retTy=*/"bool",
InterfaceMethod<
       /*desc=*/"Return the mask type if the op has a mask.",
       /*retTy=*/"::mlir::VectorType",
       /*methodName=*/"getMaskType",
       /*args=*/(ins),
       /*methodBody=*/"",
       /*defaultImplementation=*/[{
         return $_op.getMask()
             ? ::mlir::vector::detail::transferMaskType(
-                $_op.getVectorType(), $_op.getPermutationMap())
+                $_op.getVectorType(), $_op.getPermutationMapOrMinorIdentity())
             : ::mlir::VectorType();
       }]
     >,
     InterfaceMethod<
       /*desc=*/[{ Return the number of dimensions that participate in the
                   permutation map.}],
       /*retTy=*/"unsigned",
       /*methodName=*/"getTransferRank",
       /*args=*/(ins),
       /*methodBody=*/"",
       /*defaultImplementation=*/
-        "return $_op.getPermutationMap().getNumResults();"
+        "return $_op.getPermutationMapOrMinorIdentity().getNumResults();"
     >,
     InterfaceMethod<
       /*desc=*/[{ Return the number of leading shaped dimensions that do not
                   participate in the permutation map.}],
       /*retTy=*/"unsigned",
       /*methodName=*/"getLeadingShapedRank",
       /*args=*/(ins),
       /*methodBody=*/"",
InterfaceMethod<
       ```
       returns a shape [2, 4].
       }],
       /*retTy=*/"SmallVector<int64_t>",
       /*methodName=*/"getTransferChunkAccessed",
       /*args=*/(ins),
       /*methodBody=*/"",
       /*defaultImplementation=*/[{
-        SmallVector<int64_t> dimSizes($_op.getPermutationMap().getNumDims(), 0);
-        for (auto vecDims : llvm::zip($_op.getPermutationMap().getResults(),
+        SmallVector<int64_t> dimSizes($_op.getPermutationMapOrMinorIdentity().getNumDims(), 0);
+        for (auto vecDims : llvm::zip($_op.getPermutationMapOrMinorIdentity().getResults(),
                                       $_op.getVectorType().getShape())) {
           AffineExpr dim = std::get<0>(vecDims);
           int64_t size = std::get<1>(vecDims);
           // Skip broadcast.
           if (dim.isa<AffineConstantExpr>())
             continue;
           dimSizes[dim.cast<AffineDimExpr>().getPosition()] = size;
         }
         return dimSizes;
       }]
     >,
   ];
 }
 
 #endif // MLIR_INTERFACES_VECTORINTERFACES

mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp

 /// `dimValues`.
 template <typename TransferOpType>
 static void getXferIndices(OpBuilder &b, TransferOpType xferOp,
                            AffineMap offsetMap, ArrayRef<Value> dimValues,
                            SmallVector<Value, 4> &indices) {
   indices.append(xferOp.getIndices().begin(), xferOp.getIndices().end());
   Location loc = xferOp.getLoc();
   unsigned offsetsIdx = 0;
-  for (auto expr : xferOp.getPermutationMap().getResults()) {
+  for (auto expr : xferOp.getPermutationMapOrMinorIdentity().getResults()) {
     if (auto dim = expr.template dyn_cast<AffineDimExpr>()) {
       Value prevIdx = indices[dim.getPosition()];
       SmallVector<Value, 3> dims(dimValues.begin(), dimValues.end());
       dims.push_back(prevIdx);
       AffineExpr d0 = b.getAffineDimExpr(offsetMap.getNumDims());
       indices[dim.getPosition()] = makeComposedAffineApply(
           b, loc, d0 + offsetMap.getResult(offsetsIdx++), dims);
       continue;

 // Return true if the transfer op can be converted to a MMA matrix load.
 static bool transferReadSupportsMMAMatrixType(vector::TransferReadOp readOp,
                                               bool useNvGpu) {
   if (readOp.getMask() || readOp.hasOutOfBoundsDim() ||
       readOp.getVectorType().getRank() != 2)
     return false;
   if (!getMemrefConstantHorizontalStride(readOp.getShapedType()))
     return false;
-  AffineMap map = readOp.getPermutationMap();
+  AffineMap map = readOp.getPermutationMapOrMinorIdentity();
   OpBuilder b(readOp.getContext());
   AffineExpr innerDim = b.getAffineDimExpr(map.getNumDims() - 1);
   AffineExpr zero = b.getAffineConstantExpr(0);
   auto broadcastInnerDim = AffineMap::get(map.getNumDims(), 0, {zero, innerDim},
                                           readOp.getContext());
 
   if (!useNvGpu) {
     // TODO: Support transpose once it is added to GPU dialect ops.
if (writeOp.getTransferRank() == 0)
     return false;
 
   if (writeOp.getMask() || writeOp.hasOutOfBoundsDim() ||
       writeOp.getVectorType().getRank() != 2)
     return false;
   if (!getMemrefConstantHorizontalStride(writeOp.getShapedType()))
     return false;
   // TODO: Support transpose once it is added to GPU dialect ops.
-  if (!writeOp.getPermutationMap().isMinorIdentity())
+  if (!writeOp.getPermutationMapOrMinorIdentity().isMinorIdentity())
     return false;
   return true;
 }
 
 /// Return true if the constant is a splat to a 2D vector so that it can be
 /// converted to a MMA constant matrix op.
 static bool constantSupportsMMAMatrixType(arith::ConstantOp constantOp) {
   auto vecType = constantOp.getType().dyn_cast<VectorType>();
LogicalResult matchAndRewrite(vector::TransposeOp op,
     SmallVector<int64_t, 2> perm;
     op.getTransp(perm);
     SmallVector<unsigned, 2> permU;
     for (int64_t o : perm)
       permU.push_back(unsigned(o));
     AffineMap permutationMap =
         AffineMap::getPermutationMap(permU, op.getContext());
     AffineMap newMap =
-        permutationMap.compose(transferReadOp.getPermutationMap());
+        permutationMap.compose(transferReadOp.getPermutationMapOrMinorIdentity());
     rewriter.replaceOpWithNewOp<vector::TransferReadOp>(
         op, op.getType(), transferReadOp.getSource(),
         transferReadOp.getIndices(), AffineMapAttr::get(newMap),
         transferReadOp.getPadding(), transferReadOp.getMask(),
         transferReadOp.getInBoundsAttr());
     return success();
   }
 };

 }
 
 static void convertTransferReadOp(vector::TransferReadOp op,
                                   llvm::DenseMap<Value, Value> &valueMapping) {
   assert(op.getTransferRank() > 0 && "unexpected 0-d transfer");
   assert(transferReadSupportsMMAMatrixType(op, /*useNvGpu=*/false));
   Optional<int64_t> stride =
       getMemrefConstantHorizontalStride(op.getShapedType());
-  AffineMap map = op.getPermutationMap();
+  AffineMap map = op.getPermutationMapOrMinorIdentity();
   // Handle broadcast by setting the stride to 0.
   if (map.getResult(0).isa<AffineConstantExpr>()) {
     assert(map.getResult(0).cast<AffineConstantExpr>().getValue() == 0);
     stride = 0;
   }
   assert(stride);
   const char *fragType = inferFragType(op);
   gpu::MMAMatrixType type =
creatLdMatrixCompatibleLoads(vector::TransferReadOp op, OpBuilder &builder,
 
   FailureOr<nvgpu::FragmentElementInfo> regInfo =
       nvgpu::getMmaSyncRegisterType(*warpMatrixInfo);
   if (failed(regInfo))
     return failure();
 
   FailureOr<nvgpu::LdMatrixParams> params = nvgpu::getLdMatrixParams(
       *warpMatrixInfo,
-      /*transpose=*/!op.getPermutationMap().isMinorIdentity());
+      /*transpose=*/!op.getPermutationMapOrMinorIdentity().isMinorIdentity());
   if (failed(params)) {
     return op->emitError()
            << "failed to convert vector.transfer_read to ldmatrix; this op "
               "likely "
               "should not be converted to a nvgpu.ldmatrix call.";
   }
 
   // Adjust the load offset.
   auto laneId = builder.create<gpu::LaneIdOp>(loc);
   FailureOr<AffineMap> offsets =
       nvgpu::getLaneIdToLdMatrixMatrixCoord(loc, builder, *params);
   if (failed(offsets))
     return failure();
 
   VectorType vectorType = getMmaSyncVectorOperandType(*regInfo);
 
   SmallVector<Value, 4> indices;
   getXferIndices<vector::TransferReadOp>(builder, op, *offsets, {laneId},
                                          indices);
   nvgpu::LdMatrixOp newOp = builder.create<nvgpu::LdMatrixOp>(
       loc, vectorType, op.getSource(), indices,
-      !op.getPermutationMap().isMinorIdentity(), params->numTiles);
+      !op.getPermutationMapOrMinorIdentity().isMinorIdentity(), params->numTiles);
   valueMapping[op] = newOp->getResult(0);
   return success();
 }
 
 static LogicalResult
 createNonLdMatrixLoads(vector::TransferReadOp op, OpBuilder &builder,
                        llvm::DenseMap<Value, Value> &valueMapping) {
   Location loc = op.getLoc();
createNonLdMatrixLoads(vector::TransferReadOp op, OpBuilder &builder,
   Type loadedElType = regInfo->registerLLVMType;
   VectorType vectorType = getMmaSyncVectorOperandType(*regInfo);
 
   Value fill = builder.create<arith::ConstantOp>(
       op.getLoc(), vectorType.getElementType(),
       builder.getZeroAttr(vectorType.getElementType()));
   Value result = builder.create<vector::SplatOp>(op.getLoc(), fill, vectorType);
 
-  bool isTransposeLoad = !op.getPermutationMap().isMinorIdentity();
+  bool isTransposeLoad = !op.getPermutationMapOrMinorIdentity().isMinorIdentity();
 
   // If we are not transposing, then we can use vectorized loads. Otherwise, we
   // must load each element individually.
   if (!isTransposeLoad) {
     if (!loadedElType.isa<VectorType>()) {
       loadedElType = VectorType::get({1}, loadedElType);
     }
 
bool isLdMatrixCompatible =
       nvgpu::inferTileWidthInBits(*warpMatrixInfo) == 128;
 
   VectorType vecTy = op.getVectorType();
   int64_t bitWidth = vecTy.getElementType().getIntOrFloatBitWidth();
 
   // When we are transposing the B operand, ldmatrix will only work if we have
   // at least 8 rows to read and the width to read for the transpose is 128
   // bits.
-  if (!op.getPermutationMap().isMinorIdentity() &&
+  if (!op.getPermutationMapOrMinorIdentity().isMinorIdentity() &&
       (bitWidth != 16 || vecTy.getDimSize(1) < 8 ||
        vecTy.getDimSize(0) * bitWidth < 128))
     isLdMatrixCompatible = false;
 
   if (!isLdMatrixCompatible)
     return createNonLdMatrixLoads(op, b, valueMapping);
 
   return creatLdMatrixCompatibleLoads(op, b, valueMapping);

mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp

 
 /// Given a vector transfer op, calculate which dimension of the `source`
 /// memref should be unpacked in the next application of TransferOpConversion.
 /// A return value of None indicates a broadcast.
 template <typename OpTy>
 static Optional<int64_t> unpackedDim(OpTy xferOp) {
   // TODO: support 0-d corner case.
   assert(xferOp.getTransferRank() > 0 && "unexpected 0-d transfer");
-  auto map = xferOp.getPermutationMap();
+  auto map = xferOp.getPermutationMapOrMinorIdentity();
   if (auto expr = map.getResult(0).template dyn_cast<AffineDimExpr>()) {
     return expr.getPosition();
   }
   assert(xferOp.isBroadcastDim(0) &&
          "Expected AffineDimExpr or AffineConstantExpr");
   return None;
 }
 
 /// Compute the permutation map for the new (N-1)-D vector transfer op. This
 /// map is identical to the current permutation map, but the first result is
 /// omitted.
 template <typename OpTy>
 static AffineMap unpackedPermutationMap(OpBuilder &b, OpTy xferOp) {
   // TODO: support 0-d corner case.
   assert(xferOp.getTransferRank() > 0 && "unexpected 0-d transfer");
-  auto map = xferOp.getPermutationMap();
+  auto map = xferOp.getPermutationMapOrMinorIdentity();
   return AffineMap::get(map.getNumDims(), 0, map.getResults().drop_front(),
                         b.getContext());
 }
 
 /// Calculate the indices for the new vector transfer op.
 ///
 /// E.g.: transfer_read %A[%a, %b, %c, %d] ... : vector<5x4x3xf32> ...
 ///       --> transfer_read %A[%a, %b + iv, %c, %d] ... vector<4x3f32>

 /// Compute the indices into the memref for the LoadOp/StoreOp generated as
 /// part of TransferOp1dConversion. Return the memref dimension on which
 /// the transfer is operating. A return value of None indicates a broadcast.
 template <typename OpTy>
 static Optional<int64_t>
 get1dMemrefIndices(OpBuilder &b, OpTy xferOp, Value iv,
                    SmallVector<Value, 8> &memrefIndices) {
   auto indices = xferOp.getIndices();
-  auto map = xferOp.getPermutationMap();
+  auto map = xferOp.getPermutationMapOrMinorIdentity();
   assert(xferOp.getTransferRank() > 0 && "unexpected 0-d transfer");
 
   memrefIndices.append(indices.begin(), indices.end());
   assert(map.getNumResults() == 1 &&
          "Expected 1 permutation map result for 1D transfer");
   if (auto expr = map.getResult(0).template dyn_cast<AffineDimExpr>()) {
     Location loc = xferOp.getLoc();
     auto dim = expr.getPosition();

 struct TransferOp1dConversion : public VectorToSCFPattern<OpTy> {
   using VectorToSCFPattern<OpTy>::VectorToSCFPattern;
 
   LogicalResult matchAndRewrite(OpTy xferOp,
                                 PatternRewriter &rewriter) const override {
     // TODO: support 0-d corner case.
     if (xferOp.getTransferRank() == 0)
       return failure();
-    auto map = xferOp.getPermutationMap();
+    auto map = xferOp.getPermutationMapOrMinorIdentity();
     auto memRefType = xferOp.getShapedType().template dyn_cast<MemRefType>();
 
     if (!memRefType)
       return failure();
     if (xferOp.getVectorType().getRank() != 1)
       return failure();
     if (map.isMinorIdentity() && isLastMemrefDimUnitStride(memRefType))
       return failure(); // Handled by ConvertVectorToLLVM

mlir/lib/Dialect/MemRef/Transforms/FoldMemRefAliasOps.cpp

llvm::TypeSwitch<Operation *, void>(loadOp)
         rewriter.replaceOpWithNewOp<decltype(op)>(loadOp, subViewOp.getSource(),
                                                   sourceIndices);
       })
       .Case([&](vector::TransferReadOp transferReadOp) {
         rewriter.replaceOpWithNewOp<vector::TransferReadOp>(
             transferReadOp, transferReadOp.getVectorType(),
             subViewOp.getSource(), sourceIndices,
             getPermutationMapAttr(rewriter.getContext(), subViewOp,
-                                  transferReadOp.getPermutationMap()),
+                                  transferReadOp.getPermutationMapOrMinorIdentity()),
             transferReadOp.getPadding(),
             /*mask=*/Value(), transferReadOp.getInBoundsAttr());
       })
       .Default([](Operation *) { llvm_unreachable("unexpected operation."); });
   return success();
 }
 
 template <typename OpTy>
llvm::TypeSwitch<Operation *, void>(storeOp)
       .Case<AffineStoreOp, memref::StoreOp>([&](auto op) {
         rewriter.replaceOpWithNewOp<decltype(op)>(
             storeOp, storeOp.getValue(), subViewOp.getSource(), sourceIndices);
       })
       .Case([&](vector::TransferWriteOp op) {
         rewriter.replaceOpWithNewOp<vector::TransferWriteOp>(
             op, op.getValue(), subViewOp.getSource(), sourceIndices,
             getPermutationMapAttr(rewriter.getContext(), subViewOp,
-                                  op.getPermutationMap()),
+                                  op.getPermutationMapOrMinorIdentity()),
             op.getInBoundsAttr());
       })
       .Default([](Operation *) { llvm_unreachable("unexpected operation."); });
   return success();
 }
 
 template <typename OpTy>
 LogicalResult StoreOpOfExpandShapeOpFolder<OpTy>::matchAndRewrite(

mlir/lib/Dialect/Vector/IR/VectorOps.cpp

 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
 #include "mlir/Dialect/Utils/IndexingUtils.h"
 #include "mlir/Dialect/Utils/StructuredOpsUtils.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
+#include "mlir/IR/BuiltinAttributes.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/DialectImplementation.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/TypeUtilities.h"
 #include "mlir/Support/LLVM.h"
+#include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/StringSet.h"
 #include "llvm/ADT/TypeSwitch.h"
 #include "llvm/ADT/bit.h"
 #include <numeric>
 
 #include "mlir/Dialect/Vector/IR/VectorOpsDialect.cpp.inc"
 // Pull in all enum type and utility function definitions.
 #include "mlir/Dialect/Vector/IR/VectorOpsEnums.cpp.inc"
return AffineMap::getMinorIdentityMap(
       shapedType.getContext());
 }
 
 bool mlir::vector::checkSameValueRAW(vector::TransferWriteOp defWrite,
                                      vector::TransferReadOp read) {
   return !defWrite.hasOutOfBoundsDim() && !defWrite.getMask() &&
          !read.getMask() && defWrite.getIndices() == read.getIndices() &&
          defWrite.getVectorType() == read.getVectorType() &&
-         defWrite.getPermutationMap() == read.getPermutationMap();
+         defWrite.getPermutationMapOrMinorIdentity() ==
+             read.getPermutationMapOrMinorIdentity();
 }
 
 bool mlir::vector::checkSameValueWAW(vector::TransferWriteOp write,
                                      vector::TransferWriteOp priorWrite) {
   return priorWrite.getIndices() == write.getIndices() &&
          priorWrite.getMask() == write.getMask() &&
          priorWrite.getVectorType() == write.getVectorType() &&
-         priorWrite.getPermutationMap() == write.getPermutationMap();
+         priorWrite.getPermutationMapOrMinorIdentity() ==
+             write.getPermutationMapOrMinorIdentity();
 }
 
 bool mlir::vector::isDisjointTransferIndices(
     VectorTransferOpInterface transferA, VectorTransferOpInterface transferB) {
   // For simplicity only look at transfer of same type.
   if (transferA.getVectorType() != transferB.getVectorType())
     return false;
   unsigned rankOffset = transferA.getLeadingShapedRank();
for (auto attr : *inBounds) {
       }
     }
   }
   if (elideInBounds)
     elidedAttrs.push_back(op.getInBoundsAttrStrName());
   p.printOptionalAttrDict(op->getAttrs(), elidedAttrs);
 }
 
-void TransferReadOp::print(OpAsmPrinter &p) {
-  p << " " << getSource() << "[" << getIndices() << "], " << getPadding();
-  if (getMask())
-    p << ", " << getMask();
-  printTransferAttrs(p, *this);
-  p << " : " << getShapedType() << ", " << getVectorType();
-}
+static void printCustomTransferAttrs(OpAsmPrinter &printer, Operation *op,
+                                     AffineMapAttr permutationMapAttr,
+                                     ArrayAttr inBoundsAttr,
+                                     DictionaryAttr attrs) {
+  auto xferOp = cast<VectorTransferOpInterface>(op);
+  SmallVector<StringRef, 3> elidedAttrs;
+  elidedAttrs.push_back(TransferReadOp::getOperandSegmentSizeAttr());
 
-ParseResult TransferReadOp::parse(OpAsmParser &parser, OperationState &result) {
-  auto &builder = parser.getBuilder();
-  SMLoc typesLoc;
-  OpAsmParser::UnresolvedOperand sourceInfo;
-  SmallVector<OpAsmParser::UnresolvedOperand, 8> indexInfo;
-  OpAsmParser::UnresolvedOperand paddingInfo;
-  SmallVector<Type, 2> types;
-  OpAsmParser::UnresolvedOperand maskInfo;
-  // Parsing with support for paddingValue.
-  if (parser.parseOperand(sourceInfo) ||
-      parser.parseOperandList(indexInfo, OpAsmParser::Delimiter::Square) ||
-      parser.parseComma() || parser.parseOperand(paddingInfo))
-    return failure();
-  ParseResult hasMask = parser.parseOptionalComma();
-  if (hasMask.succeeded()) {
-    if (parser.parseOperand(maskInfo))
-      return failure();
-  }
-  if (parser.parseOptionalAttrDict(result.attributes) ||
-      parser.getCurrentLocation(&typesLoc) || parser.parseColonTypeList(types))
-    return failure();
-  if (types.size() != 2)
-    return parser.emitError(typesLoc, "requires two types");
-  auto indexType = builder.getIndexType();
-  auto shapedType = types[0].dyn_cast<ShapedType>();
-  if (!shapedType || !shapedType.isa<MemRefType, RankedTensorType>())
-    return parser.emitError(typesLoc, "requires memref or ranked tensor type");
-  VectorType vectorType = types[1].dyn_cast<VectorType>();
-  if (!vectorType)
-    return parser.emitError(typesLoc, "requires vector type");
-  auto permutationAttrName = TransferReadOp::getPermutationMapAttrStrName();
-  Attribute mapAttr = result.attributes.get(permutationAttrName);
-  if (!mapAttr) {
-    auto permMap = getTransferMinorIdentityMap(shapedType, vectorType);
-    // Update `mapAttr` that is used later to determine mask type.
-    mapAttr = AffineMapAttr::get(permMap);
-    result.attributes.set(permutationAttrName, mapAttr);
-  }
-  if (parser.resolveOperand(sourceInfo, shapedType, result.operands) ||
-      parser.resolveOperands(indexInfo, indexType, result.operands) ||
-      parser.resolveOperand(paddingInfo, shapedType.getElementType(),
-                            result.operands))
-    return failure();
-  if (hasMask.succeeded()) {
-    if (shapedType.getElementType().dyn_cast<VectorType>())
-      return parser.emitError(
-          maskInfo.location, "does not support masks with vector element type");
-    auto map = mapAttr.dyn_cast<AffineMapAttr>().getValue();
-    // Instead of adding the mask type as an op type, compute it based on the
-    // vector type and the permutation map (to keep the type signature small).
-    auto maskType = mlir::vector::detail::transferMaskType(vectorType, map);
-    if (parser.resolveOperand(maskInfo, maskType, result.operands))
-      return failure();
+  SmallVector<NamedAttribute> components{attrs.getValue().begin(),
+                                         attrs.getValue().end()};
+  // If there is an inBoundsAttr and it all false, it is conservatively elided
+  // from the dictionary.
+  if (!permutationMapAttr ||
+      !permutationMapAttr.getAffineMap().isMinorIdentity()) {
+    elidedAttrs.push_back(xferOp.getPermutationMapAttrStrName());
+  }
+  // If there is an inBoundsAttr and it all false, it is conservatively elided
+  // from the dictionary.
+  if (!inBoundsAttr || llvm::all_of(inBoundsAttr.getAsRange<BoolAttr>(),
+                                    [](BoolAttr b) { return !b.getValue(); })) {
+    elidedAttrs.push_back(xferOp.getInBoundsAttrStrName());
+  }
+  printer.printOptionalAttrDict(op->getAttrs(), elidedAttrs);
+}
+
+static ParseResult parseCustomTransferAttrs(OpAsmParser &parser,
+                                            AffineMapAttr &permutationMapAttr,
+                                            ArrayAttr &inBoundsAttr,
+                                            NamedAttrList &attrs) {
+  if (failed(parser.parseOptionalAttrDict(attrs)))
+    return failure();
+  if (auto attr = attrs.get("permutation_map")) {
+    permutationMapAttr = attr.dyn_cast_or_null<AffineMapAttr>();
+    attrs.erase("permutation_map");
+  }
+  if (auto attr = attrs.get("in_bounds")) {
+    inBoundsAttr = attr.dyn_cast_or_null<ArrayAttr>();
+    attrs.erase("in_bounds");
   }
-  result.addAttribute(TransferReadOp::getOperandSegmentSizeAttr(),
-                      builder.getDenseI32ArrayAttr(
-                          {1, static_cast<int32_t>(indexInfo.size()), 1,
-                           static_cast<int32_t>(hasMask.succeeded())}));
-  return parser.addTypeToList(vectorType, result.types);
+  return success();
 }
 
 LogicalResult TransferReadOp::verify() {
   // Consistency of elemental types in source and vector.
   ShapedType shapedType = getShapedType();
   VectorType vectorType = getVectorType();
   VectorType maskType = getMaskType();
   auto paddingType = getPadding().getType();
-  auto permutationMap = getPermutationMap();
+  auto permutationMap = getPermutationMapOrMinorIdentity();
   auto sourceElementType = shapedType.getElementType();
 
   if (static_cast<int64_t>(getIndices().size()) != shapedType.getRank())
     return emitOpError("requires ") << shapedType.getRank() << " indices";
 
   if (failed(verifyTransferOp(cast<VectorTransferOpInterface>(getOperation()),
                               shapedType, vectorType, maskType, permutationMap,
                               getInBounds() ? *getInBounds() : ArrayAttr())))

 }
 
 template <typename TransferOp>
 static LogicalResult foldTransferInBoundsAttribute(TransferOp op) {
   // TODO: support 0-d corner case.
   // TODO: Be less conservative.
   if (op.getTransferRank() == 0)
     return failure();
-  AffineMap permutationMap = op.getPermutationMap();
+  AffineMap permutationMap = op.getPermutationMapOrMinorIdentity();
   bool changed = false;
   SmallVector<bool, 4> newInBounds;
   newInBounds.reserve(op.getTransferRank());
   for (unsigned i = 0; i < op.getTransferRank(); ++i) {
     // Already marked as in-bounds, nothing to see here.
     if (op.isDimInBounds(i)) {
       newInBounds.push_back(true);
       continue;
public:
 
   LogicalResult matchAndRewrite(TransferReadOp xferOp,
                                 PatternRewriter &rewriter) const override {
     // TODO: support 0-d corner case.
     if (xferOp.getTransferRank() == 0)
       return failure();
     if (xferOp.hasOutOfBoundsDim())
       return failure();
-    if (!xferOp.getPermutationMap().isMinorIdentity())
+    if (!xferOp.getPermutationMapOrMinorIdentity().isMinorIdentity())
       return failure();
     if (xferOp.getMask())
       return failure();
     auto extractOp = xferOp.getSource().getDefiningOp<tensor::ExtractSliceOp>();
     if (!extractOp)
       return failure();
     if (!extractOp.hasUnitStride())
       return failure();
if (readOp.getIndices() == defWrite.getIndices() &&
         vec = defWrite.getVector();
     }
     // TODO: loop through the chain of transfer_write if we can prove that they
     // don't overlap with the transfer_read. This requires improving
     // `isDisjointTransferIndices` helper.
     if (!vec)
       return failure();
     SmallVector<unsigned> permutation;
-    AffineMap readMap = compressUnusedDims(readOp.getPermutationMap());
-    AffineMap writeMap = compressUnusedDims(defWrite.getPermutationMap());
+    AffineMap readMap =
+        compressUnusedDims(readOp.getPermutationMapOrMinorIdentity());
+    AffineMap writeMap =
+        compressUnusedDims(defWrite.getPermutationMapOrMinorIdentity());
     AffineMap map = readMap.compose(writeMap);
     if (map.getNumResults() == 0)
       return failure();
     // Calculate the permuation to apply to go from the vector stored to the
     // vector read.
     if (!map.isPermutationOfMinorIdentityWithBroadcasting(permutation))
       return failure();
 
static LogicalResult foldReadInitWrite(TransferWriteOp write,
   // If no read, bail.
   auto read = write.getVector().getDefiningOp<vector::TransferReadOp>();
   if (!read)
     return failure();
   // TODO: support 0-d corner case.
   if (read.getTransferRank() == 0)
     return failure();
   // For now, only accept minor identity. Future: composition is minor identity.
-  if (!read.getPermutationMap().isMinorIdentity() ||
-      !write.getPermutationMap().isMinorIdentity())
+  if (!read.getPermutationMapOrMinorIdentity().isMinorIdentity() ||
+      !write.getPermutationMapOrMinorIdentity().isMinorIdentity())
     return failure();
   // Bail on mismatching ranks.
   if (read.getTransferRank() != write.getTransferRank())
     return failure();
   // Bail on potential out-of-bounds accesses.
   if (read.hasOutOfBoundsDim() || write.hasOutOfBoundsDim())
     return failure();
   // Tensor types must be the same.
static LogicalResult foldReadInitWrite(TransferWriteOp write,
   results.push_back(read.getSource());
   return success();
 }
 
 static bool checkSameValueWAR(vector::TransferReadOp read,
                               vector::TransferWriteOp write) {
   return read.getSource() == write.getSource() &&
          read.getIndices() == write.getIndices() &&
-         read.getPermutationMap() == write.getPermutationMap() &&
+         read.getPermutationMapOrMinorIdentity() ==
+             write.getPermutationMapOrMinorIdentity() &&
          read.getVectorType() == write.getVectorType() && !read.getMask() &&
          !write.getMask();
 }
 /// Fold transfer_write write after read:
 /// ```
 ///    %t0 = ...
 ///    %v = vector.transfer_read %t0[%c0...] :
 ///      tensor<static_sizesxf32>, vector<static_sizesxf32>
LogicalResult matchAndRewrite(tensor::InsertSliceOp insertOp,
     if (xferOp.getMask())
       return failure();
     // Fold only if the TransferWriteOp completely overwrites the `source` with
     // a vector. I.e., the result of the TransferWriteOp is a new tensor whose
     // content is the data of the vector.
     if (!llvm::equal(xferOp.getVectorType().getShape(),
                      xferOp.getShapedType().getShape()))
       return failure();
-    if (!xferOp.getPermutationMap().isIdentity())
+    if (!xferOp.getPermutationMapOrMinorIdentity().isIdentity())
       return failure();
 
     // Bail on illegal rank-reduction: we need to check that the rank-reduced
     // dims are exactly the leading dims. I.e. the following is illegal:
     // ```
     //    %0 = vector.transfer_write %v, %t[0,0], %cst :
     //      vector<2x4xf32>, tensor<2x4xf32>
     //    %1 = tensor.insert_slice %0 into %tt[0,0,0][2,1,4][1,1,1] :
for (const auto &it :
             insertOp, "InsertSliceOp and ExtractSliceOp sizes differ");
       }
     }
 
     // Fail if the vector::TransferWriteOp may not overwrite the full tensor.
     assert(transferOp.getVectorType().hasStaticShape() &&
            "expected vector to have a static shape");
     ArrayRef<int64_t> vectorShape = transferOp.getVectorType().getShape();
-    SmallVector<int64_t> resultShape = applyPermutationMap(
-        transferOp.getPermutationMap(), transferOp.getShapedType().getShape());
+    SmallVector<int64_t> resultShape =
+        applyPermutationMap(transferOp.getPermutationMapOrMinorIdentity(),
+                            transferOp.getShapedType().getShape());
     if (transferOp.getMask() || !vectorShape.equals(resultShape)) {
       return rewriter.notifyMatchFailure(
           insertOp, "TransferWriteOp may not write the full tensor.");
     }
 
     // Swap the tensor::ExtractSliceOp in front of the vector::TransferWriteOp.
-    SmallVector<int64_t> newResultShape = applyPermutationMap(
-        transferOp.getPermutationMap(), insertOp.getSourceType().getShape());
+    SmallVector<int64_t> newResultShape =
+        applyPermutationMap(transferOp.getPermutationMapOrMinorIdentity(),
+                            insertOp.getSourceType().getShape());
     SmallVector<bool> newInBounds;
     for (const auto &en : enumerate(newResultShape))
       newInBounds.push_back(en.value() == vectorShape[en.index()]);
     auto newExtractOp = rewriter.create<tensor::ExtractSliceOp>(
         extractOp.getLoc(), insertOp.getSourceType(), insertOp.getDest(),
         insertOp.getMixedOffsets(), insertOp.getMixedSizes(),
         insertOp.getMixedStrides());
     auto newTransferWriteOp = rewriter.create<TransferWriteOp>(

mlir/lib/Dialect/Vector/Transforms/BufferizableOpInterfaceImpl.cpp

LogicalResult bufferize(Operation *op, RewriterBase &rewriter,
     auto readOp = cast<vector::TransferReadOp>(op);
     assert(readOp.getShapedType().isa<TensorType>() &&
            "only tensor types expected");
     FailureOr<Value> buffer = getBuffer(rewriter, readOp.getSource(), options);
     if (failed(buffer))
       return failure();
     replaceOpWithNewBufferizedOp<vector::TransferReadOp>(
         rewriter, readOp, readOp.getVectorType(), *buffer, readOp.getIndices(),
-        readOp.getPermutationMap(), readOp.getPadding(), readOp.getMask(),
-        readOp.getInBoundsAttr());
+        AffineMapAttr::get(readOp.getPermutationMapOrMinorIdentity()),
+        readOp.getPadding(), readOp.getMask(), readOp.getInBoundsAttr());
     return success();
   }
 };
 
 /// Bufferization of vector.transfer_write. Replace with a new
 /// vector.transfer_write that operates on a memref.
 ///
 /// Note: DstBufferizableOpInterfaceExternalModel provides many default method

mlir/lib/Dialect/Vector/Transforms/LowerVectorMask.cpp

matchAndRewriteMaskableOp(TransferReadOp readOp, MaskingOpInterface maskingOp,
     // the passthru is a broadcast of a scalar.
     if (maskingOp.hasPassthru())
       return rewriter.notifyMatchFailure(
           maskingOp, "Can't lower passthru to vector.transfer_read");
 
     // Replace the `vector.mask` operation.
     rewriter.replaceOpWithNewOp<TransferReadOp>(
         maskingOp.getOperation(), readOp.getVectorType(), readOp.getSource(),
-        readOp.getIndices(), readOp.getPermutationMap(), readOp.getPadding(),
-        maskingOp.getMask(), readOp.getInBounds().value_or(ArrayAttr()));
+        readOp.getIndices(),
+        AffineMapAttr::get(readOp.getPermutationMapOrMinorIdentity()),
+        readOp.getPadding(), maskingOp.getMask(),
+        readOp.getInBounds().value_or(ArrayAttr()));
     return success();
   }
 };
 
 /// Lowers a masked `vector.transfer_write` operation.
 struct MaskedTransferWriteOpPattern
     : public MaskOpRewritePattern<TransferWriteOp> {
 public:
   using MaskOpRewritePattern<TransferWriteOp>::MaskOpRewritePattern;
 
   LogicalResult
   matchAndRewriteMaskableOp(TransferWriteOp writeOp,
                             MaskingOpInterface maskingOp,
                             PatternRewriter &rewriter) const override {
     Type resultType =
         writeOp.getResult() ? writeOp.getResult().getType() : Type();
 
     // Replace the `vector.mask` operation.
     rewriter.replaceOpWithNewOp<TransferWriteOp>(
         maskingOp.getOperation(), resultType, writeOp.getVector(),
-        writeOp.getSource(), writeOp.getIndices(), writeOp.getPermutationMap(),
-        maskingOp.getMask(), writeOp.getInBounds().value_or(ArrayAttr()));
+        writeOp.getSource(), writeOp.getIndices(),
+        writeOp.getPermutationMapOrMinorIdentity(), maskingOp.getMask(),
+        writeOp.getInBounds().value_or(ArrayAttr()));
     return success();
   }
 };
 
 /// Populates instances of `MaskOpRewritePattern` to lower masked operations
 /// with `vector.mask`. Patterns should rewrite the `vector.mask` operation and
 /// not its nested `MaskableOpInterface`.
 void populateVectorMaskLoweringPatternsForSideEffectingOps(

mlir/lib/Dialect/Vector/Transforms/VectorDistribute.cpp

LogicalResult tryDistributeOp(RewriterBase &rewriter,
     // the rest.
     vector::TransferWriteOp newWriteOp =
         cloneWriteOp(rewriter, warpOp, writeOp, targetType);
 
     // 4. Reindex the write using the distribution map.
     auto newWarpOp =
         newWriteOp.getVector().getDefiningOp<WarpExecuteOnLane0Op>();
     rewriter.setInsertionPoint(newWriteOp);
-    AffineMap indexMap = map.compose(newWriteOp.getPermutationMap());
+    AffineMap indexMap = map.compose(newWriteOp.getPermutationMapOrMinorIdentity());
     Location loc = newWriteOp.getLoc();
     SmallVector<Value> indices(newWriteOp.getIndices().begin(),
                                newWriteOp.getIndices().end());
     for (auto it : llvm::zip(indexMap.getResults(), map.getResults())) {
       AffineExpr d0, d1;
       bindDims(newWarpOp.getContext(), d0, d1);
       auto indexExpr = std::get<0>(it).dyn_cast<AffineDimExpr>();
       if (!indexExpr)
LogicalResult matchAndRewrite(WarpExecuteOnLane0Op warpOp,
     unsigned operandIndex = operand->getOperandNumber();
     Value distributedVal = warpOp.getResult(operandIndex);
 
     SmallVector<Value, 4> indices(read.getIndices().begin(),
                                   read.getIndices().end());
     auto sequentialType = read.getResult().getType().cast<VectorType>();
     auto distributedType = distributedVal.getType().cast<VectorType>();
     AffineMap map = calculateImplicitMap(sequentialType, distributedType);
-    AffineMap indexMap = map.compose(read.getPermutationMap());
+    AffineMap indexMap = map.compose(read.getPermutationMapOrMinorIdentity());
     OpBuilder::InsertionGuard g(rewriter);
     rewriter.setInsertionPointAfter(warpOp);
     for (auto it : llvm::zip(indexMap.getResults(), map.getResults())) {
       AffineExpr d0, d1;
       bindDims(read.getContext(), d0, d1);
       auto indexExpr = std::get<0>(it).dyn_cast<AffineDimExpr>();
       if (!indexExpr)
         continue;

mlir/lib/Dialect/Vector/Transforms/VectorDropLeadUnitDim.cpp

if (shapedType.getElementType() != read.getVectorType().getElementType())
       return failure();
 
     VectorType oldType = read.getVectorType();
     VectorType newType = trimLeadingOneDims(oldType);
 
     if (newType == oldType)
       return failure();
 
-    AffineMap oldMap = read.getPermutationMap();
+    AffineMap oldMap = read.getPermutationMapOrMinorIdentity();
     ArrayRef<AffineExpr> newResults =
         oldMap.getResults().take_back(newType.getRank());
     AffineMap newMap =
         AffineMap::get(oldMap.getNumDims(), oldMap.getNumSymbols(), newResults,
                        rewriter.getContext());
 
     ArrayAttr inBoundsAttr;
     if (read.getInBounds())
if (shapedType.getElementType() != write.getVectorType().getElementType())
       return failure();
 
     VectorType oldType = write.getVectorType();
     VectorType newType = trimLeadingOneDims(oldType);
     if (newType == oldType)
       return failure();
     int64_t dropDim = oldType.getRank() - newType.getRank();
 
-    AffineMap oldMap = write.getPermutationMap();
+    AffineMap oldMap = write.getPermutationMapOrMinorIdentity();
     ArrayRef<AffineExpr> newResults =
         oldMap.getResults().take_back(newType.getRank());
     AffineMap newMap =
         AffineMap::get(oldMap.getNumDims(), oldMap.getNumSymbols(), newResults,
                        rewriter.getContext());
 
     ArrayAttr inBoundsAttr;
     if (write.getInBounds())

mlir/lib/Dialect/Vector/Transforms/VectorTransferOpTransforms.cpp

LogicalResult matchAndRewrite(vector::TransferReadOp transferReadOp,
     // TODO: support tensor types.
     if (!sourceType || !sourceType.hasStaticShape())
       return failure();
     if (sourceType.getNumElements() != vectorType.getNumElements())
       return failure();
     // TODO: generalize this pattern, relax the requirements here.
     if (transferReadOp.hasOutOfBoundsDim())
       return failure();
-    if (!transferReadOp.getPermutationMap().isMinorIdentity())
+    if (!transferReadOp.getPermutationMapOrMinorIdentity().isMinorIdentity())
       return failure();
     int reducedRank = getReducedRank(sourceType.getShape());
     if (reducedRank == sourceType.getRank())
       return failure(); // The source shape can't be further reduced.
     if (reducedRank != vectorType.getRank())
       return failure(); // This pattern requires the vector shape to match the
                         // reduced source shape.
     if (llvm::any_of(transferReadOp.getIndices(),
LogicalResult matchAndRewrite(vector::TransferWriteOp transferWriteOp,
     // TODO: support tensor type.
     if (!sourceType || !sourceType.hasStaticShape())
       return failure();
     if (sourceType.getNumElements() != vectorType.getNumElements())
       return failure();
     // TODO: generalize this pattern, relax the requirements here.
     if (transferWriteOp.hasOutOfBoundsDim())
       return failure();
-    if (!transferWriteOp.getPermutationMap().isMinorIdentity())
+    if (!transferWriteOp.getPermutationMapOrMinorIdentity().isMinorIdentity())
       return failure();
     int reducedRank = getReducedRank(sourceType.getShape());
     if (reducedRank == sourceType.getRank())
       return failure(); // The source shape can't be further reduced.
     if (reducedRank != vectorType.getRank())
       return failure(); // This pattern requires the vector shape to match the
                         // reduced source shape.
     if (llvm::any_of(transferWriteOp.getIndices(),
if (!hasMatchingInnerContigousShape(
             sourceType,
             vectorType.getShape().take_back(vectorType.getRank() - 1)))
       return failure();
     int64_t firstContiguousInnerDim =
         sourceType.getRank() - vectorType.getRank();
     // TODO: generalize this pattern, relax the requirements here.
     if (transferReadOp.hasOutOfBoundsDim())
       return failure();
-    if (!transferReadOp.getPermutationMap().isMinorIdentity())
+    if (!transferReadOp.getPermutationMapOrMinorIdentity().isMinorIdentity())
       return failure();
     if (transferReadOp.getMask())
       return failure();
     SmallVector<Value> collapsedIndices;
     if (failed(checkAndCollapseInnerZeroIndices(transferReadOp.getIndices(),
                                                 firstContiguousInnerDim,
                                                 collapsedIndices)))
       return failure();
if (!hasMatchingInnerContigousShape(
             sourceType,
             vectorType.getShape().take_back(vectorType.getRank() - 1)))
       return failure();
     int64_t firstContiguousInnerDim =
         sourceType.getRank() - vectorType.getRank();
     // TODO: generalize this pattern, relax the requirements here.
     if (transferWriteOp.hasOutOfBoundsDim())
       return failure();
-    if (!transferWriteOp.getPermutationMap().isMinorIdentity())
+    if (!transferWriteOp.getPermutationMapOrMinorIdentity().isMinorIdentity())
       return failure();
     if (transferWriteOp.getMask())
       return failure();
     SmallVector<Value> collapsedIndices;
     if (failed(checkAndCollapseInnerZeroIndices(transferWriteOp.getIndices(),
                                                 firstContiguousInnerDim,
                                                 collapsedIndices)))
       return failure();

mlir/lib/Dialect/Vector/Transforms/VectorTransferPermutationMapRewritePatterns.cpp

struct TransferReadPermutationLowering
 
   LogicalResult matchAndRewrite(vector::TransferReadOp op,
                                 PatternRewriter &rewriter) const override {
     // TODO: support 0-d corner case.
     if (op.getTransferRank() == 0)
       return failure();
 
     SmallVector<unsigned> permutation;
-    AffineMap map = op.getPermutationMap();
+    AffineMap map = op.getPermutationMapOrMinorIdentity();
     if (map.getNumResults() == 0)
       return failure();
     if (!map.isPermutationOfMinorIdentityWithBroadcasting(permutation))
       return failure();
     AffineMap permutationMap =
         map.getPermutationMap(permutation, op.getContext());
     if (permutationMap.isIdentity())
       return failure();
struct TransferWritePermutationLowering
 
   LogicalResult matchAndRewrite(vector::TransferWriteOp op,
                                 PatternRewriter &rewriter) const override {
     // TODO: support 0-d corner case.
     if (op.getTransferRank() == 0)
       return failure();
 
     SmallVector<unsigned> permutation;
-    AffineMap map = op.getPermutationMap();
+    AffineMap map = op.getPermutationMapOrMinorIdentity();
     if (map.isMinorIdentity())
       return failure();
     if (!map.isPermutationOfMinorIdentityWithBroadcasting(permutation))
       return failure();
 
     // Remove unused dims from the permutation map. E.g.:
     // E.g.:  (d0, d1, d2, d3, d4, d5) -> (d5, d3, d4)
     // comp = (d0, d1, d2) -> (d2, d0, d1)
struct TransferOpReduceRank : public OpRewritePattern<vector::TransferReadOp> {
   using OpRewritePattern::OpRewritePattern;
 
   LogicalResult matchAndRewrite(vector::TransferReadOp op,
                                 PatternRewriter &rewriter) const override {
     // TODO: support 0-d corner case.
     if (op.getTransferRank() == 0)
       return failure();
 
-    AffineMap map = op.getPermutationMap();
+    AffineMap map = op.getPermutationMapOrMinorIdentity();
     unsigned numLeadingBroadcast = 0;
     for (auto expr : map.getResults()) {
       auto dimExpr = expr.dyn_cast<AffineConstantExpr>();
       if (!dimExpr || dimExpr.getValue() != 0)
         break;
       numLeadingBroadcast++;
     }
     // If there are no leading zeros in the map there is nothing to do.

mlir/lib/Dialect/Vector/Transforms/VectorTransforms.cpp

LogicalResult matchAndRewrite(vector::TransferReadOp read,
                                 PatternRewriter &rewriter) const override {
     if (maxTransferRank && read.getVectorType().getRank() > *maxTransferRank)
       return failure();
 
     SmallVector<unsigned> broadcastedDims;
     // Permutations are handled by VectorToSCF or
     // populateVectorTransferPermutationMapLoweringPatterns.
     // We let the 0-d corner case pass-through as it is supported.
-    if (!read.getPermutationMap().isMinorIdentityWithBroadcasting(
+    if (!read.getPermutationMapOrMinorIdentity().isMinorIdentityWithBroadcasting(
             &broadcastedDims))
       return failure();
 
     auto memRefType = read.getShapedType().dyn_cast<MemRefType>();
     if (!memRefType)
       return failure();
 
     // Non-unit strides are handled by VectorToSCF.
LogicalResult matchAndRewrite(vector::TransferWriteOp write,
     if (maxTransferRank && write.getVectorType().getRank() > *maxTransferRank)
       return rewriter.notifyMatchFailure(write.getLoc(), [=](Diagnostic &diag) {
         diag << "rank exceeds maxTransferRank: " << write;
       });
 
     // Permutations are handled by VectorToSCF or
     // populateVectorTransferPermutationMapLoweringPatterns.
     if ( // pass-through for the 0-d corner case.
-        !write.getPermutationMap().isMinorIdentity())
+        !write.getPermutationMapOrMinorIdentity().isMinorIdentity())
       return rewriter.notifyMatchFailure(write.getLoc(), [=](Diagnostic &diag) {
         diag << "permutation map is not minor identity: " << write;
       });
 
     auto memRefType = write.getShapedType().dyn_cast<MemRefType>();
     if (!memRefType)
       return rewriter.notifyMatchFailure(write.getLoc(), [=](Diagnostic &diag) {
         diag << "not a memref type: " << write;
LogicalResult matchAndRewrite(vector::TransferReadOp readOp,
     // TODO: support mask.
     if (readOp.getMask())
       return failure();
 
     auto srcType = readOp.getSource().getType().dyn_cast<MemRefType>();
     if (!srcType || !srcType.hasStaticShape())
       return failure();
 
-    if (!readOp.getPermutationMap().isMinorIdentity())
+    if (!readOp.getPermutationMapOrMinorIdentity().isMinorIdentity())
       return failure();
 
     auto targetType = readOp.getVectorType();
     if (targetType.getRank() <= 1)
       return failure();
 
     SmallVector<int64_t> srcStrides;
     int64_t srcOffset;

mlir/lib/Dialect/Vector/Transforms/VectorUnroll.cpp

LogicalResult matchAndRewrite(vector::TransferReadOp readOp,
     SmallVector<int64_t> loopOrder =
         getUnrollOrder(originalSize.size(), readOp, options);
     DecomposeShapeIterator indexToOffsets(originalSize, *targetShape,
                                           loopOrder);
     for (int64_t i = 0; i < indexToOffsets.maxIndex(); i++) {
       SmallVector<int64_t> elementOffsets = indexToOffsets.getVectorOffset(i);
       SmallVector<Value> indices =
           sliceTransferIndices(elementOffsets, originalIndices,
-                               readOp.getPermutationMap(), loc, rewriter);
+                               readOp.getPermutationMapOrMinorIdentity(), loc, rewriter);
       auto slicedRead = rewriter.create<vector::TransferReadOp>(
           loc, targetType, readOp.getSource(), indices,
           readOp.getPermutationMapAttr(), readOp.getPadding(), readOp.getMask(),
           readOp.getInBoundsAttr());
 
       result = rewriter.create<vector::InsertStridedSliceOp>(
           loc, slicedRead, result, elementOffsets, strides);
     }
DecomposeShapeIterator indexToOffsets(originalSize, *targetShape,
                                           loopOrder);
     Value resultTensor;
     for (int64_t i = 0; i < indexToOffsets.maxIndex(); i++) {
       SmallVector<int64_t> elementOffsets = indexToOffsets.getVectorOffset(i);
       Value slicedVector = rewriter.create<vector::ExtractStridedSliceOp>(
           loc, writeOp.getVector(), elementOffsets, *targetShape, strides);
       SmallVector<Value> indices =
           sliceTransferIndices(elementOffsets, originalIndices,
-                               writeOp.getPermutationMap(), loc, rewriter);
+                               writeOp.getPermutationMapOrMinorIdentity(), loc, rewriter);
       Operation *slicedWrite = rewriter.create<vector::TransferWriteOp>(
           loc, slicedVector, resultTensor ? resultTensor : writeOp.getSource(),
           indices, writeOp.getPermutationMapAttr(), writeOp.getInBoundsAttr());
       // For the tensor case update the destination for the next transfer write.
       if (!slicedWrite->getResults().empty())
         resultTensor = slicedWrite->getResult(0);
     }
     if (resultTensor)

mlir/test/Conversion/VectorToLLVM/vector-to-llvm.mlir

 // CHECK: %[[mask1:.*]] = arith.constant dense<[false, false, true, false, true]>
 // CHECK: %[[cmpi:.*]] = arith.cmpi slt
 // CHECK: %[[mask2:.*]] = arith.andi %[[cmpi]], %[[mask1]]
 // CHECK: %[[r:.*]] = llvm.intr.masked.load %{{.*}}, %[[mask2]]
 // CHECK: return %[[r]]
 func.func @transfer_read_1d_mask(%A : memref<?xf32>, %base : index) -> vector<5xf32> {
   %m = arith.constant dense<[0, 0, 1, 0, 1]> : vector<5xi1>
   %f7 = arith.constant 7.0: f32
-  %f = vector.transfer_read %A[%base], %f7, %m : memref<?xf32>, vector<5xf32>
+  %f = vector.transfer_read %A[%base], %f7, %m : vector<5xi1> : memref<?xf32>, vector<5xf32>
   return %f: vector<5xf32>
 }
 
 // -----
 
 func.func @genbool_0d_f() -> vector<i1> {
   %0 = vector.constant_mask [0] : vector<i1>
   return %0 : vector<i1>

mlir/test/Dialect/Vector/lower-vector-mask.mlir

 // RUN: mlir-opt -lower-vector-mask -split-input-file %s | FileCheck %s
 
 func.func @vector_transfer_read(%t0: tensor<?xf32>, %idx: index, %m0: vector<16xi1>) -> vector<16xf32> {
   %ft0 = arith.constant 0.0 : f32
   %0 = vector.mask %m0 { vector.transfer_read %t0[%idx], %ft0 : tensor<?xf32>, vector<16xf32> } : vector<16xi1> -> vector<16xf32>
   return %0 : vector<16xf32>
 }
 
 // CHECK-LABEL:   func.func @vector_transfer_read(
 // CHECK-SAME:                                    %[[VAL_0:.*]]: tensor<?xf32>,
 // CHECK-SAME:                                    %[[VAL_1:.*]]: index,
 // CHECK-SAME:                                    %[[VAL_2:.*]]: vector<16xi1>) -> vector<16xf32> {
 // CHECK-NOT:       vector.mask
-// CHECK:           %[[VAL_4:.*]] = vector.transfer_read {{.*}}, %[[VAL_2]] : tensor<?xf32>, vector<16xf32>
+// CHECK:           %[[VAL_4:.*]] = vector.transfer_read {{.*}}, %[[VAL_2]] : vector<16xi1>{{.*}} : tensor<?xf32>, vector<16xf32>
 // CHECK:           return %[[VAL_4]] : vector<16xf32>
 // CHECK:         }
 
 // -----
 
 func.func @vector_transfer_write_on_memref(%val: vector<16xf32>, %t0: memref<?xf32>, %idx: index, %m0: vector<16xi1>) {
   vector.mask %m0 { vector.transfer_write %val, %t0[%idx] : vector<16xf32>, memref<?xf32> } : vector<16xi1>
   return

mlir/test/Dialect/Vector/ops.mlir

func.func @vector_transfer_ops(%arg0: memref<?x?xf32>,
   %4 = vector.transfer_read %arg1[%c3, %c3], %vf0 {permutation_map = affine_map<(d0, d1)->(d0, d1)>} : memref<?x?xvector<4x3xf32>>, vector<1x1x4x3xf32>
   // CHECK: vector.transfer_read %{{.*}}[%[[C3]], %[[C3]]], %{{.*}} {in_bounds = [false, true]} : memref<?x?xvector<4x3xf32>>, vector<1x1x4x3xf32>
   %5 = vector.transfer_read %arg1[%c3, %c3], %vf0 {in_bounds = [false, true]} : memref<?x?xvector<4x3xf32>>, vector<1x1x4x3xf32>
   // CHECK: vector.transfer_read %{{.*}}[%[[C3]], %[[C3]]], %{{.*}} : memref<?x?xvector<4x3xi32>>, vector<5x24xi8>
   %6 = vector.transfer_read %arg2[%c3, %c3], %v0 : memref<?x?xvector<4x3xi32>>, vector<5x24xi8>
   // CHECK: vector.transfer_read %{{.*}}[%[[C3]], %[[C3]]], %{{.*}} : memref<?x?xvector<4x3xindex>>, vector<5x48xi8>
   %7 = vector.transfer_read %arg3[%c3, %c3], %vi0 : memref<?x?xvector<4x3xindex>>, vector<5x48xi8>
   // CHECK: vector.transfer_read %{{.*}}[%[[C3]], %[[C3]]], %{{.*}}, %{{.*}} : memref<?x?xf32>, vector<5xf32>
-  %8 = vector.transfer_read %arg0[%c3, %c3], %f0, %m : memref<?x?xf32>, vector<5xf32>
+  %8 = vector.transfer_read %arg0[%c3, %c3], %f0, %m : vector<5xi1> : memref<?x?xf32>, vector<5xf32>
   // CHECK: vector.transfer_read %{{.*}}[%[[C3]], %[[C3]], %[[C3]]], %{{.*}}, %{{.*}} : memref<?x?x?xf32>, vector<5x4x8xf32>
-  %9 = vector.transfer_read %arg4[%c3, %c3, %c3], %f0, %m2 {permutation_map = affine_map<(d0, d1, d2)->(d1, d0, 0)>} : memref<?x?x?xf32>, vector<5x4x8xf32>
+  %9 = vector.transfer_read %arg4[%c3, %c3, %c3], %f0, %m : vector<5xi1> {permutation_map = affine_map<(d0, d1, d2)->(d1, d0, 0)>} : memref<?x?x?xf32>, vector<5x4x8xf32>
 
   // CHECK: vector.transfer_write
   vector.transfer_write %0, %arg0[%c3, %c3] {permutation_map = affine_map<(d0, d1)->(d0)>} : vector<128xf32>, memref<?x?xf32>
   // CHECK: vector.transfer_write
   vector.transfer_write %1, %arg0[%c3, %c3] {permutation_map = affine_map<(d0, d1)->(d1, d0)>} : vector<3x7xf32>, memref<?x?xf32>
   // CHECK: vector.transfer_write %{{.*}}, %{{.*}}[%[[C3]], %[[C3]]] : vector<1x1x4x3xf32>, memref<?x?xvector<4x3xf32>>
   vector.transfer_write %4, %arg1[%c3, %c3] {permutation_map = affine_map<(d0, d1)->(d0, d1)>} : vector<1x1x4x3xf32>, memref<?x?xvector<4x3xf32>>
   // CHECK: vector.transfer_write %{{.*}}, %{{.*}}[%[[C3]], %[[C3]]] : vector<1x1x4x3xf32>, memref<?x?xvector<4x3xf32>>

mlir/test/Integration/Dialect/Vector/CPU/test-transfer-read-2d.mlir

 
 // Vector load with mask.
 func.func @transfer_read_2d_mask(%A : memref<?x?xf32>, %base1: index, %base2: index) {
   %fm42 = arith.constant -42.0: f32
   %mask = arith.constant dense<[[1, 0, 1, 0, 1, 1, 1, 0, 1],
                           [0, 0, 1, 1, 1, 1, 1, 0, 1],
                           [1, 1, 1, 1, 1, 1, 1, 0, 1],
                           [0, 0, 1, 0, 1, 1, 1, 0, 1]]> : vector<4x9xi1>
-  %f = vector.transfer_read %A[%base1, %base2], %fm42, %mask
+  %f = vector.transfer_read %A[%base1, %base2], %fm42, %mask : vector<4x9xi1>
       {permutation_map = affine_map<(d0, d1) -> (d0, d1)>} :
     memref<?x?xf32>, vector<4x9xf32>
   vector.print %f: vector<4x9xf32>
   return
 }
 
 // Vector load with mask + transpose.
 func.func @transfer_read_2d_mask_transposed(
     %A : memref<?x?xf32>, %base1: index, %base2: index) {
   %fm42 = arith.constant -42.0: f32
   %mask = arith.constant dense<[[1, 0, 1, 0], [0, 0, 1, 0],
                           [1, 1, 1, 1], [0, 1, 1, 0],
                           [1, 1, 1, 1], [1, 1, 1, 1],
                           [1, 1, 1, 1], [0, 0, 0, 0],
                           [1, 1, 1, 1]]> : vector<9x4xi1>
-  %f = vector.transfer_read %A[%base1, %base2], %fm42, %mask
+  %f = vector.transfer_read %A[%base1, %base2], %fm42, %mask : vector<9x4xi1>
       {permutation_map = affine_map<(d0, d1) -> (d1, d0)>} :
     memref<?x?xf32>, vector<9x4xf32>
   vector.print %f: vector<9x4xf32>
   return
 }
 
 // Vector load with mask + broadcast.
 func.func @transfer_read_2d_mask_broadcast(
     %A : memref<?x?xf32>, %base1: index, %base2: index) {
   %fm42 = arith.constant -42.0: f32
   %mask = arith.constant dense<[1, 0, 1, 0, 1, 1, 1, 0, 1]> : vector<9xi1>
-  %f = vector.transfer_read %A[%base1, %base2], %fm42, %mask
+  %f = vector.transfer_read %A[%base1, %base2], %fm42, %mask : vector<9xi1>
       {permutation_map = affine_map<(d0, d1) -> (0, d1)>} :
     memref<?x?xf32>, vector<4x9xf32>
   vector.print %f: vector<4x9xf32>
   return
 }
 
 // Transpose + vector load with mask + broadcast.
 func.func @transfer_read_2d_mask_transpose_broadcast_last_dim(
     %A : memref<?x?xf32>, %base1: index, %base2: index) {
   %fm42 = arith.constant -42.0: f32
   %mask = arith.constant dense<[1, 0, 1, 1]> : vector<4xi1>
-  %f = vector.transfer_read %A[%base1, %base2], %fm42, %mask
+  %f = vector.transfer_read %A[%base1, %base2], %fm42, %mask : vector<4xi1>
       {permutation_map = affine_map<(d0, d1) -> (d1, 0)>} :
     memref<?x?xf32>, vector<4x9xf32>
   vector.print %f: vector<4x9xf32>
   return
 }
 
 // Load + transpose.
 func.func @transfer_read_2d_transposed(

mlir/test/Integration/Dialect/Vector/CPU/test-transfer-read-3d.mlir

func.func @transfer_read_3d_broadcast(%A : memref<?x?x?x?xf32>,
   vector.print %f: vector<2x5x3xf32>
   return
 }
 
 func.func @transfer_read_3d_mask_broadcast(
     %A : memref<?x?x?x?xf32>, %o: index, %a: index, %b: index, %c: index) {
   %fm42 = arith.constant -42.0: f32
   %mask = arith.constant dense<[0, 1]> : vector<2xi1>
-  %f = vector.transfer_read %A[%o, %a, %b, %c], %fm42, %mask
+  %f = vector.transfer_read %A[%o, %a, %b, %c], %fm42, %mask : vector<2xi1>
       {permutation_map = affine_map<(d0, d1, d2, d3) -> (d1, 0, 0)>}
       : memref<?x?x?x?xf32>, vector<2x5x3xf32>
   vector.print %f: vector<2x5x3xf32>
   return
 }
 
 func.func @transfer_read_3d_transposed(%A : memref<?x?x?x?xf32>,
                                   %o: index, %a: index, %b: index, %c: index) {

mlir/test/Integration/Dialect/Vector/CPU/test-transfer-read.mlir

%f = vector.transfer_read %A[%base], %fm42
     memref<?xf32>, vector<13xf32>
   vector.print %f: vector<13xf32>
   return
 }
 
 func.func @transfer_read_mask_1d(%A : memref<?xf32>, %base: index) {
   %fm42 = arith.constant -42.0: f32
   %m = arith.constant dense<[0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 0]> : vector<13xi1>
-  %f = vector.transfer_read %A[%base], %fm42, %m : memref<?xf32>, vector<13xf32>
+  %f = vector.transfer_read %A[%base], %fm42, %m : vector<13xi1> : memref<?xf32>, vector<13xf32>
   vector.print %f: vector<13xf32>
   return
 }
 
 func.func @transfer_read_inbounds_4(%A : memref<?xf32>, %base: index) {
   %fm42 = arith.constant -42.0: f32
   %f = vector.transfer_read %A[%base], %fm42
       {permutation_map = affine_map<(d0) -> (d0)>, in_bounds = [true]} :
     memref<?xf32>, vector<4xf32>
   vector.print %f: vector<4xf32>
   return
 }
 
 func.func @transfer_read_mask_inbounds_4(%A : memref<?xf32>, %base: index) {
   %fm42 = arith.constant -42.0: f32
   %m = arith.constant dense<[0, 1, 0, 1]> : vector<4xi1>
-  %f = vector.transfer_read %A[%base], %fm42, %m {in_bounds = [true]}
+  %f = vector.transfer_read %A[%base], %fm42, %m : vector<4xi1> {in_bounds = [true]}
       : memref<?xf32>, vector<4xf32>
   vector.print %f: vector<4xf32>
   return
 }
 
 func.func @transfer_write_1d(%A : memref<?xf32>, %base: index) {
   %f0 = arith.constant 0.0 : f32
   %vf0 = vector.splat %f0 : vector<4xf32>