From: River Riddle Date: Fri, 5 Apr 2019 15:19:42 +0000 (-0700) Subject: Tidy up the links in the documents and fix any broken ones. X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=465ef55088ee43cecf77e54159aedba78c377b36;p=platform%2Fupstream%2Fllvm.git Tidy up the links in the documents and fix any broken ones. -- PiperOrigin-RevId: 242127863 --- diff --git a/mlir/g3doc/ConversionToLLVMDialect.md b/mlir/g3doc/ConversionToLLVMDialect.md index 7f3f7c4..0ddf842 100644 --- a/mlir/g3doc/ConversionToLLVMDialect.md +++ b/mlir/g3doc/ConversionToLLVMDialect.md @@ -81,7 +81,7 @@ memref<1x? x vector<4xf32>> !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 diff --git a/mlir/g3doc/Dialects/Affine.md b/mlir/g3doc/Dialects/Affine.md index 8975507..fa51be4 100644 --- a/mlir/g3doc/Dialects/Affine.md +++ b/mlir/g3doc/Dialects/Affine.md @@ -4,24 +4,24 @@ This dialect provides a powerful abstraction for affine operations and analyses. [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: @@ -47,7 +47,7 @@ Example: %2 = affine.apply (i)[s0] -> (i+s0) (%42)[%n] ``` -#### 'affine.for' operation {#'affine.for'-operation} +#### 'affine.for' operation Syntax: @@ -62,9 +62,9 @@ shorthand-bound ::= ssa-id | `-`? integer-literal 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 @@ -108,7 +108,7 @@ func @simple_example(%A: memref, %B: memref) { } ``` -#### 'affine.if' operation {#'affine.if'-operation} +#### 'affine.if' operation Syntax: @@ -127,12 +127,12 @@ and the SSA values bound to the dimensions and symbols in the integer set. The [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: @@ -154,7 +154,7 @@ func @reduced_domain_example(%A, %X, %N) : (memref<10xi32>, i32, i32) { } ``` -#### `affine.terminator` operation {#'affine.terminator'-operation} +#### `affine.terminator` operation Syntax: @@ -163,11 +163,11 @@ operation ::= `"affine.terminator"() : () -> ()` ``` 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. diff --git a/mlir/g3doc/Dialects/LLVM.md b/mlir/g3doc/Dialects/LLVM.md index ca80bbd..0274ee2 100644 --- a/mlir/g3doc/Dialects/LLVM.md +++ b/mlir/g3doc/Dialects/LLVM.md @@ -20,7 +20,7 @@ obtained from the dialect object using `.getLLVMContext()` and `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 @@ -45,7 +45,7 @@ associated with the LLVM context of the dialect. In this document, we use the 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.`". @@ -250,12 +250,12 @@ Bitwise reinterpretation: `bitcast `. Selection: `select , , `. -### 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. @@ -283,7 +283,7 @@ Examples: %3 = llvm.constant(splat, 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 diff --git a/mlir/g3doc/Dialects/Vector.md b/mlir/g3doc/Dialects/Vector.md index 6d62720..30a79d7 100644 --- a/mlir/g3doc/Dialects/Vector.md +++ b/mlir/g3doc/Dialects/Vector.md @@ -4,11 +4,11 @@ This dialect provides mid-level abstraction for the MLIR super-vectorizer. [TOC] -## Operations {#operations} +## Operations -### Vector transfers {#vector-transfers} +### Vector transfers -#### `vector.transfer_read` operation {#'vector.transfer_read'-operation} +#### `vector.transfer_read` operation Syntax: @@ -117,7 +117,7 @@ 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`. -#### `vector.transfer_write` operation {#'vector.transfer_write'-operation} +#### `vector.transfer_write` operation Syntax: @@ -156,9 +156,9 @@ 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. -### Vector views {#vector-views} +### Vector views -#### `vector.type_cast` operation {#'vector.type_cast'-operation} +#### `vector.type_cast` operation Syntax: diff --git a/mlir/g3doc/LangRef.md b/mlir/g3doc/LangRef.md index 497595b..38fae21 100644 --- a/mlir/g3doc/LangRef.md +++ b/mlir/g3doc/LangRef.md @@ -25,11 +25,11 @@ document describes the human-readable textual form. [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). @@ -115,7 +115,7 @@ func @multiply(%A: memref<100x?xf32>, %B: memref) } ``` -## 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. @@ -145,7 +145,7 @@ Code examples are presented in blue boxes. example ::= `b` (`an` | `om`)* `a` ``` -### Common syntax {#common-syntax} +### Common syntax The following core grammar productions are used in this document: @@ -166,7 +166,7 @@ string-literal ::= `"` [^"\n\f\v\r]* `"` TODO define escaping rules 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: @@ -203,16 +203,16 @@ The scope of SSA values is defined based on the standard definition of 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: @@ -255,7 +255,7 @@ Example: %x = alloc()[%N] : memref<40x50xf32, #affine_map2to3> ``` -### Affine Expressions {#affine-expressions} +### Affine Expressions Syntax: @@ -306,7 +306,7 @@ examples, $$(i+j+1, j)$$, $$(i \mod 2, j+i)$$, $$(j, i/4, i \mod 4)$$, $$(2i+1, 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: @@ -333,7 +333,7 @@ dimension indices and symbols into a list of results, with affine expressions 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)). @@ -341,7 +341,7 @@ Affine maps are used for various core structures in MLIR. The restrictions we 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: @@ -376,7 +376,7 @@ Examples: 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. @@ -426,7 +426,7 @@ module-header-def ::= semi-affine-map-def 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 @@ -487,7 +487,7 @@ affine.if #set42(%i, %j)[%M, %N] { 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 @@ -525,7 +525,7 @@ ssa-use-and-type ::= ssa-use `:` type 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 @@ -548,11 +548,11 @@ Example: "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: @@ -565,17 +565,17 @@ function-type ::= type-list-parens `->` function-result-type ``` 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: @@ -588,12 +588,12 @@ The `index` type is a signless integer whose size is equal to the natural 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: @@ -613,7 +613,7 @@ bit one). 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: @@ -625,7 +625,7 @@ float-type ::= `f16` | `bf16` | `f32` | `f64` MLIR supports float types of certain widths that are widely used as indicated above. -#### Vector Type {#vector-type} +#### Vector Type Syntax: @@ -647,7 +647,7 @@ Note: hexadecimal integer literals are not allowed in vector type declarations, `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: @@ -671,7 +671,7 @@ you cannot control layout or get a pointer to the data. For low level buffer 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 @@ -705,7 +705,7 @@ tensor<0 x 42 x f32> tensor<0xf32> ``` -#### Memref Type {#memref-type} +#### Memref Type Syntax: @@ -772,12 +772,12 @@ Symbol capture example: %A = alloc (%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. @@ -793,7 +793,7 @@ Examples %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 @@ -825,7 +825,7 @@ Index map examples: 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 @@ -842,7 +842,7 @@ Layout map examples: #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 @@ -861,7 +861,7 @@ access pattern analysis, and for performance optimizations like vectorization, 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: @@ -880,7 +880,7 @@ complex complex ``` -#### Tuple Type {#tuple-type} +#### Tuple Type Syntax: @@ -893,7 +893,7 @@ each element may be of a different type. **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: @@ -908,7 +908,7 @@ tuple tuple, i5> ``` -## Attributes {#attributes} +## Attributes Syntax: @@ -923,11 +923,11 @@ dependent-attribute-name ::= (letter|[_]) (letter|digit|[_$])* 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, @@ -954,7 +954,7 @@ specific and dependent attributes. This is because an operation represents a 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: @@ -971,7 +971,7 @@ attribute-value ::= affine-map-attribute | type-attribute ``` -#### AffineMap Attribute {#affine-map-attribute} +#### AffineMap Attribute Syntax: @@ -981,7 +981,7 @@ affine-map-attribute ::= affine-map An affine-map attribute is an attribute that represents a affine-map object. -#### Array Attribute {#array-attribute} +#### Array Attribute Syntax: @@ -992,7 +992,7 @@ array-attribute ::= `[` (attribute-value (`,` attribute-value)*)? `]` An array attribute is an attribute that represents a collection of attribute values. -#### Boolean Attribute {#bool-attribute} +#### Boolean Attribute Syntax: @@ -1003,7 +1003,7 @@ bool-attribute ::= bool-literal A boolean attribute is a literal attribute that represents a one-bit boolean value, true or false. -#### Elements Attributes {#elements-attributes} +#### Elements Attributes Syntax: @@ -1017,7 +1017,7 @@ elements-attribute ::= dense-elements-attribute 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: @@ -1031,7 +1031,7 @@ constant vector or tensor value has been packed to the element bitwidth. The 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: @@ -1048,7 +1048,7 @@ it. Note: The parsed string literal must be in hexadecimal form. -##### Sparse Elements Attribute {#sparse-elements-attribute} +##### Sparse Elements Attribute Syntax: @@ -1078,7 +1078,7 @@ Example: /// [0, 0, 0, 0]] ``` -##### Splat Elements Attribute {#splat-elements-attribute} +##### Splat Elements Attribute Syntax: @@ -1090,7 +1090,7 @@ splat-elements-attribute ::= `splat` `<` ( tensor-type | vector-type ) `,` 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: @@ -1102,7 +1102,7 @@ An integer attribute is a literal attribute that represents an integral value of 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: @@ -1112,7 +1112,7 @@ integer-set-attribute ::= affine-map An integer-set attribute is an attribute that represents a integer-set object. -#### Float Attribute {#float-attribute} +#### Float Attribute Syntax: @@ -1123,7 +1123,7 @@ float-attribute ::= float-literal (`:` float-type)? 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: @@ -1134,7 +1134,7 @@ function-attribute ::= function-id `:` function-type A function attribute is a literal attribute that represents a reference to the given function object. -#### String Attribute {#string-attribute} +#### String Attribute Syntax: @@ -1144,7 +1144,7 @@ string-attribute ::= string-literal A string attribute is an attribute that represents a string literal value. -#### Type Attribute {#type-attribute} +#### Type Attribute Syntax: @@ -1154,7 +1154,7 @@ type-attribute ::= type A type attribute is an attribute that represents a [type object](#type-system). -## Module {#module} +## Module ``` {.ebnf} module ::= module-header-def* function* @@ -1169,7 +1169,7 @@ prepopulate a symbol table with known named types and mappings (e.g. for TPU) 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: @@ -1206,7 +1206,7 @@ func @count(%x: i64) -> (i64, i64) } ``` -#### Blocks {#blocks} +#### Blocks Syntax: @@ -1265,7 +1265,7 @@ of SSA is immediately apparent, and function arguments are no longer a special case: they become arguments to the entry block [[more rationale](Rationale.md#block-arguments-vs-phi-nodes)]. -### Operations {#operations} +### Operations Syntax: @@ -1332,21 +1332,21 @@ reason about IR dumps and manipulate operations in C++, the MLIR compiler 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: @@ -1363,7 +1363,7 @@ arguments in the target block. 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: @@ -1394,7 +1394,7 @@ func @select(%a : i32, %b :i32, %flag : i1) -> i32 { } ``` -##### 'return' terminator operation {#'return'-terminator-operation} +##### 'return' terminator operation Syntax: @@ -1407,9 +1407,9 @@ produces the result values. The count and types of the operands must match the 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: @@ -1428,7 +1428,7 @@ Example: %31 = call @my_add(%0, %1) : (tensor<16xf32>, tensor<16xf32>) -> tensor<16xf32> ``` -#### 'call_indirect' operation {#'call_indirect'-operation} +#### 'call_indirect' operation Syntax: @@ -1442,7 +1442,7 @@ and merged together with block arguments. The operands and result types of the call must match the specified function type. Function values can be created with the -[`constant` operation](#'constant'-operation). +[`constant` operation](#constant-operation). Example: @@ -1451,7 +1451,7 @@ Example: : (tensor<16xf32>, tensor<16xf32>) -> tensor<16xf32> ``` -#### 'dim' operation {#'dim'-operation} +#### 'dim' operation Syntax: @@ -1460,7 +1460,7 @@ operation ::= ssa-id `=` `dim` ssa-id `,` integer-literal `:` type ``` 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. @@ -1479,7 +1479,7 @@ Examples: %y = "std.dim"(%A){index: 1} : (tensor<4 x ? x f32>) -> index ``` -#### 'reshape' operation {#'reshape'-operation} +#### 'reshape' operation Syntax: @@ -1509,7 +1509,7 @@ Example: ``` -#### 'view' operation {#'view'-operation} +#### 'view' operation Syntax: @@ -1550,9 +1550,9 @@ Example: (%s1) [%0, %n1] %A : memref<16x?xf32, #map_a, hbm> ``` -### Memory Operations {#memory-operations} +### Memory Operations -#### 'alloc' operation {#'alloc'-operation} +#### 'alloc' operation Syntax: @@ -1581,7 +1581,7 @@ Example: %B = alloc(%M, %N)[%x, %y] : memref ``` -#### 'alloc_static' operation {#'alloc_static'-operation} +#### 'alloc_static' operation Syntax: @@ -1593,7 +1593,7 @@ operation ::= 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: @@ -1604,7 +1604,7 @@ Example: The `alloc_static` operation is used to represent code after buffer allocation has been performed. -#### 'dealloc' operation {#'dealloc'-operation} +#### 'dealloc' operation Syntax: @@ -1613,8 +1613,8 @@ operation ::= `dealloc` ssa-use `:` memref-type ``` 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: @@ -1686,7 +1686,7 @@ Example: dma_wait %tag[%index], %num_elements : memref<1 x i32, (d0) -> (d0), 4> ``` -#### 'extract_element' operation {#'extract_element'-operation} +#### 'extract_element' operation Syntax: @@ -1709,7 +1709,7 @@ Examples: %5 = extract_element %ut[%1, %2] : tensor<*xi32> ``` -#### 'load' operation {#'load'-operation} +#### 'load' operation Syntax: @@ -1726,10 +1726,9 @@ identifier). 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: @@ -1746,13 +1745,13 @@ 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: @@ -1767,11 +1766,11 @@ provided within brackets need to match the rank of the memref. 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: @@ -1782,13 +1781,13 @@ store %100, %A[%1, 1023] : memref<4x?xf32, #layout, hbm> **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: @@ -1807,7 +1806,7 @@ Example: %12 = tensor_load %10 : memref<4x?xf32, #layout, hbm> ``` -#### 'tensor_store' operation {#'tensor_store'-operation} +#### 'tensor_store' operation Syntax: @@ -1827,7 +1826,7 @@ Example: 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. @@ -1836,7 +1835,7 @@ TODO: "sub" etc. Let's not get excited about filling this out yet, we can define 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: @@ -1856,7 +1855,7 @@ required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers. It has no standard attributes. -#### 'addf' operation {#'addf'-operation} +#### 'addf' operation Examples: @@ -1881,7 +1880,7 @@ TODO: In the distant future, this will accept optional attributes for fast math, contraction, rounding mode, and other controls. -#### 'cmpi' operation {#'cmpi'-operation} +#### 'cmpi' operation Examples: @@ -1936,13 +1935,15 @@ point-related particularities, e.g., `-ffast-math` behavior, IEEE754 compliance, 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: @@ -1979,7 +1980,7 @@ anticipate the desire to multithread the compiler, and disallowing SSA values to 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`. @@ -2011,7 +2012,7 @@ is required to be the same type. This type may be an integer scalar type, a 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 / @@ -2068,7 +2069,7 @@ same element type, same mappings, same address space, and same rank, yet the 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: @@ -2093,7 +2094,7 @@ TODO: In the distant future, this will accept 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`. @@ -2125,7 +2126,7 @@ is required to be the same type. This type may be an integer scalar type, a 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`. @@ -2157,7 +2158,7 @@ is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers. It has no standard attributes. -#### 'select' operation {#'select'-operation} +#### 'select' operation Syntax: @@ -2188,10 +2189,10 @@ The operation applies to vectors and tensors elementwise given the _shape_ of 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: @@ -2220,7 +2221,7 @@ same element type, and the source and destination types may not be the same. 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 @@ -2233,7 +2234,7 @@ Currently, MLIR supports the following dialects: * [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 @@ -2262,7 +2263,7 @@ Examples: : (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. @@ -2291,7 +2292,7 @@ Example: 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 @@ -2311,7 +2312,7 @@ Example: !tf<"string"> ``` -### TensorFlow types {#tensorflow-types} +### TensorFlow types The TensorFlow dialect in MLIR defines several extended types: diff --git a/mlir/g3doc/OpDefinitions.md b/mlir/g3doc/OpDefinitions.md index c0c2614..2879f35 100644 --- a/mlir/g3doc/OpDefinitions.md +++ b/mlir/g3doc/OpDefinitions.md @@ -136,7 +136,7 @@ Operation definitions consists of: 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. @@ -162,9 +162,9 @@ Operation definitions consists of: 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. diff --git a/mlir/g3doc/Passes.md b/mlir/g3doc/Passes.md index fae225b..33ea6b4 100644 --- a/mlir/g3doc/Passes.md +++ b/mlir/g3doc/Passes.md @@ -4,7 +4,7 @@ This document describes the available MLIR passes and their contracts. [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. @@ -51,7 +51,7 @@ already present before the pass. 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. @@ -76,7 +76,7 @@ value is returned, packed into an LLVM IR struct type. Function calls and 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 @@ -148,17 +148,17 @@ func @loop_nest_tiled() -> memref<256x1024xf32> { 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 @@ -170,7 +170,7 @@ achieved at the expense of redundant computation through a cost model that 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. @@ -184,7 +184,7 @@ test/Transforms/memref-bound-check.mlir:19:13: error: 'load' op memref out of lo ^ ``` -## 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. @@ -230,7 +230,7 @@ func @store_load_affine_apply() -> memref<10x10xf32> { } ``` -## 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 @@ -245,7 +245,7 @@ test/Transforms/memref-dataflow-opt.mlir:232:7: note: dependence from 2 to 1 at 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 diff --git a/mlir/g3doc/Quantization.md b/mlir/g3doc/Quantization.md index 35e2783..eeb202f 100644 --- a/mlir/g3doc/Quantization.md +++ b/mlir/g3doc/Quantization.md @@ -26,7 +26,7 @@ Further, the scheme can be applied: * *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*. @@ -51,7 +51,7 @@ This makes it convenient to represent scaled values on a computer as signed 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). @@ -99,7 +99,7 @@ scope of this document, and it is safe to assume unless otherwise stated that 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. @@ -127,7 +127,7 @@ point values is indicative of common types on typical hardware but is not 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 @@ -144,7 +144,7 @@ $$ 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. @@ -153,7 +153,7 @@ $$ 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. @@ -162,7 +162,7 @@ $$ 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: @@ -176,11 +176,11 @@ 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) @@ -188,15 +188,16 @@ MLIR: * 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 @@ -208,22 +209,22 @@ Not every application of quantization will use all facilities. Specifically, the 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 @@ -234,7 +235,7 @@ 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) @@ -260,7 +261,7 @@ This scheme also naturally allows computations that are *partially quantized* 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 @@ -278,9 +279,9 @@ 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 @@ -310,7 +311,7 @@ TODO: This op set is still evolving and needs to be completed. * 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. @@ -318,7 +319,7 @@ 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 diff --git a/mlir/g3doc/Rationale.md b/mlir/g3doc/Rationale.md index 42b32cf9..de9b68c 100644 --- a/mlir/g3doc/Rationale.md +++ b/mlir/g3doc/Rationale.md @@ -8,7 +8,7 @@ about their consistency or readability. [TOC] -## Abstract {#abstract} +## Abstract MLIR is a compiler intermediate representation with similarities to traditional three-address SSA representations (like @@ -26,7 +26,7 @@ MLIR stands for one of "Multi-Level IR" or "Multi-dimensional Loop IR" or 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 @@ -101,12 +101,12 @@ memory hierarchy. * 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 @@ -114,7 +114,7 @@ n-ranked tensor. This disallows the equivalent of pointer arithmetic or the 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). @@ -126,7 +126,7 @@ have an SSA representation -- [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 @@ -166,7 +166,7 @@ type - memref<8x%Nxf32>. We went for the current approach in MLIR because it 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. @@ -203,7 +203,7 @@ and described in 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 @@ -231,7 +231,7 @@ More information about this split is available in an old [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 @@ -253,7 +253,7 @@ fixed-width integer types, at the SSA value level. It has an additional benefit 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 @@ -273,7 +273,7 @@ introduced. 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 @@ -299,7 +299,7 @@ We are a long way from this sort of thing being a priority to care about in 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 @@ -313,7 +313,7 @@ and cannot be padded by this operation, it is impossible to compare two values 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 @@ -337,7 +337,7 @@ possible to store the predicate as string attribute, it would have rendered 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 @@ -364,7 +364,7 @@ However, this control flow granularity is not available in the ML functions 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 @@ -390,19 +390,19 @@ to re-examine the representation for quantized arithmetic when we have that 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: @@ -427,7 +427,7 @@ 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 @@ -436,7 +436,7 @@ appears in function signatures and generic assembly forms of operations). 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 @@ -455,7 +455,7 @@ identify and parse them. 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 @@ -485,7 +485,7 @@ to think of these types as existing within the namespace of the dialect. If a 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` @@ -519,12 +519,12 @@ specific needs. The custom assembly form can de-duplicate information from the 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 @@ -585,7 +585,7 @@ inside the called function (`@search_body`) is necessary to determine if the `%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. @@ -627,7 +627,7 @@ func @inner_nest2(%m, %n) -> i32 { } ``` -### 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 @@ -697,7 +697,7 @@ func @conv2d(memref<16x1024x1024x3xf32, #lm0, vmem> %input, 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 @@ -705,7 +705,7 @@ consideration on demand. We will revisit these discussions when we have more 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 @@ -738,7 +738,7 @@ At a high level, we have two alternatives here: 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 @@ -791,7 +791,7 @@ func @matmul(%A, %B, %C, %M, %N, %K) : (...) { // %M, N, K are symbols ``` -### 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 @@ -853,7 +853,7 @@ bb0 (%0, %1: memref<128xf32>, i64): } ``` -### 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, @@ -890,7 +890,7 @@ func @dma_hbm_to_vmem(memref<1024 x f32, #layout_map0, hbm> %a, ``` -### 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 @@ -911,7 +911,7 @@ func @dma_hbm_to_vmem(memref<1024 x f32, #layout_map0, hbm> %a, 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 @@ -925,7 +925,7 @@ arguments and a yield like terminator in for/if instructions. The abandoned design of supporting escaping scalars is as follows: -#### For Instruction {#for-instruction} +#### For Instruction Syntax: @@ -955,7 +955,7 @@ func int32 @sum(%A : memref, %N : i32) -> (i32) { } ``` -#### If/else Instruction {#if-else-instruction} +#### If/else Instruction Syntax: diff --git a/mlir/g3doc/WritingAPass.md b/mlir/g3doc/WritingAPass.md index eac8bc2..7202153 100644 --- a/mlir/g3doc/WritingAPass.md +++ b/mlir/g3doc/WritingAPass.md @@ -13,14 +13,14 @@ See [MLIR Rewrites](QuickstartRewrites.md) for a quick start on graph rewriting 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 @@ -70,7 +70,7 @@ static PassRegistration pass( "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 @@ -231,7 +231,7 @@ void MyPass::runOn*() { } ``` -## 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 @@ -286,7 +286,7 @@ Function Pipeline 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