!llvm.type<"{<4 x float>*, i64}">
```
-### Function Types {#function-types}
+### Function Types
Function types get converted to LLVM function types. The arguments are converted
individually according to these rules. The result types need to accommodate the
[TOC]
-## Restrictions on Dimension and Symbols {#restrictions-on-dimensions-and-symbols}
+## Restrictions on Dimensions and Symbols
The affine dialect imposes certain restrictions on dimension and symbolic
identifiers to enable powerful analysis and transformation. A symbolic
identifier can be bound to an SSA value that is either an argument to the
function, a value defined at the top level of that function (outside of all
loops and if operations), the result of a
-[`constant` operation](../LangRef.md#'constant'-operation), or the result of an
-[`affine.apply` operation](#'affine.apply'-operation) that recursively takes as
+[`constant` operation](../LangRef.md#constant-operation), or the result of an
+[`affine.apply` operation](#affineapply-operation) that recursively takes as
arguments any symbolic identifiers. Dimensions may be bound not only to anything
that a symbol is bound to, but also to induction variables of enclosing
-[`affine.for` operations](#'affine.for'-operation), and the result of an
-[`affine.apply` operation](#'affine.apply'-operation) (which recursively may use
+[`affine.for` operations](#affinefor-operation), and the result of an
+[`affine.apply` operation](#affineapply-operation) (which recursively may use
other dimensions and symbols).
-## Operations {#operations}
+## Operations
-#### 'affine.apply' operation {#'affine.apply'-operation}
+#### 'affine.apply' operation
Syntax:
%2 = affine.apply (i)[s0] -> (i+s0) (%42)[%n]
```
-#### 'affine.for' operation {#'affine.for'-operation}
+#### 'affine.for' operation
Syntax:
The `affine.for` operation represents an affine loop nest. It has one region
containing its body. This region must contain one block that terminates with
-[`affine.terminator`](#'affine.terminator"-operation). *Note:* when `affine.for`
-is printed in custom format, the terminator is omitted. The block has one
-argument of [`index`](../LangRef.md#index-type) type that represents the induction
+[`affine.terminator`](#affineterminator-operation). *Note:* when `affine.for` is
+printed in custom format, the terminator is omitted. The block has one argument
+of [`index`](../LangRef.md#index-type) type that represents the induction
variable of the loop.
The `affine.for` operation executes its body a number of times iterating from a
}
```
-#### 'affine.if' operation {#'affine.if'-operation}
+#### 'affine.if' operation
Syntax:
[same restrictions](#restrictions-on-dimensions-and-symbols) hold for these SSA
values as for all bindings of SSA values to dimensions and symbols.
-The `if` operation contains two regions for the "then" and "else" clauses. The
-latter may be empty (i.e. contain no blocks), meaning the absence of the else
-clause. When non-empty, both regions must contain exactly one block terminating
-with [`affine.terminator`](#'affine.terminator'-operation). *Note:* when `if` is
-printed in custom format, the terminator is omitted. These blocks must not have
-any arguments.
+The `affine.if` operation contains two regions for the "then" and "else"
+clauses. The latter may be empty (i.e. contain no blocks), meaning the absence
+of the else clause. When non-empty, both regions must contain exactly one block
+terminating with [`affine.terminator`](#affineterminator-operation). *Note:*
+when `affine.if` is printed in custom format, the terminator is omitted. These
+blocks must not have any arguments.
Example:
}
```
-#### `affine.terminator` operation {#'affine.terminator'-operation}
+#### `affine.terminator` operation
Syntax:
```
Affine terminator is a special terminator operation for blocks inside affine
-loops ([`for`](#'for'-operation)) and branches ([`if`](#'if'-operation)). It
-unconditionally transmits the control flow to the successor of the operation
-enclosing the region.
+loops ([`affine.for`](#affinefor-operation)) and branches
+([`affine.if`](#affineif-operation)). It unconditionally transmits the control
+flow to the successor of the operation enclosing the region.
-*Rationale*: bodies of affine operations are [blocks](../LangRef.md#block) that
+*Rationale*: bodies of affine operations are [blocks](../LangRef.md#blocks) that
must have terminators. Loops and branches represent structured control flow and
should not accept arbitrary branches as terminators.
`getLLVMModule()`. All LLVM IR objects that interact with the LLVM IR dialect
must exist in the dialect's context.
-## Types {#types}
+## Types
The LLVM IR dialect defines a single MLIR type, `LLVM::LLVMType`, that can wrap
any existing LLVM IR type. Its syntax is as follows
term "wrapped LLVM IR type" to refer to the LLVM IR dialect type containing a
specific LLVM IR type.
-## Operations {#operations}
+## Operations
All operations in the LLVM IR dialect have a custom form in MLIR. The mnemonic
of an operation is that used in LLVM IR prefixed with "`llvm.`".
Selection: `select <condition>, <lhs>, <rhs>`.
-### Pseudo-operations {#pseudo-operations}
+### Pseudo-operations
These operations do not have LLVM IR counterparts but are necessary to map LLVM
IR into MLIR.
-#### `llvm.constant` {#constant-operation}
+#### `llvm.constant`
Unlike LLVM IR, MLIR does not have first-class constant values. Therefore, all
constants must be created as SSA values before being used in other operations.
%3 = llvm.constant(splat<vector<4xf32>, 1.0>) : !llvm<"<4 x float>">
```
-#### `llvm.undef` {#undef-operation}
+#### `llvm.undef`
Unlike LLVM IR, MLIR does not have first-class undefined values. Such values
must be created as SSA values using `llvm.undef`. This operation has no operands
[TOC]
-## Operations {#operations}
+## Operations
-### Vector transfers {#vector-transfers}
+### Vector transfers
-#### `vector.transfer_read` operation {#'vector.transfer_read'-operation}
+#### `vector.transfer_read` operation
Syntax:
broadcast is required. On a GPU this broadcast could be implemented using a
warp-shuffle if loop `j` were mapped to `threadIdx.x`.
-#### `vector.transfer_write` operation {#'vector.transfer_write'-operation}
+#### `vector.transfer_write` operation
Syntax:
ensure the memory writes are valid. Different lowerings may be pertinent
depending on the hardware support.
-### Vector views {#vector-views}
+### Vector views
-#### `vector.type_cast` operation {#'vector.type_cast'-operation}
+#### `vector.type_cast` operation
Syntax:
[TOC]
-## High-Level Structure {#high-level-structure}
+## High-Level Structure
The top-level unit of code in MLIR is a [Module](#module). A module contains a
list of [Functions](#functions). Functions are represented as a composition of
-[operations](#operations) and contain a Control Flow Graph (CFG) of
+[Operations](#operations) and contain a Control Flow Graph (CFG) of
[Blocks](#blocks), which contain operations and end with
[terminator operations](#terminator-operations) (like branches).
}
```
-## Notation {#notation}
+## Notation
MLIR has a simple and unambiguous grammar, allowing it to reliably round-trip
through a textual form. This is important for development of the compiler - e.g.
example ::= `b` (`an` | `om`)* `a`
```
-### Common syntax {#common-syntax}
+### Common syntax
The following core grammar productions are used in this document:
Not listed here, but MLIR does support comments. They use standard BCPL syntax,
starting with a `//` and going until the end of the line.
-### Identifiers and keywords {#identifiers-and-keywords}
+### Identifiers and keywords
Syntax:
identifiers in mapping functions are in scope for the mapping body. Function
identifiers and mapping identifiers are visible across the entire module.
-## Polyhedral Structures {#polyhedral-structures}
+## Polyhedral Structures
MLIR uses techniques from polyhedral compilation to make dependence analysis and
loop transformations efficient and reliable. This section introduces some of the
core concepts that are used throughout the document.
-### Dimensions and Symbols {#dimensions-and-symbols}
+### Dimensions and Symbols
Dimensions and symbols are the two kinds of identifiers that can appear in the
-polyhedral structures, and are always of '[index](#index-type)' type. Dimensions
+polyhedral structures, and are always of [`index`](#index-type) type. Dimensions
are declared in parentheses and symbols are declared in square brackets.
Examples:
%x = alloc()[%N] : memref<40x50xf32, #affine_map2to3>
```
-### Affine Expressions {#affine-expressions}
+### Affine Expressions
Syntax:
j)$$ are two-dimensional affine functions of $$(i, j)$$, but $$(i \cdot j,
i^2)$$, $$(i \mod j, i/j)$$ are not affine functions of $$(i, j)$$.
-### Affine Maps {#affine-maps}
+### Affine Maps
Syntax:
combining the indices and symbols. Affine maps distinguish between
[indices and symbols](#dimensions-and-symbols) because indices are inputs to the
affine map when the latter may be called through an operation, such as
-[affine.apply](Dialects/Affine.md#'affine.apply'-operation) operation, whereas
+[affine.apply](Dialects/Affine.md#affineapply-operation) operation, whereas
symbols are bound when an affine mapping is established (e.g. when a memref is
formed, establishing a memory [layout map](#layout-map)).
impose on their form allows powerful analysis and transformation, while keeping
the representation closed with respect to several operations of interest.
-#### Named affine mappings {#named-affine-mappings}
+#### Named affine mappings
Syntax:
size (10, s0)>
```
-### Semi-affine maps {#semi-affine-maps}
+### Semi-affine maps
Semi-affine maps are extensions of affine maps to allow multiplication,
`floordiv`, `ceildiv`, and `mod` with respect to symbolic identifiers.
semi-affine-map ::= semi-affine-map-id | semi-affine-map-inline
```
-### Integer Sets {#integer-sets}
+### Integer Sets
An integer set is a conjunction of affine constraints on a list of identifiers.
The identifiers associated with the integer set are separated out into two
MLIR provides a first class set of polyhedral operations and analyses within the
[affine dialect](Dialects/Affine.md).
-## Type System {#type-system}
+## Type System
Each SSA value in MLIR has a type defined by the type system below. There are a
number of primitive types (like integers) and also aggregate types for tensors
ssa-use-and-type-list ::= ssa-use-and-type (`,` ssa-use-and-type)*
```
-### Type Aliases {#type-aliases}
+### Type Aliases
``` {.ebnf}
type-alias-def ::= '!' alias-name '=' 'type' type
"foo"(%x) : !avx.m128 -> ()
```
-### Builtin Types {#builtin-types}
+### Builtin Types
Builtin types consist of only the types needed for the validity of the IR.
-#### Function Type {#function-type}
+#### Function Type
Syntax:
```
MLIR supports first-class functions: the
-[`constant` operation](#'constant'-operation) produces the address of a function
+[`constant` operation](#constant-operation) produces the address of a function
as an SSA value. This SSA value may be passed to and returned from functions,
merged across control flow boundaries with [block arguments](#blocks), and
-called with the [`call_indirect` operation](#'call_indirect'-operation).
+called with the [`call_indirect` operation](#call-indirect-operation).
Function types are also used to indicate the arguments and results of
[operations](#operations).
-### Standard Types {#standard-types}
+### Standard Types
-#### Index Type {#index-type}
+#### Index Type
Syntax:
machine word of the target ([rationale](Rationale.md#signless-types)) and is
used by the affine constructs in MLIR. Unlike fixed-size integers. It cannot be
used as an element of vector, tensor or memref type
-([rationale](Rationale.md#index-type-disallowed-in-aggregate-types)).
+([rationale](Rationale.md#index-type-disallowed-in-vectortensormemref-types)).
**Rationale:** integers of platform-specific bit widths are practical to express
sizes, dimensionalities and subscripts.
-#### Integer Type {#integer-type}
+#### Integer Type
Syntax:
TODO: Need to decide on a representation for quantized integers
([initial thoughts](Rationale.md#quantized-integer-operations)).
-#### Floating Point Types {#floating-point-types}
+#### Floating Point Types
Syntax:
MLIR supports float types of certain widths that are widely used as indicated
above.
-#### Vector Type {#vector-type}
+#### Vector Type
Syntax:
`vector<0x42xi32>` is invalid because it is interpreted as a 2D vector with
shape `(0, 42)` and zero shapes are not allowed.
-#### Tensor Type {#tensor-type}
+#### Tensor Type
Syntax:
access, MLIR has a [`memref` type](#memref-type). This abstracted runtime
representation holds both the tensor data values as well as information about
the (potentially dynamic) shape of the tensor. The
-[`dim` operation](#'dim'-operation) returns the size of a dimension from a value
+[`dim` operation](#dim-operation) returns the size of a dimension from a value
of tensor type.
Note: hexadecimal integer literals are not allowed in tensor type declarations
tensor<0xf32>
```
-#### Memref Type {#memref-type}
+#### Memref Type
Syntax:
%A = <strong>alloc</strong> (%n)[%o] : <16x?xf32, #imapA>
```
-##### Index Space {#index-space}
+##### Index Space
A memref dimension list defines an index space within which the memref can be
indexed to access data.
-##### Index {#index}
+##### Index
Data is accessed through a memref type using a multidimensional index into the
multidimensional index space defined by the memref's dimension list.
%v = load %A[%i, %j] : memref<16x32xf32, #imapA, hbm>
```
-##### Index Map {#index-map}
+##### Index Map
An index map is a one-to-one [semi-affine map](#semi-affine-maps) that
transforms a multidimensional index from one index space to another. For
size (M, N)
```
-##### Layout Map {#layout-map}
+##### Layout Map
A layout map is a [semi-affine map](#semi-affine-maps) which encodes logical to
physical index space mapping, by mapping input dimensions to their ordering from
#layout_map_col_major = (i, j), [M, N] -> (j, i) size (M, N)
```
-##### Affine Map Composition {#affine-map-composition}
+##### Affine Map Composition
A memref specifies a semi-affine map composition as part of its type. A
semi-affine map composition is a composition of semi-affine maps beginning with
copy elision and in-place updates. If an affine map composition is not specified
for the memref, the identity affine map is assumed.
-#### Complex Type {#complex-type}
+#### Complex Type
Syntax:
complex<i32>
```
-#### Tuple Type {#tuple-type}
+#### Tuple Type
Syntax:
**Rationale:** Though this type is first class in the type system, MLIR provides
no standard operations for operating on `tuple` types
-[rationale](Rationale.md#tuple-type).
+([rationale](Rationale.md#tuple-types)).
Examples:
tuple<i32, f32, tensor<i1>, i5>
```
-## Attributes {#attributes}
+## Attributes
Syntax:
Attributes are the mechanism for specifying constant data in MLIR in places
where a variable is never allowed - e.g. the index of a
-[`dim` operation](#'dim'-operation), or the stride of a convolution. They
-consist of a name and a [concrete attribute value](#attribute-values). It is
-possible to attach attributes to operations, functions, and function arguments.
-The set of expected attributes, their structure, and their interpretation are
-all contextually dependent on what they are attached to.
+[`dim` operation](#dim-operation), or the stride of a convolution. They consist
+of a name and a [concrete attribute value](#attribute-values). It is possible to
+attach attributes to operations, functions, and function arguments. The set of
+expected attributes, their structure, and their interpretation are all
+contextually dependent on what they are attached to.
There are two main classes of attributes; dependent and dialect. Dependent
attributes derive their structure and meaning from what they are attached to,
distinct semantic context, and can thus provide a single source of meaning to
dependent attributes.
-### Attribute Values {#attribute-values}
+### Attribute Values
Attributes values are represented by the following forms:
| type-attribute
```
-#### AffineMap Attribute {#affine-map-attribute}
+#### AffineMap Attribute
Syntax:
An affine-map attribute is an attribute that represents a affine-map object.
-#### Array Attribute {#array-attribute}
+#### Array Attribute
Syntax:
An array attribute is an attribute that represents a collection of attribute
values.
-#### Boolean Attribute {#bool-attribute}
+#### Boolean Attribute
Syntax:
A boolean attribute is a literal attribute that represents a one-bit boolean
value, true or false.
-#### Elements Attributes {#elements-attributes}
+#### Elements Attributes
Syntax:
An elements attribute is a literal attribute that represents a constant
[vector](#vector-type) or [tensor](#tensor-type) value.
-##### Dense Elements Attribute {#dense-elements-attribute}
+##### Dense Elements Attribute
Syntax:
element type of the vector or tensor constant must be of integer, index, or
floating point type.
-##### Opaque Elements Attribute {#opaque-elements-attribute}
+##### Opaque Elements Attribute
Syntax:
Note: The parsed string literal must be in hexadecimal form.
-##### Sparse Elements Attribute {#sparse-elements-attribute}
+##### Sparse Elements Attribute
Syntax:
/// [0, 0, 0, 0]]
```
-##### Splat Elements Attribute {#splat-elements-attribute}
+##### Splat Elements Attribute
Syntax:
A splat elements attribute is an elements attribute that represents a tensor or
vector constant where all elements have the same value.
-#### Integer Attribute {#integer-attribute}
+#### Integer Attribute
Syntax:
the specified integer or index type. The default type for this attribute, if one
is not specified, is a 64-bit integer.
-#### Integer Set Attribute {#integer-set-attribute}
+#### Integer Set Attribute
Syntax:
An integer-set attribute is an attribute that represents a integer-set object.
-#### Float Attribute {#float-attribute}
+#### Float Attribute
Syntax:
A float attribute is a literal attribute that represents a floating point value
of the specified [float type](#floating-point-types).
-#### Function Attribute {#function-attribute}
+#### Function Attribute
Syntax:
A function attribute is a literal attribute that represents a reference to the
given function object.
-#### String Attribute {#string-attribute}
+#### String Attribute
Syntax:
A string attribute is an attribute that represents a string literal value.
-#### Type Attribute {#type-attribute}
+#### Type Attribute
Syntax:
A type attribute is an attribute that represents a [type object](#type-system).
-## Module {#module}
+## Module
``` {.ebnf}
module ::= module-header-def* function*
and will define the set of operations that are allowed (allowing the verifier to
detect common errors).
-## Functions {#functions}
+## Functions
MLIR functions have a signature (including argument and result types) and
associated attributes according to the following grammar:
}
```
-#### Blocks {#blocks}
+#### Blocks
Syntax:
case: they become arguments to the entry block
[[more rationale](Rationale.md#block-arguments-vs-phi-nodes)].
-### Operations {#operations}
+### Operations
Syntax:
infrastructure uses C++ templates to make working with them convenient and safe.
The details of this are not described in this document.
-## Standard Operations {#standard-operations}
+## Standard Operations
TODO: shape, which returns a 1D tensor, and can take an unknown rank tensor as
input.
TODO: rank, which returns an index.
-#### Terminator operations {#terminator-operations}
+#### Terminator operations
Terminator operations are required at the end of each block. They may contain a
list of successors, i.e. other blocks to which the control flow will proceed.
Currently, all terminator operations must be registered in some known
[dialect](#dialects), unlike regular operations.
-##### 'br' terminator operation {#'br'-terminator-operation}
+##### 'br' terminator operation
Syntax:
The MLIR branch operation is not allowed to target the entry block for a
function.
-##### 'cond_br' terminator operation {#'cond_br'-terminator-operation}
+##### 'cond_br' terminator operation
Syntax:
}
```
-##### 'return' terminator operation {#'return'-terminator-operation}
+##### 'return' terminator operation
Syntax:
result types of the enclosing function. It is legal for multiple blocks in a
single function to return.
-### Core Operations {#core-operations}
+### Core Operations
-#### 'call' operation {#'call'-operation}
+#### 'call' operation
Syntax:
%31 = call @my_add(%0, %1) : (tensor<16xf32>, tensor<16xf32>) -> tensor<16xf32>
```
-#### 'call_indirect' operation {#'call_indirect'-operation}
+#### 'call_indirect' operation
Syntax:
call must match the specified function type.
Function values can be created with the
-[`constant` operation](#'constant'-operation).
+[`constant` operation](#constant-operation).
Example:
: (tensor<16xf32>, tensor<16xf32>) -> tensor<16xf32>
```
-#### 'dim' operation {#'dim'-operation}
+#### 'dim' operation
Syntax:
```
The `dim` operation takes a memref or tensor operand and a dimension index, and
-returns an ['index'](#index-type) that is the size of that dimension.
+returns an [`index`](#index-type) that is the size of that dimension.
The `dim` operation is represented with a single integer attribute named
`index`, and the type specifies the type of the memref or tensor operand.
%y = "std.dim"(%A){index: 1} : (tensor<4 x ? x f32>) -> index
```
-#### 'reshape' operation {#'reshape'-operation}
+#### 'reshape' operation
Syntax:
```
-#### 'view' operation {#'view'-operation}
+#### 'view' operation
Syntax:
(%s1) [%0, %n1] %A : memref<16x?xf32, #map_a, hbm>
```
-### Memory Operations {#memory-operations}
+### Memory Operations
-#### 'alloc' operation {#'alloc'-operation}
+#### 'alloc' operation
Syntax:
%B = alloc(%M, %N)[%x, %y] : memref<?x?xf32, #layout_map1, vmem>
```
-#### 'alloc_static' operation {#'alloc_static'-operation}
+#### 'alloc_static' operation
Syntax:
Allocates a new memref of specified type with a fixed base pointer location in
memory. 'alloc_static' does not support types that have dynamic shapes or that
require dynamic symbols in their layout function (use the
-[`alloc'`operation](#'alloc'-operation) in those cases).
+[`alloc` operation](#alloc-operation) in those cases).
Example:
The `alloc_static` operation is used to represent code after buffer allocation
has been performed.
-#### 'dealloc' operation {#'dealloc'-operation}
+#### 'dealloc' operation
Syntax:
```
Delineates the end of the lifetime of the memory corresponding to a memref
-allocation. It is paired with an [`alloc`](#'alloc'-operation) or
-[`alloc_static`](#'alloc_static'-operation) operation.
+allocation. It is paired with an [`alloc`](#alloc-operation) or
+[`alloc_static`](#alloc-static-operation) operation.
Example:
dma_wait %tag[%index], %num_elements : memref<1 x i32, (d0) -> (d0), 4>
```
-#### 'extract_element' operation {#'extract_element'-operation}
+#### 'extract_element' operation
Syntax:
%5 = extract_element %ut[%1, %2] : tensor<*xi32>
```
-#### 'load' operation {#'load'-operation}
+#### 'load' operation
Syntax:
In an `affine.if` or `affine.for` body, the indices of a load are restricted to
SSA values bound to surrounding loop induction variables,
[symbols](#dimensions-and-symbols), results of a
-[`constant` operation](#'constant'-operation), or the result of an
-`affine.apply` operation that can in turn take as arguments all of the
-aforementioned SSA values or the recursively result of such an `affine.apply`
-operation.
+[`constant` operation](#constant-operation), or the result of an `affine.apply`
+operation that can in turn take as arguments all of the aforementioned SSA
+values or the recursively result of such an `affine.apply` operation.
Example:
**Context:** The `load` and `store` operations are specifically crafted to fully
resolve a reference to an element of a memref, and (in affine `affine.if` and
`affine.for` operations) the compiler can follow use-def chains (e.g. through
-[`affine.apply`](Dialects/Affine.md#'affine.apply'-operation) operations) to
+[`affine.apply`](Dialects/Affine.md#affineapply-operation) operations) to
precisely analyze references at compile-time using polyhedral techniques. This
is possible because of the
[restrictions on dimensions and symbols](Dialects/Affine.md#restrictions-on-dimensions-and-symbols)
in these contexts.
-#### 'store' operation {#'store'-operation}
+#### 'store' operation
Syntax:
In an affine context, the indices of a store are restricted to SSA values bound
to surrounding loop induction variables,
-[symbols](Dialect/Affine.md#restrictions-on-dimensions-and-symbols), results of
-a [`constant` operation](#'constant'-operation), or the result of an
-[`affine.apply`](Dialect/Affine.md#'affine.apply'-operation) operation that can
-in turn take as arguments all of the aforementioned SSA values or the
-recursively result of such an `affine.apply` operation.
+[symbols](Dialects/Affine.md#restrictions-on-dimensions-and-symbols), results of
+a [`constant` operation](#constant-operation), or the result of an
+[`affine.apply`](Dialects/Affine.md#affineapply-operation) operation that can in
+turn take as arguments all of the aforementioned SSA values or the recursively
+result of such an `affine.apply` operation.
Example:
**Context:** The `load` and `store` operations are specifically crafted to fully
resolve a reference to an element of a memref, and (in polyhedral `affine.if`
and `affine.for` operations) the compiler can follow use-def chains (e.g.
-through [`affine.apply`](Dialects/Affine.md#'affine.apply'-operation)
-operations) to precisely analyze references at compile-time using polyhedral
-techniques. This is possible because of the
-[restrictions on dimensions and symbols](Dialect/Affine.md#restrictions-on-dimensions-and-symbols)
+through [`affine.apply`](Dialects/Affine.md#affineapply-operation) operations)
+to precisely analyze references at compile-time using polyhedral techniques.
+This is possible because of the
+[restrictions on dimensions and symbols](Dialects/Affine.md#restrictions-on-dimensions-and-symbols)
in these contexts.
-#### 'tensor_load' operation {#'tensor_load'-operation}
+#### 'tensor_load' operation
Syntax:
%12 = tensor_load %10 : memref<4x?xf32, #layout, hbm>
```
-#### 'tensor_store' operation {#'tensor_store'-operation}
+#### 'tensor_store' operation
Syntax:
tensor_store %8, %10 : memref<4x?xf32, #layout, hbm>
```
-### Arithmetic Operations {#arithmetic-operations}
+### Arithmetic Operations
Basic arithmetic in MLIR is specified by standard operations described in this
section.
these on demand. We should be highly informed by and learn from the operations
supported by HLO and LLVM.
-#### 'addi' operation {#'addi'-operation}
+#### 'addi' operation
Examples:
whose element type is integer, or a tensor of integers. It has no standard
attributes.
-#### 'addf' operation {#'addf'-operation}
+#### 'addf' operation
Examples:
optional attributes for fast math, contraction, rounding mode, and other
controls.
-#### 'cmpi' operation {#'cmpi'-operation}
+#### 'cmpi' operation
Examples:
etc ([rationale](Rationale.md#splitting-floating-point-vs-integer-operations)).
The type of comparison is specified as attribute to avoid introducing ten
similar operations, taking into account that they are often implemented using
-the same operation downstream ([rationale](Rationale.md#cmpi-predicate)). The
+the same operation downstream
+([rationale](Rationale.md#specifying-comparison-kind-as-attribute)). The
separation between signed and unsigned order comparisons is necessary because of
integers being signless. The comparison operation must know how to interpret
values with the foremost bit being set: negatives in two's complement or large
-positives ([rationale](Rationale.md#sign-in-cmpi)).
+positives
+([rationale](Rationale.md#specifying-sign-in-integer-comparison-operations)).
-#### 'constant' operation {#'constant'-operation}
+#### 'constant' operation
Syntax:
directly reference a function simplifies this
([rationale](Rationale.md#multithreading-the-compiler)).
-#### 'divis' operation {#'divis'-operation}
+#### 'divis' operation
Signed integer division. Rounds towards zero. Treats the leading bit as sign,
i.e. `6 / -2 = -3`.
vector whose element type is integer, or a tensor of integers. It has no
standard attributes.
-#### 'diviu' operation {#'diviu'-operation}
+#### 'diviu' operation
Unsigned integer division. Rounds towards zero. Treats the leading bit as the
most significant, i.e. for `i16` given two's complement representation, `6 /
source and destination types may not be the same. The operation is invalid if
converting to a mismatching constant dimension.
-#### 'mulf' operation {#'mulf'-operation}
+#### 'mulf' operation
Examples:
optional attributes for fast math, contraction, rounding mode, and other
controls.
-#### 'remis' operation {#'remis'-operation}
+#### 'remis' operation
Signed integer division remainder. Treats the leading bit as sign, i.e. `6 %
-2 = 0`.
vector whose element type is integer, or a tensor of integers. It has no
standard attributes.
-#### 'remiu' operation {#'remiu'-operation}
+#### 'remiu' operation
Unsigned integer division remainder. Treats the leading bit as the most
significant, i.e. for `i16`, `6 % -2 = 6 % (2^16 - 2) = 6`.
vector whose element type is integer, or a tensor of integers. It has no
standard attributes.
-#### 'select' operation {#'select'-operation}
+#### 'select' operation
Syntax:
all operands is identical. The choice is made for each element individually
based on the value at the same position as the element in the condition operand.
-The `select` operation combined with [`cmpi`](#'cmpi'-operation) can be used to
+The `select` operation combined with [`cmpi`](#cmpi-operation) can be used to
implement `min` and `max` with signed or unsigned comparison semantics.
-#### 'tensor_cast' operation {#'tensor_cast'-operation}
+#### 'tensor_cast' operation
Syntax:
They must either have the same rank, or one may be an unknown rank. The
operation is invalid if converting to a mismatching constant dimension.
-## Dialects {#dialects}
+## Dialects
MLIR supports multiple dialects containing a set of operations and types defined
together, potentially outside of the main tree. Dialects are produced and
* [Vector dialect](Dialects/Vector.md)
* [TensorFlow dialect](#tensorflow-operations)
-### TensorFlow operations {#tensorflow-operations}
+### TensorFlow operations
MLIR operations can represent arbitrary TensorFlow operations with a reversible
mapping. Switch and merge nodes are represented with the MLIR control flow
: (tensor<*xf32>, tensor<*xf32>) -> tensor<*xf32>
```
-### Target specific operations {#target-specific-operations}
+### Target specific operations
We expect to expose many target-specific (such as TPU-specific) operations
directly through to MLIR.
These operations only work when targeting LLVM as a backend (e.g. for CPUs and
GPUs), and are required to align with the LLVM definition of these intrinsics.
-### Dialect specific types {#dialect-specific-types}
+### Dialect specific types
Similarly to operations, dialects may define custom extensions to the type
system. These extensions fit within the same type system as described in the
!tf<"string">
```
-### TensorFlow types {#tensorflow-types}
+### TensorFlow types
The TensorFlow dialect in MLIR defines several extended types:
These are the results of other operations and mostly only known at
runtime. They can have a fixed type or correspond to set of possible
- types. See [Type constraints](type-constraints) specification below.
+ types. See [Type constraints](#type-constraints) specification below.
1. Attributes.
Traits of the operations. They are operation properties that affect syntax
or semantics. MLIR C++ models various traits in the `mlir::OpTrait`
- namespace. In TableGen, we have the corresponding `OpTrait` class to
- wrap around any C++ trait symbol and use it in operation definition.
- For example, `NoSideEffect` is just a definition that expands to
+ namespace. In TableGen, we have the corresponding `OpTrait` class to wrap
+ around any C++ trait symbol and use it in operation definition. For example,
+ `NoSideEffect` is just a definition that expands to
`OpTrait<"HasNoSideEffect">`; having such a definition makes the trait
inside TableGen more integrated and easier to parse as a declarative
language.
[TOC]
-## Affine control lowering (`-lower-affine`) {#lower-affine-apply}
+## Affine control lowering (`-lower-affine`)
Convert operations related to affine control into a graph of blocks using
operations from the standard dialect.
they do not depend on the loop iterator value or on the result of
`affine.apply`.
-## Conversion from Standard to LLVM IR dialect (`-convert-to-llvmir`) {#convert-to-llvmir}
+## Conversion from Standard to LLVM IR dialect (`-convert-to-llvmir`)
Convert standard operations into the LLVM IR dialect operations.
returns are updated accordingly. Block argument types are updated to use LLVM IR
types.
-## DMA generation (`-dma-generate`) {#dma-generate}
+## DMA generation (`-dma-generate`)
Replaces all loads and stores on memref's living in 'slowMemorySpace' by
introducing DMA operations (strided DMA if necessary) to transfer data to/from
Performs tiling or blocking of loop nests. It currently works on perfect loop
nests.
-## Loop unroll (`-loop-unroll`) {#loop-unroll}
+## Loop unroll (`-loop-unroll`)
This pass implements loop unrolling. It is able to unroll loops with arbitrary
bounds, and generate a cleanup loop when necessary.
-## Loop unroll and jam (`-loop-unroll-jam`) {#loop-unroll-jam}
+## Loop unroll and jam (`-loop-unroll-jam`)
This pass implements unroll and jam for loops. It works on both perfect or
imperfect loop nests.
-## Loop fusion (`-loop-fusion`) {#loop-fusion}
+## Loop fusion (`-loop-fusion`)
Performs fusion of loop nests using a slicing-based approach. The fused loop
nests, when possible, are rewritten to access significantly smaller local
evaluates available choices such as the depth at which a source slice should be
materialized in the designation slice.
-## Memref bound checking (`-memref-bound-check`) {#memref-bound-check}
+## Memref bound checking (`-memref-bound-check`)
Checks all load's and store's on memref's for out of bound accesses, and reports
any out of bound accesses (both overrun and underrun) with location information.
^
```
-## Memref dataflow optimization (`-memref-dataflow-opt`) {#memref-dataflow-opt}
+## Memref dataflow optimization (`-memref-dataflow-opt`)
This pass performs store to load forwarding for memref's to eliminate memory
accesses and potentially the entire memref if all its accesses are forwarded.
}
```
-## Memref dependence analysis (`-memref-dependence-check`) {#memref-dependence-check}
+## Memref dependence analysis (`-memref-dependence-check`)
This pass performs dependence analysis to determine dependences between pairs of
memory operations (load's and store's) on memref's. Dependence analysis exploits
store %cf9, %m[%idx] : memref<10xf32>
```
-## Pipeline data transfer (`-pipeline-data-transfer`) {#pipeline-data-transfer}
+## Pipeline data transfer (`-pipeline-data-transfer`)
This pass performs a transformation to overlap non-blocking DMA operations in a
loop with computations through double buffering. This is achieved by advancing
* *per-axis* (also called *per-channel*) : Applying individually to each index
along a specific axis of a tensor type.
-### Fixed point values {#fixed-point}
+### Fixed point values
[Fixed point](https://en.wikipedia.org/wiki/Fixed-point_arithmetic) values are a
[Real](https://en.wikipedia.org/wiki/Real_number) number divided by a *scale*.
integers, and perform arithmetic on those signed integers, because the results
will be correct scaled values.
-### Affine values {#affine}
+### Affine values
Mathematically speaking, affine values are the result of
[adding a Real-valued *zero point*, to a scaled value](https://en.wikipedia.org/wiki/Affine_transformation#Representation).
rounding should be according to the IEEE754 default of RNE (where hardware
permits).
-### Converting between Real and fixed point or affine {#converting-between}
+### Converting between Real and fixed point or affine
To convert a Real value to a fixed point value, you must know the scale. To
convert a Real value to an affine value, you must know the scale and zero point.
constrained to particular bit depths or a requirement that the entire range of
an N-bit integer is used.
-#### Affine to Real {#affine-to-real}
+#### Affine to Real
To convert an output tensor of affine elements represented by uint8
or uint16 to a tensor of Real-valued elements (usually represented with a
In the above, we assume that the result of subtraction is in 32-bit signed
integer format, and that $$roundToNearestFloat$$ returns a Single.
-#### Affine to fixed point {#affine-to-fixed-point}
+#### Affine to fixed point
When the affine and fixed point scales are the same, subtract the zero point
from the affine value to get the equivalent fixed point value.
scaled\_value = affine\_value_{non\mbox{-}negative} - zero\_point_{non\mbox{-}negative}
$$
-#### Fixed point to affine {#fixed-point-to-affine}
+#### Fixed point to affine
When the affine and fixed point scales are the same, add the zero point to the
fixed point value to get the equivalent affine value.
affine\_value_{non\mbox{-}negative} = scaled\_value + zero\_point_{non\mbox{-}negative}
$$
-## Usage within MLIR {#usage-within-mlir}
+## Usage within MLIR
There are several components to the quantization system being developed within
MLIR:
* [Type conversion ops](#quantized-type-conversion-ops) for converting
between types based on a QuantizedType and its *expressed* and *storage*
sub-types.
- * [Instrumentation ops](#instrumentation-ops) for assigning
+ * [Instrumentation ops](#instrumentation-and-constraint-ops) for assigning
instrumentation points within the computation where runtime statistics
may help guide the quantization process.
-* [Integration with simulated quantization at training time](#fake-quant)
+* [Integration with simulated quantization at training time](#integration-with-simulated-quantization-at-training-time)
* [TFLite native quantization](#tflite-native-quantization)
* Passes and tools exist to convert directly from the *TensorFlow* dialect
to the TFLite quantized op-set.
-* [*FxpMath* dialect](#fxp-math-dialect) containing (experimental) generalized
+* [*FxpMath* dialect](#fxpmath-dialect) containing (experimental) generalized
representations of fixed-point math ops and conversions:
* [Real math ops](#real-math-ops) representing common combinations of
arithmetic operations that closely match corresponding fixed-point math
concepts (as opposed to being spread across multiple ops as is typical
in source dialects).
- * [Fixed-point math ops](#fxp-math-ops) that for carrying out computations
- on integers, as are typically needed by uniform quantization schemes.
+ * [Fixed-point math ops](#fixed-point-math-ops) that for carrying out
+ computations on integers, as are typically needed by uniform
+ quantization schemes.
* Passes to lower from real math ops to fixed-point math ops.
* [Solver tools](#solver-tools) which can (experimentally and generically
TensorFlow to TensorFlow Lite conversion uses the QuantizedTypes but has its own
ops for type conversion and expression of the backing math.
-## Quantization Dialect {#quantization-dialect}
+## Quantization Dialect
-### Quantized type {#quantized-type}
+### Quantized type
TODO : Flesh this section out.
* QuantizedType base class
* UniformQuantizedType
-### Quantized type conversion ops {#quantized-type-conversion-ops}
+### Quantized type conversion ops
* qcast : Convert from an expressed type to QuantizedType
* dcast : Convert from a QuantizedType to its expressed type
* scast : Convert between a QuantizedType and its storage type
-### Instrumentation and constraint ops {#instrumentation-ops}
+### Instrumentation and constraint ops
TODO : These ops are not defined yet
fixed-point values, underlying storage type, or whether to constrain to
power of two scales.
-## Integration with simulated quantization at training time {#fake-quant}
+## Integration with simulated quantization at training time
TensorFlow has historically used the
[tf.quantization.fake_quant_\*](https://www.tensorflow.org/api_docs/python/tf/quantization/fake_quant_with_min_max_args)
where the parts which could not be reduced to integral ops are still carried out
in floating point with appropriate conversions at the boundaries.
-## TFLite Native Quantization {#tflite-native-quantization}
+## TFLite Native Quantization
TODO : Flesh this out
-> tfl.Q) and replaces with (op). Also replace (constant_float -> tfl.Q)
with (constant_quant).
-## FxpMath Dialect {#fxp-math-dialect}
+## FxpMath Dialect
-### Real math ops {#real-math-ops}
+### Real math ops
Note that these all support explicit clamps, which allows for simple fusions and
representation of some common sequences quantization-compatible math. Of
* CMPLZ
* CMPGZ
-### Fixed-point math ops {#fxp-math-ops}
+### Fixed-point math ops
TODO: This op set only has enough ops to lower a simple power-of-two
RealAddEwOp.
* RoundingDivideByPotFxpOp
* SaturatingAddFxpOp
-## Solver tools {#solver-tools}
+## Solver tools
Solver tools exist to analyze an MLIR-computation, expressed in either a
supported source dialect or in the *real math ops* set and solve for appropriate
[TOC]
-## Abstract {#abstract}
+## Abstract
MLIR is a compiler intermediate representation with similarities to traditional
three-address SSA representations (like
provides the rationale behind MLIR -- its actual
[specification document](LangRef.md) and other content is hosted elsewhere.
-## Introduction and Motivation {#introduction-and-motivation}
+## Introduction and Motivation
The Multi-Level Intermediate Representation (MLIR) is intended for easy
expression and optimization of computations involving deep loop nests and dense
* MLIR allows us to build modular and reusable target independent and target
dependent passes - since each pass/emitter can read in another's output.
-## Design Decisions {#design-decisions}
+## Design Decisions
This section sheds light on some of the design decisions -- some of these are
indirectly implied by the specification document.
-### Loads and stores {#loads-and-stores}
+### Loads and stores
The 'load' and 'store' instructions are specifically crafted to fully resolve to
an element of a memref. These instructions take as arguments n+1 indices for an
ability to index into the same memref in other ways (something which C arrays
allow for example). Furthermore, in an affine construct, the compiler can follow
use-def chains (e.g. through
-[affine.apply instructions](Dialects/Affine.md#'affine.apply'-operation)) to
+[affine.apply instructions](Dialects/Affine.md#affineapply-operation)) to
precisely analyze references at compile-time using polyhedral techniques. This
is possible because of the
[restrictions on dimensions and symbols](Dialects/Affine.md#restrictions-on-dimensions-and-symbols).
[an extension](#mlfunction-extensions-for-"escaping-scalars") to allow that is
described later in this doc.
-### Symbols and types {#symbols-and-types}
+### Symbols and types
The current MLIR disallows use of symbols in types. For example, when a tensor
or memref dimension is statically unknown, it is denoted in the type as '?'. An
simplifies the design --- types remain immutable when the values of symbols
change.
-### Block Arguments vs PHI nodes {#block-arguments-vs-phi-nodes}
+### Block Arguments vs PHI nodes
MLIR Functions represent SSA using "[block arguments](LangRef.md#blocks)" rather
than [PHI instructions](http://llvm.org/docs/LangRef.html#i-phi) used in LLVM.
interest
[starts here](https://www.google.com/url?q=https://youtu.be/Ntj8ab-5cvE?t%3D596&sa=D&ust=1529450150971000&usg=AFQjCNFQHEWL7m8q3eO-1DiKw9zqC2v24Q).
-### Signless types {#signless-types}
+### Signless types
Integers in the MLIR type system have a bitwidth (note that the `int` type is a
symbolic width equal to the machine word size), but they do not have an
[talk on youtube](https://www.youtube.com/watch?v=VeRaLPupGks) talking about
LLVM 2.0.
-### Index type disallowed in vector/tensor/memref types {#index-type-disallowed-in-aggregate-types}
+### Index type disallowed in vector/tensor/memref types
Index types are not allowed as elements of `vector`, `tensor` or `memref` type.
Index types are intended to be used for platform-specific "size" values and may
of supporting smaller integer types, e.g. `i8` or `i16`, for small indices
instead of (presumably larger) `index` type.
-### Bit width of a non-primitive types and `index` is undefined {#bit-width-of-a-compound-type}
+### Bit width of a non-primitive types and `index` is undefined
The bit width of a compound type is not defined by MLIR, it may be defined by a
specific lowering pass. In MLIR, bit width is a property of certain primitive
The bit width is not defined for dialect-specific types at MLIR level. Dialects
are free to define their own quantities for type sizes.
-### Splitting floating point vs integer operations {#splitting-floating-point-vs-integer-operations}
+### Splitting floating point vs integer operations
The MLIR operation set is likely to
[follow LLVM](http://llvm.org/docs/LangRef.html#binary-operations) and split
MLIR, but since we have experience and know the right way to do this, we'd
rather design it in from the beginning.
-### Specifying sign in integer comparison operations {#sign-in-cmpi}
+### Specifying sign in integer comparison operations
Since integers are [signless](#signless-types), it is necessary to define the
sign for integer comparison operations. This sign indicates how to treat the
whose bit representations would differ while the values are interpreted as
equal.
-### Specifying comparison kind as attribute {#cmpi-predicate}
+### Specifying comparison kind as attribute
Unlike arithmetic, comparison operators share several common properties, e.g.
they cannot be considered associative. In practice, comparisons are sometimes
impossible to implement switching logic based on the comparison kind and made
attribute validity checks (one out of ten possible kinds) more complex.
-### 'select' operation to implement min/max {#select-operation}
+### 'select' operation to implement min/max
Although `min` and `max` operations are likely to occur as a result of
transforming affine loops in ML functions, we did not make them first-class
where min/max, and thus `select`, are likely to appear. In addition, simpler
control flow may be beneficial for optimization in general.
-### Quantized integer operations {#quantized-integer-operations}
+### Quantized integer operations
We haven't designed integer quantized operations in MLIR, but experience from
TensorFlow suggests that it is better to put information about the quantization
experience. When we do, we should chat with benoitjacob@ and
[read the paper](https://arxiv.org/abs/1712.05877).
-### Dialect type extensions {#dialect-type-extensions}
+### Dialect type extensions
This section describes the design decisions that shaped the dialect extensible
type system present in MLIR.
-#### Reserving dialect type kinds {#reserving-type-kinds}
+#### Reserving dialect type kinds
Dialects that wish to define type extensions must reserve a range of type kinds
within a '.def' file within the core IR library. This means that every dialect
wishing to define custom types must modify this file, but it guarantees that all
type casting checkings are performed in O(1) time.
-#### Interactions between dialects {#type-interactions-between-dialects}
+#### Interactions between dialects
There are two different interactions between dialects that are important to
understand. When types of a dialect are:
invariants, e.g. if the standard tensor type can contain a specific type
of that dialect.
-#### Separating builtin and standard types {#separating-builtin-and-standard-types}
+#### Separating builtin and standard types
Following the separation between the built-in and standard dialect, it makes
sense to separate built-in types and standard dialect types. Built-in types are
Integer, float, vector, memref and tensor types, while important, are not
necessary for IR validity.
-#### Unregistered types {#unregistered-types}
+#### Unregistered types
MLIR supports unregistered operations in generic assembly form. MLIR also
supports a similar concept for types. When parsing, if the dialect for dialect
This representation was chosen for several reasons:
-##### Dialects must provide custom type parsers {#dialect-type-custom-parser}
+##### Dialects must provide custom type parsers
Dialect type parsing cannot plug into the existing parser infrastructure as
operations do with the OpAsmParser/Printer. Operations have a defined syntax
dialect wishes to assign a canonical name to a type, it can be done via
[type aliases](LangRef.md#type-aliases).
-### Tuple types {#tuple-type}
+### Tuple types
The MLIR type system provides first class support for defining
[tuple types](LangRef.md#tuple-type). This is due to the fact that `Tuple`
operation to derive a more concise form, thus better facilitating the
comprehension of the IR.
-## Examples {#examples}
+## Examples
This section describes a few very simple examples that help understand how MLIR
represents computation.
-### Non-affine control flow {#non-affine-control-flow}
+### Non-affine control flow
```mlir {.mlir}
// A simple linear search in every row of a matrix
`%i` loop could be parallelized: such function access analysis is calling
context sensitive.
-### Non-affine loop bounds {#non-affine-loop-bounds}
+### Non-affine loop bounds
Loop bounds that are not affine lead to a nesting of functions as shown below.
}
```
-### Reference 2D Convolution {#reference-2d-convolution}
+### Reference 2D Convolution
The following example illustrates a reference implementation of a 2D
convolution, which uses an integer set `#domain` to represent valid input data
TODO (Add more examples showing the IR for a variety of interesting cases)
-## Design alternatives and extensions {#design-alternatives-and-extensions}
+## Design alternatives and extensions
This is a list of some design alternatives and extensions that we discussed in
detail but did not include in the spec or postponed them for future
implementation experience and learn more about the challenges and limitations of
our current design in practice.
-### Polyhedral code representation alternatives: schedule lists vs schedules trees vs affine loop/if forms {#mlfunction-representation-alternatives-polyhedral-schedule-lists-vs-polyhedral-schedules-trees-vs-affine-loop-if-forms}
+### Polyhedral code representation alternatives: schedule lists vs schedules trees vs affine loop/if forms
The current MLIR uses a representation of polyhedral schedules using a tree of
if/for loops. We extensively debated the tradeoffs involved in the typical
AffineLoopTreeFunction, Polyhedral Schedule Tree function, and external
functions.
-#### Schedule Tree Representation for MLFunctions {#schedule-tree-representation-for-mlfunctions}
+#### Schedule Tree Representation for MLFunctions
This representation is based on a simplified form of the domain/schedule
representation used by the polyhedral compiler community. Domains represent what
```
-### Affine Relations {#affine-relations}
+### Affine Relations
The current MLIR spec includes affine maps and integer sets, but not affine
relations. Affine relations are a natural way to model read and write access
}
```
-### Read/Write/May_Read/May_Write sets for External Functions {#read-write-may_read-may_write-sets-for-external-functions}
+### Read/Write/May_Read/May_Write sets for External Functions
Having read, write, may_read, and may_write sets for external functions which
include opaque ones, high-performance vendor libraries such as CuDNN, CuB, MKL,
```
-### Memref Extensions {#memref-extensions}
+### Memref Extensions
1. Arbitrary polyhedral shapes for tensors: e.g., triangular shapes in tensor
dimensions where there is symmetry: use integer set (affine constraints) to
representation. 2(b) requires no change, but impacts how cost models look at
index and layout maps.
-### `affine.if` and `affine.for` Extensions for "Escaping Scalars" {#extensions-for-"escaping-scalars"}
+### `affine.if` and `affine.for` Extensions for "Escaping Scalars"
We considered providing a representation for SSA values that are live out of
`if/else` conditional bodies and loop carried in `affine.for` loops. We
The abandoned design of supporting escaping scalars is as follows:
-#### For Instruction {#for-instruction}
+#### For Instruction
Syntax:
}
```
-#### If/else Instruction {#if-else-instruction}
+#### If/else Instruction
Syntax:
in MLIR. If your transformation involves pattern matching operation DAGs, this
is a great place to start.
-## Pass Types {#pass-types}
+## Pass Types
MLIR provides different pass classes for several different granularities of
transformation. Depending on the granularity of the transformation being
-performed, a pass may derive from [FunctionPass](#function-pass) or
-[ModulePass](#module-pass); with each requiring a different set of constraints.
+performed, a pass may derive from [FunctionPass](#functionpass) or
+[ModulePass](#modulepass); with each requiring a different set of constraints.
-### FunctionPass {#function-pass}
+### FunctionPass
A function pass operates on a per-function granularity, executing on
non-external functions within a module in no particular order. Function passes
"flag-name-to-invoke-pass-via-mlir-opt", "Pass description here");
```
-### ModulePass {#module-pass}
+### ModulePass
A module pass operates on a per-module granularity, executing on the entire
program as a unit. As such, module passes are able to add/remove/modify any
}
```
-## Pass Manager {#pass-manager}
+## Pass Manager
Above we introduced the different types of passes and their constraints. Now
that we have our pass we need to be able to run it over a specific module. This
MyModulePass2
```
-## Pass Registration {pass-registration}
+## Pass Registration
Briefly shown in the example definitions of the various
[pass types](#pass-types) is the `PassRegistration` class. This is a utility to
projects / platform / upstream / llvm.git / commitdiff