MLIR

Multi-Level IR Compiler Framework

'vector' Dialect

MLIR supports multi-dimensional vector types and custom operations on those types. A generic, retargetable, higher-order vector type (n-D with n > 1) is a structured type, that carries semantic information useful for transformations. This document discusses retargetable abstractions that exist in MLIR today and operate on ssa-values of type vector along with pattern rewrites and lowerings that enable targeting specific instructions on concrete targets. These abstractions serve to separate concerns between operations on memref (a.k.a buffers) and operations on vector values. This is not a new proposal but rather a textual documentation of existing MLIR components along with a rationale.

Positioning in the Codegen Infrastructure

The following diagram, recently presented with the StructuredOps abstractions , captures the current codegen paths implemented in MLIR in the various existing lowering paths.

The following diagram seeks to isolate vector dialects from the complexity of the codegen paths and focus on the payload-carrying ops that operate on std and vector types. This diagram is not to be taken as set in stone and representative of what exists today but rather illustrates the layering of abstractions in MLIR.

vector Abstractions in MLIR

This  separates concerns related to (a) defining efficient operations on vector types from (b) program analyses + transformations on memref, loops and other types of structured ops (be they HLO, LHLO, Linalg or other ). Looking a bit forward in time, we can put a stake in the ground and venture that the higher level of vector-level primitives we build and target from codegen (or some user/language level), the simpler our task will be, the more complex patterns can be expressed and the better performance will be.

Components of a Generic Retargetable Vector-Level Dialect

The existing MLIR vector-level dialects are related to the following bottom-up abstractions:

  1. Representation in LLVMIR via data structures, instructions and intrinsics. This is referred to as the LLVM level.
  2. Set of machine-specific operations and types that are built to translate almost 1-1 with the HW ISA. This is referred to as the Hardware Vector level; a.k.a HWV. For instance, we have (a) a NVVM dialect (for CUDA) with tensor core ops, (b) accelerator-specific dialects (internal), a potential (future) CPU dialect to capture LLVM intrinsics more closely and other dialects for specific hardware. Ideally this should be auto-generated as much as possible from the LLVM level.
  3. Set of virtual, machine-agnostic, operations that are informed by costs at the HWV-level. This is referred to as the Virtual Vector level; a.k.a VV. This is the level that higher-level abstractions (codegen, automatic vectorization, potential vector language, …) targets.

The existing generic, retargetable, vector-level dialect is related to the following top-down rewrites and conversions:

  1. MLIR Rewrite Patterns applied by the MLIR PatternRewrite infrastructure to progressively lower to implementations that match closer and closer to the HWV. Some patterns are “in-dialect” VV -> VV and some are conversions VV -> HWV.
  2. Virtual Vector -> Hardware Vector lowering is specified as a set of MLIR lowering patterns that are specified manually for now.
  3. Hardware Vector -> LLVM lowering is a mechanical process that is written manually at the moment and that should be automated, following the LLVM -> Hardware Vector ops generation as closely as possible.

Short Description of the Existing Infrastructure

LLVM level

On CPU, the n-D vector type currently lowers to !llvm<array<vector>>. More concretely, vector<4x8x128xf32> lowers to !llvm<[4 x [ 8 x [ 128 x float ]]]>. There are tradeoffs involved related to how one can access subvectors and how one uses llvm.extractelement, llvm.insertelement and llvm.shufflevector. A deeper dive section discusses the current lowering choices and tradeoffs.

Hardware Vector Ops

Hardware Vector Ops are implemented as one dialect per target. For internal hardware, we are auto-generating the specific HW dialects. For GPU, the NVVM dialect adds operations such as mma.sync, shfl and tests. For CPU things are somewhat in-flight because the abstraction is close to LLVMIR. The jury is still out on  whether a generic CPU dialect is concretely needed, but it seems reasonable to have the same levels of abstraction for all targets and perform cost-based lowering decisions in MLIR even for LLVM. Specialized CPU dialects that would capture specific features not well captured by LLVM peephole optimizations of on different types that core MLIR supports (e.g. Scalable Vectors) are welcome future extensions.

Virtual Vector Ops

Some existing Standard and Vector Dialect on n-D vector types comprise:

%2 = std.addf %0, %1 : vector<3x7x8xf32>  // -> vector<3x7x8xf32>
%2 = std.mulf %0, %1 : vector<3x7x8xf32>  // -> vector<3x7x8xf32>
%2 = std.splat %1    : vector<3x7x8xf32>  // -> vector<3x7x8xf32>

%1 = vector.extract %0[1]: vector<3x7x8xf32>                 // -> vector<7x8xf32>
%1 = vector.extract %0[1, 5]: vector<3x7x8xf32>            // -> vector<8xf32>
%2 = vector.outerproduct %0, %1: vector<4xf32>, vector<8xf32>     // -> vector<4x8xf32>
%3 = vector.outerproduct %0, %1, %2: vector<4xf32>, vector<8xf32> // fma when adding %2
%3 = vector.strided_slice %0 {offsets = [2, 2], sizes = [2, 2], strides = [1, 1]}:
   vector<4x8x16xf32> // Returns a slice of type vector<2x2x16xf32>

%2 = vector.transfer_read %A[%0, %1]
  {permutation_map = (d0, d1) -> (d0)}: memref<7x?xf32>, vector<4xf32>

vector.transfer_write %f1, %A[%i0, %i1, %i2, %i3]
  {permutation_map = (d0, d1, d2, d3) -> (d3, d1, d0)} :
    vector<5x4x3xf32>, memref<?x?x?x?xf32>

