[mlir] [VectorOps] A "reference" lowering of vector.transpose to LLVM IR

[aartbik]

Summary

Makes the vector.tranpose runnable on CPU.

[rriddle]

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

91

You need mlir after the ticks

[aartbik]

tagged ''' with mlir where applicable

[aartbik]

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

91

like so?

[nicolasvasilache]

LG, thanks Aart!

[mehdi_amini]

mlir/lib/Dialect/Vector/VectorTransforms.cpp

867

FYI the comment seems like it wasn't updated

[aartbik]

mlir/lib/Dialect/Vector/VectorTransforms.cpp

867

Good catch. Will fix.

mlir/include/mlir/Dialect/Vector/VectorOps.h

 /// operations as well (through DCE and folding). If tuple values
 /// "leak" coming in, however, some tuple related ops will remain.
 void populateVectorSlicesLoweringPatterns(OwningRewritePatternList &patterns,
                                           MLIRContext *context);
 
 /// Collect a set of transformation patterns that are related to contracting
 /// or expanding vector operations:
 ///   ContractionOpLowering,
-///   ShapeCastOp2DDownCastRewritePattern, ShapeCastOp2DUpCastRewritePattern
+///   ShapeCastOp2DDownCastRewritePattern,
+///   ShapeCastOp2DUpCastRewritePattern
+///   TransposeOpLowering
 ///   OuterproductOpLowering
 /// These transformation express higher level vector ops in terms of more
 /// elementary extraction, insertion, reduction, product, and broadcast ops.
 void populateVectorContractLoweringPatterns(
     OwningRewritePatternList &patterns, MLIRContext *context,
     VectorTransformsOptions vectorTransformOptions = VectorTransformsOptions());
 
 /// Returns the integer type required for subscripts in the vector dialect.

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

*) "parallel": Batch dimensions are iterator type "parallel", and
                    are unrelated to each other and do not co-iterate, which
                    should be expressed in their indexing maps.
 
     An indexing map attribute list must be specified with an entry for lhs, rhs
     and acc arguments. An indexing map attribute specifies a mapping from each
     iterator in the iterator type list, to each dimension of an N-D vector.
 
     Examples:
-
+    ```mlir
       // Simple dot product (K = 0).
       #contraction_accesses = [
        affine_map<(i) -> (i)>,
        affine_map<(i) -> (i)>,
        affine_map<(i) -> ()>
       ]
       #contraction_trait = {
         indexing_maps = #contraction_accesses,
```mlir
 
       // 4D vector contraction with two contracting dimensions and optional
       // vector mask arguments.
       %lhs_mask = vector.constant_mask [7, 8, 16, 15] : vector<7x8x16x15xi1>
       %rhs_mask = vector.constant_mask [8, 16, 7, 5] : vector<8x16x7x5xi1>
 
       %5 = vector.contract #contraction_trait %0, %1, %2, %lhs_mask, %rhs_mask
          : vector<7x8x16x15xf32>, vector<8x16x7x5xf32> into vector<8x15x8x5xf32>
+    ```
   }];
   let builders = [OpBuilder<
     "Builder *builder, OperationState &result, Value lhs, Value rhs, "
     "Value acc, ArrayAttr indexingMaps, ArrayAttr iteratorTypes">,
     OpBuilder<
       "Builder *builder, OperationState &result, Value lhs, Value rhs, "
       "Value acc, ArrayRef<ArrayRef<AffineExpr>> indexingExprs, "
       "ArrayRef<StringRef> iteratorTypes">];
let description = [{
     accumulator.
 
     Note that these operations are restricted to 1-D vectors to remain
     close to the corresponding LLVM intrinsics:
 
     http://llvm.org/docs/LangRef.html#experimental-vector-reduction-intrinsics
 
     Examples:
-    ```
+    ```mlir
       %1 = vector.reduction "add", %0 : vector<16xf32> into f32
 
       %3 = vector.reduction "xor", %2 : vector<4xi32> into i32
 
       %4 = vector.reduction "mul", %0, %1 : vector<16xf32> into f32
     ```
   }];
   let extraClassDeclaration = [{
* a k-D vector <s_1 x .. x s_k x type> can be broadcast to
                <duplication>         <potential stretch>
        ```
     The source operand is duplicated over all the missing leading dimensions
     and stretched over the trailing dimensions where the source has a non-equal
     dimension of 1. These rules imply that any scalar broadcast (k=0) to any
     shaped vector with the same element type is always legal.
 
     Examples:
-    ```
+    ```mlir
       %0 = constant 0.0 : f32
       %1 = vector.broadcast %0 : f32 to vector<16xf32>
       %2 = vector.broadcast %1 : vector<16xf32> to vector<4x16xf32>
     ```
   }];
   let extraClassDeclaration = [{
     Type getSourceType() { return source().getType(); }
     VectorType getVectorType() {
* the two operands and the result must have the same rank and trailing
               v2 : <t_1 x t_2 x .. x t_k x type>
       we have s_i = t_i for all 1 < i <= k
     * the mask length equals the leading dimension size of the result
     * numbering the input vector indices left to right across the operands, all
       mask values must be within range, viz. given two k-D operands v1 and v2
       above, all mask values are in the range [0,s_1+t_1)
 
     Examples:
-    ```
+    ```mlir
     %0 = vector.shuffle %a, %b[0, 3]
                : vector<2xf32>, vector<2xf32>       ; yields vector<2xf32>
     %1 = vector.shuffle %c, %b[0, 1, 2]
                : vector<2x16xf32>, vector<1x16xf32> ; yields vector<3x16xf32>
     %2 = vector.shuffle %a, %b[3, 2, 1, 0]
                : vector<2xf32>, vector<2xf32>       ; yields vector<4xf32>
 
     ```
def Vector_ExtractElementOp :
   let description = [{
     Takes an 1-D vector and a dynamic index position and extracts the
     scalar at that position. Note that this instruction resembles
     vector.extract, but is restricted to 1-D vectors and relaxed
     to dynamic indices. It is meant to be closer to LLVM's version:
     https://llvm.org/docs/LangRef.html#extractelement-instruction
 
     Example:
-    ```
+    ```mlir
       %c = constant 15 : i32
       %1 = vector.extractelement %0[%c : i32]: vector<16xf32>
     ```
   }];
   let extraClassDeclaration = [{
     VectorType getVectorType() {
       return vector().getType().cast<VectorType>();
     }
Vector_Op<"extract", [NoSideEffect,
     Arguments<(ins AnyVector:$vector, I64ArrayAttr:$position)>,
     Results<(outs AnyType)> {
   let summary = "extract operation";
   let description = [{
     Takes an n-D vector and a k-D position and extracts the (n-k)-D vector at
     the proper position. Degenerates to an element type in the 0-D case.
 
     Examples:
-    ```
+    ```mlir
       %1 = vector.extract %0[3]: vector<4x8x16xf32>
       %2 = vector.extract %0[3, 3, 3]: vector<4x8x16xf32>
     ```
   }];
   let builders = [OpBuilder<
     "Builder *builder, OperationState &result, Value source,"
     "ArrayRef<int64_t>">];
   let extraClassDeclaration = [{
let description = [{
     'sizes' except for slices at dimension boundaries when 'vector' dimension
     sizes are not a multiple of 'sizes'.
 
     Each slice is returned at the tuple element index corresponding to the
     linear index of the slice w.r.t the unrolling scheme represented by 'sizes'.
     Currently, only unit strides are supported.
 
     Examples:
-    ```
+    ```mlir
       %0 = vector.transfer_read ...: vector<4x2xf32>
 
       %1 = vector.extract_slices %0, [2, 2], [1, 1]
         : vector<4x2xf32> into tuple<vector<2x2xf32>, vector<2x2xf32>>
 
       // Example with partial slices at dimension boundaries.
       %2 = vector.transfer_read ...: vector<4x3xf32>
 
let description = [{
     pointwise multiply-and-accumulate: `$result = `$lhs * $rhs + $acc`.
     All operands and result have the same vector type. The semantics
     of the operation correspond to those of the `llvm.fma`
     [intrinsic](https://llvm.org/docs/LangRef.html#int-fma). In the
     particular case of lowering to LLVM, this is guaranteed to lower
     to the `llvm.fma.*` intrinsic.
 
     Example:
-
-    ```
+    ```mlir
       %3 = vector.fma %0, %1, %2: vector<8x16xf32>
     ```
   }];
   // Fully specified by traits.
   let verifier = ?;
   let assemblyFormat = "$lhs `,` $rhs `,` $acc attr-dict `:` type($lhs)";
   let builders = [OpBuilder<
     "Builder *b, OperationState &result, Value lhs, Value rhs, Value acc",
let description = [{
     Takes a scalar source, an 1-D destination vector and a dynamic index
     position and inserts the source into the destination at the proper
     position.  Note that this instruction resembles vector.insert, but
     is restricted to 1-D vectors and relaxed to dynamic indices. It is
     meant to be closer to LLVM's version:
     https://llvm.org/docs/LangRef.html#insertelement-instruction
 
     Example:
-    ```
+    ```mlir
       %c = constant 15 : i32
       %f = constant 0.0f : f32
       %1 = vector.insertelement %f, %0[%c : i32]: vector<16xf32>
     ```
   }];
   let extraClassDeclaration = [{
     Type getSourceType() { return source().getType(); }
     VectorType getDestVectorType() {
Vector_Op<"insert", [NoSideEffect,
      Results<(outs AnyVector:$res)> {
   let summary = "insert operation";
   let description = [{
     Takes an n-D source vector, an (n+k)-D destination vector and a k-D position
     and inserts the n-D source into the (n+k)-D destination at the proper
     position. Degenerates to a scalar source type when n = 0.
 
     Examples:
-    ```
+    ```mlir
       %2 = vector.insert %0, %1[3]:
         vector<8x16xf32> into vector<4x8x16xf32>
       %5 = vector.insert %3, %4[3, 3, 3]:
         f32 into vector<4x8x16xf32>
     ```
   }];
   let assemblyFormat = [{
     $source `,` $dest $position attr-dict `:` type($source) `into` type($dest)
let description = [{
     'sizes' except for slices at dimension boundaries when 'vector' dimension
     sizes are not a multiple of 'sizes'.
 
     Each slice in 'vectors' is at the tuple element index corresponding to the
     linear index of the slice w.r.t the unrolling scheme represented by 'sizes'.
     Currently, only unit strides are supported.
 
     Examples:
-    ```
+    ```mlir
       %0 = vector.extract_slices %0, [2, 2], [1, 1]
         : vector<4x2xf32> into tuple<vector<2x2xf32>, vector<2x2xf32>>
 
       %1 = vector.insert_slices %0, [2, 2], [1, 1]
         : tuple<vector<2x2xf32>, vector<2x2xf32>> into vector<4x2xf32>
 
       // Example with partial slices at dimension boundaries.
       %3 = vector.extract_slices %2, [2, 2], [1, 1]
let description = [{
 
     At the moment strides must contain only 1s.
 
     Returns an n-D vector that is a copy of the n-D destination vector in which
     the last k-D dimensions contain the k-D source vector elements strided at
     the proper location as specified by the offsets.
 
     Examples:
-    ```
+    ```mlir
       %2 = vector.insert_strided_slice %0, %1
           {offsets = [0, 0, 2], strides = [1, 1]}:
         vector<2x4xf32> into vector<16x4x8xf32>
     ```
   }];
 
   let assemblyFormat = [{
     $source `,` $dest attr-dict `:` type($source) `into` type($dest)
let description = [{
     An optional extra 2-D vector argument may be specified in which case the
     operation returns the sum of the outer-product and the extra vector. In this
     multiply-accumulate scenario, the rounding mode is that obtained by
     guaranteeing that a fused-multiply add operation is emitted. When lowered to
     the LLVMIR dialect, this form emits `llvm.intr.fma`, which is guaranteed to
     lower to actual `fma` instructions on x86.
 
     Examples:
-
-    ```
+    ```mlir
       %2 = vector.outerproduct %0, %1: vector<4xf32>, vector<8xf32>
       return %2: vector<4x8xf32>
 
       %3 = vector.outerproduct %0, %1, %2:
         vector<4xf32>, vector<8xf32>, vector<4x8xf32>
       return %3: vector<4x8xf32>
     ```
   }];
let description = [{
     across fixed vector shapes. For example, if a vector has a valid data
     shape [6] with fixed vector size [8], then the valid data elements are
     assumed to be stored at the beginning of the vector with the remaining
     vector elements undefined.
 
     In the examples below, valid data elements are represented by an alphabetic
     character, and undefined data elements are represented by '-'.
 
-    Example
-
+    Example:
+    ```mlir
       vector<1x8xf32> with valid data shape [6], fixed vector sizes [8]
 
                 input: [a, b, c, d, e, f]
 
            layout map: (d0) -> (d0 floordiv 8, d0 mod 8)
 
         vector layout: [a, b, c, d, e, f, -, -]
 
-    Example
-
+    ```
+    Example:
+    ```mlir
       vector<2x8xf32> with valid data shape [10], fixed vector sizes [8]
 
                 input: [a, b, c, d, e, f, g, h, i, j]
 
            layout map: (d0) -> (d0 floordiv 8, d0 mod 8)
 
         vector layout: [[a, b, c, d, e, f, g, h],
                         [i, j, -, -, -, -, -, -]]
-
-    Example
-
+    ```
+    Example:
+    ```mlir
       vector<2x2x2x3xf32> with valid data shape [3, 5], fixed vector sizes
       [2, 3]
 
                 input: [[a, b, c, d, e],
                         [f, g, h, i, j],
                         [k, l, m, n, o]]
 
            layout map: (d0, d1) -> (d0 floordiv 3, d1 floordiv 5,
                                     d0 mod 3, d1 mod 5)
 
         vector layout: [[[[a, b, c],
                           [f, g, h]]
                          [[d, e, -],
                           [i, j, -]]],
                         [[[k, l, m],
                           [-, -, -]]
                          [[n, o, -],
                           [-, -, -]]]]
-
-    Example
-
+    ```
+    Example:
+    ```mlir
       %1 = vector.reshape %0, [%c3, %c6], [%c2, %c9], [4]
         : vector<3x2x4xf32> to vector<2x3x4xf32>
 
              input: [[a, b, c, d, e, f],
                      [g, h, i, j, k, l],
                      [m, n, o, p, q, r]]
 
         layout map: (d0, d1) -> (d0, d1 floordiv 4, d1 mod 4)
 
 
       Input vector:  [[[a, b, c, d],
                        [e, f, -, -]],
                       [[g, h, i, j],
                        [k, l, -, -]],
                       [[m, n, o, p],
                        [q, r, -, -]]]
 
       Output vector:  [[[a, b, c, d],
                         [e, f, g, h],
                         [i, -, -, -]],
                        [[j, k, l, m],
                         [n, o, p, q],
                         [r, -, -, -]]]
+    ```
   }];
 
   let extraClassDeclaration = [{
     VectorType getInputVectorType() {
       return vector().getType().cast<VectorType>();
     }
     VectorType getOutputVectorType() {
       return getResult().getType().cast<VectorType>();
let description = [{
     At the moment strides must contain only 1s.
     // TODO(ntv) support non-1 strides.
 
     Returns an n-D vector where the first k-D dimensions match the `sizes`
     attribute. The returned subvector contains the elements starting at offset
     `offsets` and ending at `offsets + sizes`.
 
     Examples:
-    ```
+    ```mlir
       %1 = vector.strided_slice %0
           {offsets = [0, 2], sizes = [2, 4], strides = [1, 1]}:
         vector<4x8x16xf32> to vector<2x4x16xf32>
     ```
 
     // TODO(ntv) Evolve to a range form syntax similar to:
     %1 = vector.strided_slice %0[0:2:1][2:4:1]
       vector<4x8x16xf32> to vector<2x4x16xf32>
let description = [{
 
     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`.
 
     Syntax
-    ```
+    ```mlir
     operation ::= ssa-id `=` `vector.transfer_read` ssa-use-list
       `{` attribute-entry `} :` memref-type `,` vector-type
     ```
 
     Examples:
-
     ```mlir
     // Read the slice `%A[%i0, %i1:%i1+256, %i2:%i2+32]` into vector<32x256xf32>
     // and pad with %f0 to handle the boundary case:
     %f0 = constant 0.0f : f32
     for %i0 = 0 to %0 {
       affine.for %i1 = 0 to %1 step 256 {
         affine.for %i2 = 0 to %2 step 32 {
           %v = vector.transfer_read %A[%i0, %i1, %i2], (%f0)
let description = [{
     mid-level abstraction that supports super-vectorization with non-effecting
     padding for full-tile-only code. It is the responsibility of
     `vector.transfer_write`'s implementation to ensure the memory writes are
     valid. Different lowerings may be pertinent depending on the hardware
     support.
 
     Syntax:
 
-    ```
+    ```mlir
     operation ::= `vector.transfer_write` ssa-use-list `{` attribute-entry `} :
       ` vector-type ', ' memref-type '
     ```
 
     Examples:
 
     ```mlir
     // write vector<16x32x64xf32> into the slice
let description = [{
     Performs a conversion from a memref with scalar element to a memref with a
     *single* vector element, copying the shape of the memref to the vector. This
     is the minimal viable operation that is required to makeke
     super-vectorization operational. It can be seen as a special case of the
     `view` operation but scoped in the super-vectorization context.
 
     Syntax:
 
-    ```
+    ```mlir
     operation ::= `vector.type_cast` ssa-use : memref-type to memref-type
     ```
 
     Example:
 
     ```mlir
     %A  = alloc() : memref<5x4x3xf32>
     %VA = vector.type_cast %A : memref<5x4x3xf32> to memref<vector<5x4x3xf32>>
let description = [{
     are set to '0' or '1', based on whether the element indices are contained
     within a hyper-rectangular region specified by the 'mask_dim_sizes'
     array attribute argument. Each element of the 'mask_dim_sizes' array,
     specifies an exclusive upper bound [0, mask-dim-size-element-value)
     for a unique dimension in the vector result. The conjunction of the ranges
     define a hyper-rectangular region within which elements values are set to 1
     (otherwise element values are set to 0).
 
-    Example: create a constant vector mask of size 4x3xi1 with elements in range
-             0 <= row <= 2 and 0 <= col <= 1 are set to 1 (others to 0).
+    Example:
+    ```
+      create a constant vector mask of size 4x3xi1 with elements in range
 
       %1 = vector.constant_mask [3, 2] : vector<4x3xi1>
 
       print %1
                     columns
                   0    1    2
                 |------------
               0 | 1    1    0
         rows  1 | 1    1    0
               2 | 1    1    0
               3 | 0    0    0
+    ```
   }];
 
   let extraClassDeclaration = [{
     static StringRef getMaskDimSizesAttrName() { return "mask_dim_sizes"; }
   }];
   let assemblyFormat = "$mask_dim_sizes attr-dict `:` type(results)";
 }
 
 def Vector_CreateMaskOp :
   Vector_Op<"create_mask", [NoSideEffect]>,
     Arguments<(ins Variadic<Index>:$operands)>, Results<(outs VectorOf<[I1]>)> {
   let summary = "creates a vector mask";
   let description = [{
     Creates and returns a vector mask where elements of the result vector
     are set to '0' or '1', based on whether the element indices are contained
     within a hyper-rectangular region specified by the operands. Specifically,
     each operand specifies a range [0, operand-value) for a unique dimension in
     the vector result. The conjunction of the operand ranges define a
     hyper-rectangular region within which elements values are set to 1
     (otherwise element values are set to 0).
 
-    Example: create a vector mask of size 4x3xi1 where elements in range
-             0 <= row <= 2 and 0 <= col <= 1 are set to 1 (others to 0).
+    Example:
+    ```
+      create a vector mask of size 4x3xi1 where elements in range
 
       %1 = vector.create_mask %c3, %c2 : vector<4x3xi1>
 
       print %1
                     columns
                   0    1    2
                 |------------
               0 | 1    1    0
         rows  1 | 1    1    0
               2 | 1    1    0
               3 | 0    0    0
+    ```
   }];
 
   let hasCanonicalizer = 1;
   let assemblyFormat = "$operands attr-dict `:` type(results)";
 }
 
 def Vector_TupleOp :
   Vector_Op<"tuple", [NoSideEffect]>,
     Arguments<(ins Variadic<AnyVector>:$vectors)>,
     Results<(outs TupleOf<[AnyVector]>)> {
   let summary = "make tuple of vectors operation";
   let description = [{
     Returns a tuple of its operands 'vectors'.
 
     Note that this operation is used during the vector op unrolling
     transformation and should be removed before lowering to lower-level
     dialects.
 
-
     Examples:
-    ```
+    ```mlir
       %0 = vector.transfer_read ... : vector<2x2xf32>
       %1 = vector.transfer_read ... : vector<2x1xf32>
       %2 = vector.transfer_read ... : vector<2x2xf32>
       %3 = vector.transfer_read ... : vector<2x1xf32>
 
       %4 = vector.tuple %0, %1, %2, %3
         : vector<2x2xf32>, vector<2x1xf32>, vector<2x2xf32>, vector<2x1xf32>
 
PredOpTrait<"operand and result have same element type",
                  TCresVTEtIsSameAsOpBase<0, 0>>]>,
     Arguments<(ins AnyVector:$vector, I64ArrayAttr:$transp)>,
     Results<(outs AnyVector:$result)> {
   let summary = "vector transpose operation";
   let description = [{
     Takes a n-D vector and returns the transposed n-D vector defined by
     the permutation of ranks in the n-sized integer array attribute.
     In the operation
-
+    ```mlir
       %1 = vector.tranpose %0, [i_1, .., i_n]
         : vector<d_1 x .. x d_n x f32>
         to vector<d_trans[0] x .. x d_trans[n-1] x f32>
-
+    ```
     the transp array [i_1, .., i_n] must be a permutation of [0, .., n-1].
 
     Example:
-
+    ```mlir
     %1 = vector.tranpose %0, [1, 0] : vector<2x3xf32> to vector<3x2xf32>
 
      [ [a, b, c],       [ [a, d],
        [d, e, f] ]  ->    [b, e],
                           [c, f] ]
+    ```
   }];
   let extraClassDeclaration = [{
     VectorType getVectorType() {
       return vector().getType().cast<VectorType>();
     }
     VectorType getResultType() {
       return result().getType().cast<VectorType>();
     }
def Vector_TupleGetOp :
   let description = [{
     Returns the tuple element of 'vectors' at 'index'.
 
     Note that this operation is used during the vector op unrolling
     transformation and should be removed before lowering to lower-level
     dialects.
 
     Examples:
-    ```
+    ```mlir
       %4 = vector.tuple %0, %1, %2, %3
         : vector<2x2xf32>, vector<2x1xf32>, vector<2x2xf32>, vector<2x1xf32>>
 
       %5 = vector.tuple_get %4, 1
         : tuple<vector<2x2xf32>, vector<2x1xf32>,
                 vector<2x2xf32>, vector<2x1xf32>>
     ```
   }];

 def Vector_PrintOp :
   Vector_Op<"print", []>, Arguments<(ins AnyType:$source)> {
   let summary = "print operation (for testing and debugging)";
   let description = [{
     Prints the source vector (or scalar) to stdout in human readable
     format (for testing and debugging). No return value.
 
     Examples:
-    ```
+    ```mlir
       %0 = constant 0.0 : f32
       %1 = vector.broadcast %0 : f32 to vector<4xf32>
       vector.print %1 : vector<4xf32>
 
       when lowered to LLVM, the vector print is unrolled into
       elementary printing method calls that at runtime will yield
 
       ( 0.0, 0.0, 0.0, 0.0 )
let description = [{
 
     The ‘vector.matrix_multiply’ op treats `lhs` as matrix with <lhs_rows> rows
     and <lhs_columns> columns, `rhs` as matrix with <lhs_columns> rows and
     <rhs_columns> and multiplies them. The result matrix is returned embedded in
     the result vector.
 
     Example:
 
-    ```
+    ```mlir
       %C = vector.matrix_multiply %A, %B
         { lhs_rows = 4: i32, lhs_columns = 16: i32 , rhs_columns = 3: i32 } :
         (vector<64xf64>, vector<48xf64>) -> vector<12xf64>
     ```
   }];
   let builders = [
    OpBuilder<"Builder *builder, OperationState &result, Value lhs, Value rhs, "
              "unsigned lhsRows, unsigned lhsColumns, unsigned rhsColumns",

mlir/lib/Dialect/Vector/VectorTransforms.cpp

for (int64_t i = 0; i < tupleSize; ++i) {
           loc, tupleGet, result, elementOffsets, strides);
     }
 
     rewriter.replaceOp(op, result);
     return success();
   }
 };
 
 /// Progressive lowering of OuterProductOp.
 /// One:
+///   %x = vector.transpose %y, [1, 0]
+/// is replaced by:
+///   %z = constant dense<0.000000e+00>
+///   %0 = vector.extract %y[0, 0]
+///   %1 = vector.insert %0, %z [0, 0]
+///   ..
+///   %x = vector.insert .., .. [.., ..]
+class TransposeOpLowering : public OpRewritePattern<vector::TransposeOp> {
+public:
+  using OpRewritePattern<vector::TransposeOp>::OpRewritePattern;
+
+  LogicalResult matchAndRewrite(vector::TransposeOp op,
+                                PatternRewriter &rewriter) const override {
+    auto loc = op.getLoc();
+
+    VectorType resType = op.getResultType();
+    Type eltType = resType.getElementType();
+
+    // Set up convenience transposition table.
+    SmallVector<int64_t, 4> transp;
+    for (auto attr : op.transp())
+      transp.push_back(attr.cast<IntegerAttr>().getInt());
+
+    // Generate fully unrolled extract/insert ops.
+    Value zero = rewriter.create<ConstantOp>(loc, eltType,
+                                             rewriter.getZeroAttr(eltType));
+    Value result = rewriter.create<SplatOp>(loc, resType, zero);
+    SmallVector<int64_t, 4> lhs(transp.size(), 0);
+    SmallVector<int64_t, 4> rhs(transp.size(), 0);
+    rewriter.replaceOp(op, expandIndices(loc, resType, 0, transp, lhs, rhs,
+                                         op.vector(), result, rewriter));
+    return success();
+  }
+
+private:
+  // Builds the indices arrays for the lhs and rhs. Generates the extract/insert
+  // operation when al ranks are exhausted.
+  Value expandIndices(Location loc, VectorType resType, int64_t pos,
+                      SmallVector<int64_t, 4> &transp,
+                      SmallVector<int64_t, 4> &lhs,
+                      SmallVector<int64_t, 4> &rhs, Value input, Value result,
+                      PatternRewriter &rewriter) const {
+    if (pos >= resType.getRank()) {
+      auto ridx = rewriter.getI64ArrayAttr(rhs);
+      auto lidx = rewriter.getI64ArrayAttr(lhs);
+      Type eltType = resType.getElementType();
+      Value e = rewriter.create<vector::ExtractOp>(loc, eltType, input, ridx);
+      return rewriter.create<vector::InsertOp>(loc, resType, e, result, lidx);
+    }
+    for (int64_t d = 0, e = resType.getDimSize(pos); d < e; ++d) {
+      lhs[pos] = d;
+      rhs[transp[pos]] = d;
+      result = expandIndices(loc, resType, pos + 1, transp, lhs, rhs, input,
+                             result, rewriter);
+    }
+    return result;
+  }
+};
+
+/// Progressive lowering of OuterProductOp.
+/// One:
 ///   %x = vector.outerproduct %lhs, %rhs, %acc
 /// is replaced by:
 ///   %z = zero-result
 ///   %0 = vector.extract %lhs[0]
 ///   %1 = vector.broadcast %0
 ///   %2 = vector.extract %acc[0]
 ///   %3 = vector.fma %1, %arg1, %2
 ///   %4 = vector.insert %3, %z[0]
void mlir::vector::populateVectorSlicesLoweringPatterns(
     OwningRewritePatternList &patterns, MLIRContext *context) {
   patterns.insert<ExtractSlicesOpLowering, InsertSlicesOpLowering>(context);
 }
 
 void mlir::vector::populateVectorContractLoweringPatterns(
     OwningRewritePatternList &patterns, MLIRContext *context,
     VectorTransformsOptions parameters) {
   patterns.insert<ShapeCastOp2DDownCastRewritePattern,
-                  ShapeCastOp2DUpCastRewritePattern, OuterProductOpLowering>(
-      context);
+                  ShapeCastOp2DUpCastRewritePattern, TransposeOpLowering,
+                  OuterProductOpLowering>(context);
   patterns.insert<ContractionOpLowering>(parameters, context);
 }

mlir/test/Dialect/Vector/vector-contract-transforms.mlir

 
 func @outerproduct_acc(%arg0: vector<2xf32>,
                        %arg1: vector<3xf32>,
                        %arg2: vector<2x3xf32>) -> vector<2x3xf32> {
   %0 = vector.outerproduct %arg0, %arg1, %arg2 : vector<2xf32>, vector<3xf32>
   return %0: vector<2x3xf32>
 }
 
+// CHECK-LABEL: func @transpose23
+// CHECK-SAME: %[[A:.*]]: vector<2x3xf32>
+// CHECK:      %[[Z:.*]] = constant dense<0.000000e+00> : vector<3x2xf32>
+// CHECK:      %[[T0:.*]] = vector.extract %[[A]][0, 0] : vector<2x3xf32>
+// CHECK:      %[[T1:.*]] = vector.insert %[[T0]], %[[Z]] [0, 0] : f32 into vector<3x2xf32>
+// CHECK:      %[[T2:.*]] = vector.extract %[[A]][1, 0] : vector<2x3xf32>
+// CHECK:      %[[T3:.*]] = vector.insert %[[T2]], %[[T1]] [0, 1] : f32 into vector<3x2xf32>
+// CHECK:      %[[T4:.*]] = vector.extract %[[A]][0, 1] : vector<2x3xf32>
+// CHECK:      %[[T5:.*]] = vector.insert %[[T4]], %[[T3]] [1, 0] : f32 into vector<3x2xf32>
+// CHECK:      %[[T6:.*]] = vector.extract %[[A]][1, 1] : vector<2x3xf32>
+// CHECK:      %[[T7:.*]] = vector.insert %[[T6]], %[[T5]] [1, 1] : f32 into vector<3x2xf32>
+// CHECK:      %[[T8:.*]] = vector.extract %[[A]][0, 2] : vector<2x3xf32>
+// CHECK:      %[[T9:.*]] = vector.insert %[[T8]], %[[T7]] [2, 0] : f32 into vector<3x2xf32>
+// CHECK:      %[[T10:.*]] = vector.extract %[[A]][1, 2] : vector<2x3xf32>
+// CHECK:      %[[T11:.*]] = vector.insert %[[T10]], %[[T9]] [2, 1] : f32 into vector<3x2xf32>
+// CHECK:      return %[[T11]] : vector<3x2xf32>
+
+func @transpose23(%arg0: vector<2x3xf32>) -> vector<3x2xf32> {
+  %0 = vector.transpose %arg0, [1, 0] : vector<2x3xf32> to vector<3x2xf32>
+  return %0 : vector<3x2xf32>
+}
+
 // Shape up and downcasts for 2-D vectors, for supporting conversion to
 // llvm.matrix operations
 // CHECK-LABEL: func @shape_casts
 func @shape_casts(%a: vector<2x2xf32>) -> (vector<4xf32>, vector<2x2xf32>) {
   // CHECK: %[[cst:.*]] = constant dense<0.000000e+00> : vector<4xf32>
   // CHECK: %[[cst22:.*]] = constant dense<0.000000e+00> : vector<2x2xf32>
   // CHECK: %[[ex0:.*]] = vector.extract %{{.*}}[0] : vector<2x2xf32>
   //