The list of Vector is currently undergoing evolutions and is best kept track of by following the evolution of the VectorOps.td ODS file (markdown documentation is automatically generated locally when building and populates the Vector doc ). Recent extensions are driven by concrete use cases of interest. A notable such use case is the vector.contract op which applies principles of the StructuredOps abstraction to vector types.

Virtual Vector Rewrite Patterns

The following rewrite patterns exist at the VV->VV level:

  1. The now retired MaterializeVector pass used to legalize ops on a coarse-grained virtual vector to a finer-grained virtual vector by unrolling. This has been rewritten as a retargetable unroll-and-jam pattern on vector ops and vector types.
  2. The lowering of vector_transfer ops legalizes vector load/store ops to permuted loops over scalar load/stores. This should evolve to loops over vector load/stores + mask operations as they become available vector ops at the VV level.

The general direction is to add more Virtual Vector level ops and implement more useful VV -> VV rewrites as composable patterns that the PatternRewrite infrastructure can apply iteratively.

Virtual Vector to Hardware Vector Lowering

For now, VV -> HWV are specified in C++ (see for instance the SplatOpLowering for n-D vectors or the VectorOuterProductOp lowering ).

Simple conversion tests are available for the LLVM target starting from the Virtual Vector Level.

Rationale

Hardware as vector Machines of Minimum Granularity

Higher-dimensional vectors are ubiquitous in modern HPC hardware. One way to think about Generic Retargetable vector-Level Dialect is that it operates on vector types that are a multiples of a “good” vector size so the HW can efficiently implement a set of high-level primitives (e.g. vector<8x8x8x16xf32> when HW vector size is say vector<4x8xf32>).

Some notable vector sizes of interest include:

  1. CPU: vector<HW_vector_size * k>, vector<core_count * k’ x HW_vector_size * k> and vector<socket_count x core_count * k’ x HW_vector_size * k>
  2. GPU: vector<warp_size * k>, vector<warp_size * k x float4> and vector<warp_size * k x 4 x 4 x 4> for tensor_core sizes,
  3. Other accelerators: n-D vector as first-class citizens in the HW.

Depending on the target, ops on sizes that are not multiples of the HW vector size may either produce slow code (e.g. by going through LLVM legalization) or may not legalize at all (e.g. some unsupported accelerator X combination of ops and types).

Transformations Problems Avoided

A vector<16x32x64xf32> virtual vector is a coarse-grained type that can be “unrolled” to HW-specific sizes. The multi-dimensional unrolling factors are carried in the IR by the vector type. After unrolling, traditional instruction-level scheduling can be run.

The following key transformations (along with the supporting analyses and structural constraints) are completely avoided by operating on a vector ssa-value abstraction:

  1. Loop unroll and unroll-and-jam.
  2. Loop and load-store restructuring for register reuse.
  3. Load to store forwarding and Mem2reg.
  4. Coarsening (raising) from finer-grained vector form.

Note that “unrolling” in the context of vectors corresponds to partial loop unroll-and-jam and not full unrolling. As a consequence this is expected to compose with SW pipelining where applicable and does not result in ICache blow up.

The Big Out-Of-Scope Piece: Automatic Vectorization

One important piece not discussed here is automatic vectorization (automatically raising from scalar to n-D vector ops and types). The TL;DR is that when the first “super-vectorization” prototype was implemented, MLIR was nowhere near as mature as it is today. As we continue building more abstractions in VV -> HWV, there is an opportunity to revisit vectorization in MLIR.

Since this topic touches on codegen abstractions, it is technically out of the scope of this survey document but there is a lot to discuss in light of structured op type representations and how a vectorization transformation can be reused across dialects. In particular, MLIR allows the definition of dialects at arbitrary levels of granularity and lends itself favorably to progressive lowering. The argument can be made that automatic vectorization on a loops + ops abstraction is akin to raising structural information that has been lost. Instead, it is possible to revisit vectorization as simple pattern rewrites, provided the IR is in a suitable form. For instance, vectorizing a linalg.generic op whose semantics match a matmul can be done quite easily with a pattern . In fact this pattern is trivial to generalize to any type of contraction when targeting the vector.contract op, as well as to any field (+/*, min/+, max/+, or/and, logsumexp/+ …) . In other words, by operating on a higher level of generic abstractions than affine loops, non-trivial transformations become significantly simpler and composable at a finer granularity.

Irrespective of the existence of an auto-vectorizer, one can build a notional vector language based on the VectorOps dialect and build end-to-end models with expressing vectors in the IR directly and simple pattern-rewrites. EDSC s provide a simple way of driving such a notional language directly in C++.

Bikeshed Naming Discussion

There are arguments against naming an n-D level of abstraction vector because most people associate it with 1-D vectors. On the other hand, vectors are first-class n-D values in MLIR. The alternative name Tile has been proposed, which conveys higher-D meaning. But it also is one of the most overloaded terms in compilers and hardware. For now, we generally use the n-D vector name and are open to better suggestions.

DeeperDive

This section describes the tradeoffs involved in lowering the MLIR n-D vector type and operations on it to LLVM-IR. Putting aside the LLVM Matrix proposal for now, this assumes LLVM only has built-in support for 1-D vector. The relationship with the LLVM Matrix proposal is discussed at the end of this document.

MLIR does not currently support dynamic vector sizes (i.e. SVE style) so the discussion is limited to static rank and static vector sizes (e.g. vector<4x8x16x32xf32>). This section discusses operations on vectors in LLVM and MLIR.

LLVM instructions are prefixed by the llvm. dialect prefix (e.g. llvm.insertvalue). Such ops operate exclusively on 1-D vectors and aggregates following the LLVM LangRef . MLIR operations are prefixed by the vector. dialect prefix (e.g. vector.insertelement). Such ops operate exclusively on MLIR n-D vector types.

Alternatives For Lowering an n-D Vector Type to LLVM

Consider a vector of rank n with static sizes {s_0, ... s_{n-1}} (i.e. an MLIR vector<s_0x...s_{n-1}xf32>). Lowering such an n-D MLIR vector type to an LLVM descriptor can be done by either:

  1. Flattening to a 1-D vector: !llvm<"(s_0*...*s_{n-1})xfloat"> in the MLIR LLVM dialect.
  2. Nested aggregate type of 1-D vector: !llvm<"[s_0x[s_1x[...<s_{n-1}xfloat>]]]"> in the MLIR LLVM dialect.
  3. A mix of both.

There are multiple tradeoffs involved in choosing one or the other that we discuss. It is important to note that “a mix of both” immediately reduces to “nested aggregate type of 1-D vector” with a vector.cast %0: vector<4x8x16x32xf32> to vector<4x4096xf32> operation, that flattens the most “k” minor dimensions.

Constraints Inherited from LLVM (see LangRef)

The first constraint was already mentioned: LLVM only supports 1-D vector types natively. Additional constraints are related to the difference in LLVM between vector and aggregate types:

 “Aggregate Types are a subset of derived types that can contain multiple
 member types. Arrays and structs are aggregate types. Vectors are not
 considered to be aggregate types.”.

This distinction is also reflected in some of the operations. For 1-D vectors, the operations llvm.extractelement, llvm.insertelement, and llvm.shufflevector apply, with direct support for dynamic indices. For n-D vectors with n>1, and thus aggregate types at LLVM level, the more restrictive operations llvm.extractvalue and llvm.insertvalue apply, which only accept static indices. There is no direct shuffling support for aggregate types.

The next sentence illustrates a recurrent tradeoff, also found in MLIR, between “value types” (subject to SSA use-def chains) and “memory types” (subject to aliasing and side-effects):

“Structures in memory are accessed using ‘load’ and ‘store’ by getting a
pointer to a field with the llvm.getelementptr instruction. Structures in
registers are accessed using the llvm.extractvalue and llvm.insertvalue
instructions.”

When transposing this to MLIR, llvm.getelementptr works on pointers to n-D vectors in memory. For n-D, vectors values that live in registers we can use vector.extract and vector.insert which do not accept dynamic indices. Note that this is consistent with hardware considerations as discussed below.

An alternative is to use an LLVM 1-D vector type for which one can use llvm.extractelement, llvm.insertelement and llvm.shufflevector. These operations accept dynamic indices. The implication is that one has to use a flattened lowering of an MLIR n-D vector to an LLVM 1-D vector.

There are multiple tradeoffs involved that mix implications on the programming model, execution on actual HW and what is visible or hidden from codegen. They are discussed in the following sections.

Nested Aggregate

Pros:

  1. Natural encoding n-D vector -> (n-1)-D aggregate over 1-D vector.
  2. No need for linearization / delinearization logic inserted everywhere.
  3. llvm.insertvalue, llvm.extractvalue of (n-k)-D aggregate is natural.
  4. llvm.insertelement, llvm.extractelement, llvm.shufflevector over 1-D vector type is natural.

Cons:

  1. llvm.insertvalue / llvm.extractvalue does not accept dynamic indices but only static ones.
  2. Dynamic indexing on the non-most-minor dimension requires roundtrips to memory.
  3. Special intrinsics and native instructions in LLVM operate on 1-D vectors. This is not expected to be a practical limitation thanks to a vector.cast %0: vector<4x8x16x32xf32> to vector<4x4096xf32> operation, that flattens the most minor dimensions (see the bigger picture in implications on codegen).

Flattened 1-D Vector Type

Pros:

  1. insertelement / extractelement / shufflevector with dynamic indexing is possible over the whole lowered n-D vector type.
  2. Supports special intrinsics and native operations.

Cons:

  1. Requires linearization/delinearization logic everywhere, translations are complex.
  2. Hides away the real HW structure behind dynamic indexing: at the end of the day, HW vector sizes are generally fixed and multiple vectors will be needed to hold a vector that is larger than the HW.
  3. Unlikely peephole optimizations will result in good code: arbitrary dynamic accesses, especially at HW vector boundaries unlikely to result in regular patterns.

Discussion

HW Vectors and Implications on the SW and the Programming Model

As of today, the LLVM model only support 1-D vector types. This is unsurprising because historically, the vast majority of HW only supports 1-D vector registers. We note that multiple HW vendors are in the process of evolving to higher-dimensional physical vectors.

In the following discussion, let’s assume the HW vector size is 1-D and the SW vector size is n-D, with n >= 1. The same discussion would apply with2-DHWvectorsize andn >= 2. In this context, most HW exhibit a vector register file. The number of such vectors is fixed. Depending on the rank and sizes of the SW vector abstraction and the HW vector sizes and number of registers, an n-DSW vector type may be materialized by a mix of multiple1-D` HW vector registers + memory locations at a given point in time.

The implication of the physical HW constraints on the programming model are that one cannot index dynamically across hardware registers: a register file can generally not be indexed dynamically. This is because the register number is fixed and one either needs to unroll explicitly to obtain fixed register numbers or go through memory. This is a constraint familiar to CUDA programmers: when declaring a private float a[4]; and subsequently indexing with a dynamic value results in so-called local memory usage (i.e. roundtripping to memory).

Implication on codegen

MLIR n-D vector types are currently represented as (n-1)-D arrays of 1-D vectors when lowered to LLVM. This introduces the consequences on static vs dynamic indexing discussed previously: extractelement, insertelement and shufflevector on n-D vectors in MLIR only support static indices. Dynamic indices are only supported on the most minor 1-D vector but not the outer (n-1)-D. For other cases, explicit load / stores are required.

The implications on codegen are as follows:

  1. Loops around vector values are indirect addressing of vector values, they must operate on explicit load / store operations over n-D vector types.
  2. Once an n-D vector type is loaded into an SSA value (that may or may not live in n registers, with or without spilling, when eventually lowered), it may be unrolled to smaller k-D vector types and operations that correspond to the HW. This level of MLIR codegen is related to register allocation and spilling that occur much later in the LLVM pipeline.
  3. HW may support >1-D vectors with intrinsics for indirect addressing within these vectors. These can be targeted thanks to explicit vector_cast operations from MLIR k-D vector types and operations to LLVM 1-D vectors + intrinsics.

Alternatively, we argue that directly lowering to a linearized abstraction hides away the codegen complexities related to memory accesses by giving a false impression of magical dynamic indexing across registers. Instead we prefer to make those very explicit in MLIR and allow codegen to explore tradeoffs. Different HW will require different tradeoffs in the sizes involved in steps 1., 2. and 3.

Decisions made at the MLIR level will have implications at a much later stage in LLVM (after register allocation). We do not envision to expose concerns related to modeling of register allocation and spilling to MLIR explicitly. Instead, each target will expose a set of “good” target operations and n-D vector types, associated with costs that PatterRewriters at the MLIR level will be able to target. Such costs at the MLIR level will be abstract and used for ranking, not for accurate performance modeling. In the future such costs will be learned.

Implication on Lowering to Accelerators

To target accelerators that support higher dimensional vectors natively, we can start from either 1-D or n-D vectors in MLIR and use vector.cast to flatten the most minor dimensions to 1-D vector<Kxf32> where K is an appropriate constant. Then, the existing lowering to LLVM-IR immediately applies, with extensions for accelerator-specific intrinsics.

It is the role of an Accelerator-specific vector dialect (see codegen flow in the figure above) to lower the vector.cast. Accelerator -> LLVM lowering would then consist of a bunch of Accelerator -> Accelerator rewrites to perform the casts composed with Accelerator -> LLVM conversions + intrinsics that operate on 1-D vector<Kxf32>.

Some of those rewrites may need extra handling, especially if a reduction is involved. For example, vector.cast %0: vector<K1x...xKnxf32> to vector<Kxf32> when K != K1 * … * Kn and some arbitrary irregular vector.cast %0: vector<4x4x17xf32> to vector<Kxf32> may introduce masking and intra-vector shuffling that may not be worthwhile or even feasible, i.e. infinite cost.

However vector.cast %0: vector<K1x...xKnxf32> to vector<Kxf32> when K = K1 * … * Kn should be close to a noop.

As we start building accelerator-specific abstractions, we hope to achieve retargetable codegen: the same infra is used for CPU, GPU and accelerators with extra MLIR patterns and costs.

Implication on calling external functions that operate on vectors

It is possible (likely) that we additionally need to linearize when calling an external function.

Relationship to LLVM matrix type proposal.

The LLVM matrix proposal was formulated 1 year ago but seemed to be somewhat stalled until recently. In its current form, it is limited to 2-D matrix types and operations are implemented with LLVM intrinsics. In contrast, MLIR sits at a higher level of abstraction and allows the lowering of generic operations on generic n-D vector types from MLIR to aggregates of 1-D LLVM vectors. In the future, it could make sense to lower to the LLVM matrix abstraction also for CPU even though MLIR will continue needing higher level abstractions.

On the other hand, one should note that as MLIR is moving to LLVM, this document could become the unifying abstraction that people should target for

1-D vectors and the LLVM matrix proposal can be viewed as a subset of this work.

Conclusion

The flattened 1-D vector design in the LLVM matrix proposal is good in a HW-specific world with special intrinsics. This is a good abstraction for register allocation, Instruction-Level-Parallelism and SoftWare-Pipelining/Modulo Scheduling optimizations at the register level. However MLIR codegen operates at a higher level of abstraction where we want to target operations on coarser-grained vectors than the HW size and on which unroll-and-jam is applied and patterns across multiple HW vectors can be matched.

This makes “nested aggregate type of 1-D vector” an appealing abstraction for lowering from MLIR because:

  1. it does not hide complexity related to the buffer vs value semantics and the memory subsystem and
  2. it does not rely on LLVM to magically make all the things work from a too low-level abstraction.

The use of special intrinsics in a 1-D LLVM world is still available thanks to an explicit vector.cast op.

Operations

vector.broadcast (vector::BroadcastOp)

broadcast operation

Syntax:

operation ::= `vector.broadcast` $source attr-dict `:` type($source) `to` type($vector)

Broadcasts the scalar or k-D vector value in the source operand to a n-D result vector such that the broadcast makes sense, i.e., the source operand is duplicated to match the given rank and sizes in the result vector. The legality rules are:

  • the source operand must have the same element type as the result type
  • a k-D vector <s_1 x .. x s_k x type> can be broadcast to a n-D vector <t_1 x .. x t_n x type> if
    • k <= n, and
    • the sizes in the trailing dimensions n-k < i <= n with j=i+k-n match exactly as s_j = t_i or s_j = 1:
        t_1 x   ..  t_n-k x t_n-k+1 x .. x t_i x .. x t_n
                            s_1     x .. x s_j x .. x s_k
            <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.

Example:

%0 = constant 0.0 : f32
%1 = vector.broadcast %0 : f32 to vector<16xf32>
%2 = vector.broadcast %1 : vector<16xf32> to vector<4x16xf32>

Operands:

OperandDescription
sourceany type

Results:

ResultDescription
vectorvector of any type values

vector.constant_mask (vector::ConstantMaskOp)

creates a constant vector mask

Syntax:

operation ::= `vector.constant_mask` $mask_dim_sizes attr-dict `:` type(results)

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 ‘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).
%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

Attributes:

AttributeMLIR TypeDescription
mask_dim_sizesArrayAttr64-bit integer array attribute

Results:

ResultDescription
«unnamed»vector of 1-bit signless integer values

vector.contract (vector::ContractionOp)

vector contraction operation

Computes the sum of products of vector elements along contracting dimension pairs from 2 vectors of rank M and N respectively, adds this intermediate result to the accumulator argument of rank K, and returns a vector result of rank K (where K = num_lhs_free_dims + num_rhs_free_dims + num_batch_dims (see dimension type descriptions below)). For K = 0 (no free or batch dimensions), the accumulator and output are a scalar.

Optional vector mask arguments (produced by CreateMaskOp or ConstantMaskOp) specify the dynamic dimension sizes of valid data within the lhs/rhs vector arguments.

An iterator type attribute list must be specified, where each element of the list represents an iterator with one of the following types:

*) “reduction”: reduction dimensions are present in the lhs and rhs arguments but not in the output (and accumulator argument). These are the dimensions along which the vector contraction op computes the sum of products, and contracting dimension pair dimension sizes must match between lhs/rhs. *) “parallel”: Batch dimensions are iterator type “parallel”, and are non-contracting dimensions present in the lhs, rhs and output. The lhs/rhs co-iterate along the batch dimensions, which should be expressed in their indexing maps.

           Free dimensions are iterator type "parallel", and are
           non-contraction, non-batch dimensions accessed by either the
           lhs or rhs (but not both). The lhs and rhs free dimensions
           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.

Example:

// Simple dot product (K = 0).
#contraction_accesses = [
 affine_map<(i) -> (i)>,
 affine_map<(i) -> (i)>,
 affine_map<(i) -> ()>
]
#contraction_trait = {
  indexing_maps = #contraction_accesses,
  iterator_types = ["reduction"]
}
%3 = vector.contract #contraction_trait %0, %1, %2
  : vector<10xf32>, vector<10xf32> into f32

// 2D vector contraction with one contracting dimension (matmul, K = 2).
#contraction_accesses = [
  affine_map<(i, j, k) -> (i, k)>,
  affine_map<(i, j, k) -> (k, j)>,
  affine_map<(i, j, k) -> (i, j)>
]
#contraction_trait = {
  indexing_maps = #contraction_accesses,
  iterator_types = ["parallel", "parallel", "reduction"]
}

%3 = vector.contract #contraction_trait %0, %1, %2
  : vector<4x3xf32>, vector<3x7xf32> into vector<4x7xf32>

// 4D to 3D vector contraction with two contracting dimensions and
// one batch dimension (K = 3).
#contraction_accesses = [
  affine_map<(b0, f0, f1, c0, c1) -> (c0, b0, c1, f0)>,
  affine_map<(b0, f0, f1, c0, c1) -> (b0, c1, c0, f1)>,
  affine_map<(b0, f0, f1, c0, c1) -> (b0, f0, f1)>
]
#contraction_trait = {
  indexing_maps = #contraction_accesses,
  iterator_types = ["parallel", "parallel", "parallel",
                    "reduction", "reduction"]
}

%4 = vector.contract #contraction_trait %0, %1, %2
    : vector<7x8x16x15xf32>, vector<8x16x7x5xf32> into vector<8x15x5xf32>

// 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>

Attributes:

AttributeMLIR TypeDescription
indexing_mapsArrayAttrAffineMap array attribute
iterator_typesArrayAttrarray attribute

Operands:

OperandDescription
lhsvector of any type values
rhsvector of any type values
accany type
masksvector of 1-bit signless integer values

Results:

ResultDescription
«unnamed»any type

vector.create_mask (vector::CreateMaskOp)

creates a vector mask

Syntax:

operation ::= `vector.create_mask` $operands attr-dict `:` type(results)

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).
%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

Operands:

OperandDescription
operandsindex

Results:

ResultDescription
«unnamed»vector of 1-bit signless integer values

vector.extractelement (vector::ExtractElementOp)

extractelement operation

Syntax:

operation ::= `vector.extractelement` $vector `[` $position `:` type($position) `]` attr-dict `:` type($vector)

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:

%c = constant 15 : i32
%1 = vector.extractelement %0[%c : i32]: vector<16xf32>

Operands:

OperandDescription
vectorvector of any type values
positionsignless integer

Results:

ResultDescription
resultany type

vector.extract (vector::ExtractOp)

extract operation

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.

Example:

%1 = vector.extract %0[3]: vector<4x8x16xf32>
%2 = vector.extract %0[3, 3, 3]: vector<4x8x16xf32>

Attributes:

AttributeMLIR TypeDescription
positionArrayAttr64-bit integer array attribute

Operands:

OperandDescription
vectorvector of any type values

Results:

ResultDescription
«unnamed»any type

vector.extract_slices (vector::ExtractSlicesOp)

vector extract slices operation

Syntax:

operation ::= `vector.extract_slices` $vector `,` $sizes `,` $strides attr-dict `:` type($vector) `into`
              type(results)

Takes an N-d vector and returns a tuple of vector slices of ‘vector’, based on ‘sizes’ and ‘strides’ parameters.

The arguments ‘sizes’ and ‘strides’ represent a specification for generating the unrolling of ‘vector’ shape, which has all slices of shape ‘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.

Example:

%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>

%3 = vector.extract_slices %2, [2, 2], [1, 1]
  : vector<4x3xf32> into tuple<vector<2x2xf32>, vector<2x1xf32>,
                               vector<2x2xf32>, vector<2x1xf32>>

Attributes:

AttributeMLIR TypeDescription
sizesArrayAttr64-bit integer array attribute
stridesArrayAttr64-bit integer array attribute

Operands:

OperandDescription
vectorvector of any type values

Results:

ResultDescription
«unnamed»tuple with any combination of vector of any type values values

vector.fma (vector::FMAOp)

vector fused multiply-add

Syntax:

operation ::= `vector.fma` $lhs `,` $rhs `,` $acc attr-dict `:` type($lhs)

Multiply-add expressions operate on n-D vectors and compute a fused 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:

%3 = vector.fma %0, %1, %2: vector<8x16xf32>

Operands:

OperandDescription
lhsvector of any type values
rhsvector of any type values
accvector of any type values

Results:

ResultDescription
resultvector of any type values

vector.insertelement (vector::InsertElementOp)

insertelement operation

Syntax:

operation ::= `vector.insertelement` $source `,` $dest `[` $position `:` type($position) `]` attr-dict `:`
              type($result)

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:

%c = constant 15 : i32
%f = constant 0.0f : f32
%1 = vector.insertelement %f, %0[%c : i32]: vector<16xf32>

Operands:

OperandDescription
sourceany type
destvector of any type values
positionsignless integer

Results:

ResultDescription
resultvector of any type values

vector.insert (vector::InsertOp)

insert operation

Syntax:

operation ::= `vector.insert` $source `,` $dest $position attr-dict `:` type($source) `into` type($dest)

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.

Example:

%2 = vector.insert %0, %1[3] : vector<8x16xf32> into vector<4x8x16xf32>
%5 = vector.insert %3, %4[3, 3, 3] : f32 into vector<4x8x16xf32>

Attributes:

AttributeMLIR TypeDescription
positionArrayAttr64-bit integer array attribute

Operands:

OperandDescription
sourceany type
destvector of any type values

Results:

ResultDescription
resvector of any type values

vector.insert_slices (vector::InsertSlicesOp)

vector insert slices operation

Syntax:

operation ::= `vector.insert_slices` $vectors `,` $sizes `,` $strides attr-dict `:` type($vectors) `into`
              type(results)

Takes a tuple of vector slices and inserts them into the vector result according to the ‘sizes’ and ‘strides’ parameters.

The arguments ‘sizes’ and ‘strides’ represent a specification for generating the unrolling of ‘vector’ shape, which has all slices of shape ‘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.

Example:

%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]
  : vector<4x3xf32> into tuple<vector<2x2xf32>, vector<2x1xf32>,
                               vector<2x2xf32>, vector<2x1xf32>>

%4 = vector.insert_slices %3, [2, 2], [1, 1]
  : tuple<vector<2x2xf32>, vector<2x1xf32>,
          vector<2x2xf32>, vector<2x1xf32>> into vector<4x3xf32>

Attributes:

AttributeMLIR TypeDescription
sizesArrayAttr64-bit integer array attribute
stridesArrayAttr64-bit integer array attribute

Operands:

OperandDescription
vectorstuple with any combination of vector of any type values values

Results:

ResultDescription
«unnamed»vector of any type values

vector.insert_strided_slice (vector::InsertStridedSliceOp)

strided_slice operation

Syntax:

operation ::= `vector.insert_strided_slice` $source `,` $dest attr-dict `:` type($source) `into` type($dest)

Takes a k-D source vector, an n-D destination vector (n >= k), n-sized offsets integer array attribute, a k-sized strides integer array attribute and inserts the k-D source vector as a strided subvector at the proper offset into the n-D destination vector.

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.

Example:

%2 = vector.insert_strided_slice %0, %1
    {offsets = [0, 0, 2], strides = [1, 1]}:
  vector<2x4xf32> into vector<16x4x8xf32>

Attributes:

AttributeMLIR TypeDescription
offsetsArrayAttr64-bit integer array attribute
stridesArrayAttr64-bit integer array attribute

Operands:

OperandDescription
sourcevector of any type values
destvector of any type values

Results:

ResultDescription
resvector of any type values

vector.matrix_multiply (vector::MatmulOp)

Vector matrix multiplication op that operates on flattened 1-D MLIR vectors

Syntax:

operation ::= `vector.matrix_multiply` $lhs `,` $rhs attr-dict `:` `(` type($lhs) `,` type($rhs) `)` `->` type($res)

This is the counterpart of llvm.matrix.multiply in MLIR. It serves the purposes of more progressive lowering and localized type conversion.

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:

%C = vector.matrix_multiply %A, %B
  { lhs_rows = 4: i32, lhs_columns = 16: i32 , rhs_columns = 3: i32 } :
  (vector<64xf64>, vector<48xf64>) -> vector<12xf64>

Attributes:

AttributeMLIR TypeDescription
lhs_rowsIntegerAttr32-bit signless integer attribute
lhs_columnsIntegerAttr32-bit signless integer attribute
rhs_columnsIntegerAttr32-bit signless integer attribute

Operands:

OperandDescription
lhsvector of signless integer or signed integer or floating-point values of ranks 1
rhsvector of signless integer or signed integer or floating-point values of ranks 1

Results:

ResultDescription
resvector of signless integer or signed integer or floating-point values of ranks 1

vector.outerproduct (vector::OuterProductOp)

vector outerproduct with optional fused add

Takes 2 1-D vectors and returns the 2-D vector containing the outer-product.

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.

Example:

%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>

Operands:

OperandDescription
lhsvector of any type values
rhsvector of any type values
accvector of any type values

Results:

ResultDescription
«unnamed»vector of any type values

vector.print (vector::PrintOp)

print operation (for testing and debugging)

Syntax:

operation ::= `vector.print` $source attr-dict `:` type($source)

Prints the source vector (or scalar) to stdout in human readable format (for testing and debugging). No return value.

Example:

%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 )

on stdout when linked with a small runtime support library,
which only needs to provide a few printing methods (single
value for all data types, opening/closing bracket, comma,
newline).

Operands:

OperandDescription
sourceany type

vector.reduction (vector::ReductionOp)

reduction operation

Reduces an 1-D vector “horizontally” into a scalar using the given operation (add/mul/min/max for int/fp and and/or/xor for int only). Some reductions (add/mul for fp) also allow an optional fused 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

Example:

%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

Attributes:

AttributeMLIR TypeDescription
kindStringAttrstring attribute

Operands:

OperandDescription
vectorvector of any type values
accany type

Results:

ResultDescription
destany type

vector.reshape (vector::ReshapeOp)

vector reshape operation

Syntax:

operation ::= `vector.reshape` $vector `,` `[` $input_shape `]` `,` `[` $output_shape `]` `,`
              $fixed_vector_sizes attr-dict `:` type($vector) `to` type($result)

Reshapes its vector operand from ‘input_shape’ to ‘output_shape’ maintaining fixed vector dimension ‘fixed_vector_sizes’ on the innermost vector dimensions.

The parameters ‘input_shape’ and ‘output_shape’ represent valid data shapes 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

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

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

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

%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, -, -, -]]]

Attributes:

AttributeMLIR TypeDescription
fixed_vector_sizesArrayAttr64-bit integer array attribute

Operands:

OperandDescription
vectorvector of any type values
input_shapeindex
output_shapeindex

Results:

ResultDescription
resultvector of any type values

vector.shape_cast (vector::ShapeCastOp)

shape_cast casts between vector shapes

Syntax:

operation ::= `vector.shape_cast` $source attr-dict `:` type($source) `to` type($result)

The shape_cast operation casts between an n-D source vector shape and a k-D result vector shape (the element type remains the same).

If reducing rank (n > k), result dimension sizes must be a product of contiguous source dimension sizes. If expanding rank (n < k), source dimensions must factor into a contiguous sequence of destination dimension sizes. Each source dim is expanded (or contiguous sequence of source dims combined) in source dimension list order (i.e. 0 <= i < n), to produce a contiguous sequence of result dims (or a single result dim), in result dimension list order (i.e. 0 <= j < k). The product of all source dimension sizes and all result dimension sizes must match.

If the source/result types are a tuple of vectors, the casting operation described above is applied to each source/result tuple element pair.

It is currently assumed that this operation does not require moving data, and that it will be folded away before lowering vector operations.

There is an exception to the folding expectation when targeting llvm.intr.matrix operations. We need a type conversion back and forth from a 2-D MLIR vector to a 1-D flattened LLVM vector.shape_cast lowering to LLVM is supported in that particular case, for now.

Example:

// Example casting to a lower vector rank.
%1 = vector.shape_cast %0 : vector<5x1x4x3xf32> to vector<20x3xf32>

// Example casting to a higher vector rank.
%3 = vector.shape_cast %2 : vector<10x12x8xf32> to vector<5x2x3x4x8xf32>

// Example casting a tuple of vectors of same rank, where tuple elements
// may have different shapes.
%5 = vector.shape_cast %4 : tuple<vector<3x4x2xf32>, vector<3x3x2xf32>> to
                            tuple<vector<12x2xf32>, vector<9x2xf32>>

Operands:

OperandDescription
sourcevector of any type values or tuple with any combination of vector of any type values values

Results:

ResultDescription
resultvector of any type values or tuple with any combination of vector of any type values values

vector.shuffle (vector::ShuffleOp)

shuffle operation

The shuffle operation constructs a permutation (or duplication) of elements from two input vectors, returning a vector with the same element type as the input and a length that is the same as the shuffle mask. The two input vectors must have the same element type, rank, and trailing dimension sizes and shuffles their values in the leading dimension (which may differ in size) according to the given mask. The legality rules are:

  • the two operands must have the same element type as the result
  • the two operands and the result must have the same rank and trailing dimension sizes, viz. given two k-D operands v1 : <s_1 x s_2 x .. x s_k x type> and 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)

Example:

%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>

Attributes:

AttributeMLIR TypeDescription
maskArrayAttr64-bit integer array attribute

Operands:

OperandDescription
v1vector of any type values
v2vector of any type values

Results:

ResultDescription
vectorvector of any type values

vector.strided_slice (vector::StridedSliceOp)

strided_slice operation

Syntax:

operation ::= `vector.strided_slice` $vector attr-dict `:` type($vector) `to` type(results)

Takes an n-D vector, k-D offsets integer array attribute, a k-sized sizes integer array attribute, a k-sized strides integer array attribute and extracts the n-D subvector at the proper offset.

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.

Example:

%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>

Attributes:

AttributeMLIR TypeDescription
offsetsArrayAttr64-bit integer array attribute
sizesArrayAttr64-bit integer array attribute
stridesArrayAttr64-bit integer array attribute

Operands:

OperandDescription
vectorvector of any type values

Results:

ResultDescription
«unnamed»vector of any type values

vector.transfer_read (vector::TransferReadOp)

Reads a supervector from memory into an SSA vector value.

The vector.transfer_read op performs a blocking read from a slice within a MemRef supplied as its first operand into a vector of the same base elemental type.

A memref 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, supplied as the operands 2 .. 1 + rank(memref). The permutation_map attribute is an affine-map which specifies the transposition on the slice to match the vector shape. The size of the slice is specified by the size of the vector, given as the return type. An ssa-value of the same elemental type as the MemRef is provided as the last operand to specify padding in the case of out-of-bounds accesses. 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 code.

More precisely, let’s dive deeper into the permutation_map for the following 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.

A notional lowering of vector.transfer_read could generate code resembling:

// %expr1, %expr2, %expr3, %expr4 defined before this point
%tmp = alloc() : vector<3x4x5xf32>
%view_in_tmp = "element_type_cast"(%tmp) : memref<1xvector<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>
}}}
%c0 = constant 0 : index
%vec = load %view_in_tmp[%c0] : vector<3x4x5xf32>

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:

// %expr1, %expr2, %expr3, %expr4 defined before this point
%tmp = alloc() : vector<3x4x5xf32>
%view_in_tmp = "element_type_cast"(%tmp) : memref<1xvector<3x4x5xf32>>
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>
}}
%c0 = constant 0 : index
%tmpvec = load %view_in_tmp[%c0] : vector<3x4x5xf32>
%vec = broadcast %tmpvec, 1 : 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.

Syntax

operation ::= ssa-id `=` `vector.transfer_read` ssa-use-list
  `{` attribute-entry `} :` memref-type `,` vector-type

Example:

// 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)
           {permutation_map: (d0, d1, d2) -> (d2, d1)} :
           memref<?x?x?xf32>, vector<32x256xf32>
}}}

// Read the slice `%A[%i0, %i1]` (i.e. the element `%A[%i0, %i1]`) into
// vector<128xf32>. The underlying implementation will require a 1-D vector
// broadcast:
for %i0 = 0 to %0 {
  affine.for %i1 = 0 to %1 {
    %3 = vector.transfer_read %A[%i0, %i1]
         {permutation_map: (d0, d1) -> (0)} :
         memref<?x?xf32>, vector<128xf32>
  }
}

// Read from a memref with vector element type.
%4 = vector.transfer_read %arg1[%c3, %c3], %vf0
  {permutation_map = (d0, d1)->(d0, d1)}
    : memref<?x?xvector<4x3xf32>>, vector<1x1x4x3xf32>

Attributes:

AttributeMLIR TypeDescription
permutation_mapAffineMapAttrAffineMap attribute

Operands:

OperandDescription
memrefmemref of any type values
indicesindex
paddingany type

Results:

ResultDescription
vectorvector of any type values

vector.transfer_write (vector::TransferWriteOp)

The vector.transfer_write op writes a supervector to memory.

Syntax:

operation ::= `vector.transfer_write` $vector `,` $memref `[` $indices `]` attr-dict `:` type($vector) `,`
              type($memref)

The vector.transfer_write performs a blocking write from a vector , supplied as its first operand, into a slice within a MemRef of the same base elemental type, supplied as its second operand.

A vector memref operand 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, supplied as the operands 3 .. 2 + rank(memref). The permutation_map attribute is an affine-map which specifies the transposition on the slice to match the vector shape. The size of the slice is specified by the size of the vector. This operation is called ‘write’ by opposition to ‘store’ because the super-vector granularity is generally not representable with a single hardware register. A vector.transfer_write is thus a 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.

Example:

// write vector<16x32x64xf32> into the slice
//   `%A[%i0, %i1:%i1+32, %i2:%i2+64, %i3:%i3+16]`:
for %i0 = 0 to %0 {
  affine.for %i1 = 0 to %1 step 32 {
    affine.for %i2 = 0 to %2 step 64 {
      affine.for %i3 = 0 to %3 step 16 {
        %val = `ssa-value` : vector<16x32x64xf32>
        vector.transfer_write %val, %A[%i0, %i1, %i2, %i3]
          {permutation_map: (d0, d1, d2, d3) -> (d3, d1, d2)} :
          vector<16x32x64xf32>, memref<?x?x?x?xf32>
}}}}

// write to a memref with vector element type.
vector.transfer_write %4, %arg1[%c3, %c3]
  {permutation_map = (d0, d1)->(d0, d1)}
    : vector<1x1x4x3xf32>, memref<?x?xvector<4x3xf32>>

Attributes:

AttributeMLIR TypeDescription
permutation_mapAffineMapAttrAffineMap attribute

Operands:

OperandDescription
vectorvector of any type values
memrefmemref of any type values
indicesindex

vector.transpose (vector::TransposeOp)

vector transpose operation

Syntax:

operation ::= `vector.transpose` $vector `,` $transp attr-dict `:` type($vector) `to` type($result)

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

%1 = vector.transpose %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:

%1 = vector.transpose %0, [1, 0] : vector<2x3xf32> to vector<3x2xf32>

 [ [a, b, c],       [ [a, d],
   [d, e, f] ]  ->    [b, e],
                      [c, f] ]

Attributes:

AttributeMLIR TypeDescription
transpArrayAttr64-bit integer array attribute

Operands:

OperandDescription
vectorvector of any type values

Results:

ResultDescription
resultvector of any type values

vector.tuple_get (vector::TupleGetOp)

vector tuple get operation

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.

Example:

%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>>

Attributes:

AttributeMLIR TypeDescription
indexIntegerAttrarbitrary integer attribute

Operands:

OperandDescription
vectorstuple with any combination of vector of any type values values

Results:

ResultDescription
«unnamed»vector of any type values

vector.tuple (vector::TupleOp)

make tuple of vectors operation

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.

Example:

%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>

Operands:

OperandDescription
vectorsvector of any type values

Results:

ResultDescription
«unnamed»tuple with any combination of vector of any type values values

vector.type_cast (vector::TypeCastOp)

type_cast op converts a scalar memref to a vector memref

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:

operation ::= `vector.type_cast` ssa-use : memref-type to memref-type

Example:

%A  = alloc() : memref<5x4x3xf32>
%VA = vector.type_cast %A : memref<5x4x3xf32> to memref<vector<5x4x3xf32>>

Operands:

OperandDescription
memrefstatically shaped memref of any type values

Results:

ResultDescription
«unnamed»memref of any type values