Examples:
-```mlir {.mlir}
+```mlir
memref<f32> -> !llvm.type<"{ float*, float*, i64 }">
memref<1 x f32> -> !llvm.type<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
memref<? x f32> -> !llvm.type<"{ float*, float*, i64, [1 x i64], [1 x i64] }">
Examples:
-```mlir {.mlir}
+```mlir
// unranked descriptor
memref<*xf32> -> !llvm.type<"{i64, i8*}">
```
Examples:
-```mlir {.mlir}
+```mlir
// zero-ary function type with no results.
() -> ()
// is converted to a zero-ary function with `void` result
Examples:
-```mlir {.mlir}
+```mlir
// zero-ary function type with no results.
func @foo() -> ()
// gets LLVM type void().
Example:
-```mlir {.mlir}
+```mlir
func @foo(%arg0: i32, %arg1: i64) -> (i32, i64) {
return %arg0, %arg1 : i32, i64
}
Example:
-```mlir {.mlir}
+```mlir
cond_br %0, ^bb1(%1 : i32), ^bb1(%2 : i32)
^bb1(%3 : i32)
"use"(%3) : (i32) -> ()
leads to a new basic block being inserted,
-```mlir {.mlir}
+```mlir
cond_br %0, ^bb1(%1 : i32), ^dummy
^bb1(%3 : i32):
"use"(%3) : (i32) -> ()
before the conversion to the LLVM IR dialect:
-```mlir {.mlir}
+```mlir
llvm.cond_br %0, ^bb1(%1 : !llvm.type<"i32">), ^dummy
^bb1(%3 : !llvm.type<"i32">):
"use"(%3) : (!llvm.type<"i32">) -> ()
An access to a zero-dimensional memref is converted into a plain load:
-```mlir {.mlir}
+```mlir
// before
%0 = load %m[] : memref<f32>
An access to a memref with indices:
-```mlir {.mlir}
+```mlir
%0 = load %m[1,2,3,4] : memref<10x?x13x?xf32>
```
is transformed into the equivalent of the following code:
-```mlir {.mlir}
+```mlir
// obtain the buffer pointer
%b = llvm.extractvalue %m[0] : !llvm.type<"{float*, i64, i64}">
### CallSite Location
-``` {.ebnf}
+```
callsite-location ::= 'callsite' '(' location 'at' location ')'
```
### FileLineCol Location
-``` {.ebnf}
+```
filelinecol-location ::= string-literal ':' integer-literal ':' integer-literal
```
### Fused Location
-``` {.ebnf}
+```
fused-location ::= `fused` fusion-metadata? '[' location (location ',')* ']'
fusion-metadata ::= '<' attribute-value '>'
```
### Name Location
-``` {.ebnf}
+```
name-location ::= string-literal ('(' location ')')?
```
### Unknown Location
-``` {.ebnf}
+```
unknown-location ::= `unknown`
```
A few examples are shown below:
-```mlir {.mlir}
+```mlir
// Expect an error on the same line.
func @bad_branch() {
br ^missing // expected-error {{reference to an undefined block}}
Examples:
-```mlir {.mlir}
+```mlir
// A 2d to 3d affine mapping.
// d0/d1 are dimensions, s0 is a symbol
#affine_map2to3 = (d0, d1)[s0] -> (d0, d1 + s0, d1 - s0)
Syntax:
-``` {.ebnf}
+```
// Uses of SSA values that are passed to dimensional identifiers.
dim-use-list ::= `(` ssa-use-list? `)`
Example:
-```mlir {.mlir}
+```mlir
#affine_map2to3 = (d0, d1)[s0] -> (d0, d1 + s0, d1 - s0)
// Binds %N to the s0 symbol in affine_map2to3.
%x = alloc()[%N] : memref<40x50xf32, #affine_map2to3>
Syntax:
-``` {.ebnf}
+```
affine-expr ::= `(` affine-expr `)`
| affine-expr `+` affine-expr
| affine-expr `-` affine-expr
Syntax:
-``` {.ebnf}
+```
affine-map-inline
::= dim-and-symbol-id-lists `->` multi-dim-affine-expr
```
Syntax:
-``` {.ebnf}
+```
affine-map-id ::= `#` suffix-id
// Definitions of affine maps are at the top of the file.
Examples:
-```mlir {.mlir}
+```mlir
// Affine map out-of-line definition and usage example.
#affine_map42 = (d0, d1)[s0] -> (d0, d0 + d1 + s0 floordiv 2)
Syntax of semi-affine expressions:
-``` {.ebnf}
+```
semi-affine-expr ::= `(` semi-affine-expr `)`
| semi-affine-expr `+` semi-affine-expr
| semi-affine-expr `-` semi-affine-expr
Syntax of semi-affine maps:
-``` {.ebnf}
+```
semi-affine-map-inline
::= dim-and-symbol-id-lists `->` multi-dim-semi-affine-expr
```
the top of the file and given a name with a semi-affine map definition, and used
by name.
-``` {.ebnf}
+```
semi-affine-map-id ::= `#` suffix-id
// Definitions of semi-affine maps are at the top of file.
Syntax of affine constraints:
-``` {.ebnf}
+```
affine-constraint ::= affine-expr `>=` `0`
| affine-expr `==` `0`
affine-constraint-conjunction ::= affine-constraint (`,` affine-constraint)*
top of the file and given a name with an integer set definition, and used by
name.
-``` {.ebnf}
+```
integer-set-id ::= `#` suffix-id
integer-set-inline
Example:
-```mlir {.mlir}
+```mlir
// A example two-dimensional integer set with two symbols.
#set42 = (d0, d1)[s0, s1]
: (d0 >= 0, -d0 + s0 - 1 >= 0, d1 >= 0, -d1 + s1 - 1 >= 0)
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `affine.apply` affine-map dim-and-symbol-use-list
```
Example:
-```mlir {.mlir}
+```mlir
#map10 = (d0, d1) -> (d0 floordiv 8 + d1 floordiv 128)
...
%1 = affine.apply #map10 (%s, %t)
Syntax:
-``` {.ebnf}
+```
operation ::= `affine.for` ssa-id `=` lower-bound `to` upper-bound
(`step` integer-literal)? `{` op* `}`
Example showing reverse iteration of the inner loop:
-```mlir {.mlir}
+```mlir
#map57 = (d0)[s0] -> (s0 - d0 - 1)
func @simple_example(%A: memref<?x?xf32>, %B: memref<?x?xf32>) {
Syntax:
-``` {.ebnf}
+```
operation ::= `affine.if` if-op-cond `{` op* `}` (`else` `{` op* `}`)?
if-op-cond ::= integer-set dim-and-symbol-use-list
```
Example:
-```mlir {.mlir}
+```mlir
#set = (d0, d1)[s0]: (d0 - 10 >= 0, s0 - d0 - 9 >= 0,
d1 - 10 >= 0, s0 - d1 - 9 >= 0)
func @reduced_domain_example(%A, %X, %N) : (memref<10xi32>, i32, i32) {
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `affine.load` ssa-use `[` multi-dim-affine-map-of-ssa-ids `]` `:` memref-type
```
-The `affine.load` op reads an element from a memref, where the index
-for each memref dimension is an affine expression of loop induction
-variables and symbols. The output of 'affine.load' is a new value with the
-same type as the elements of the memref. An affine expression of loop IVs
-and symbols must be specified for each dimension of the memref. The keyword
-'symbol' can be used to indicate SSA identifiers which are symbolic.
+
+The `affine.load` op reads an element from a memref, where the index for each
+memref dimension is an affine expression of loop induction variables and
+symbols. The output of 'affine.load' is a new value with the same type as the
+elements of the memref. An affine expression of loop IVs and symbols must be
+specified for each dimension of the memref. The keyword 'symbol' can be used to
+indicate SSA identifiers which are symbolic.
Example:
-```mlir {.mlir}
+```mlir
Example 1:
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `affine.store` ssa-use, ssa-use `[` multi-dim-affine-map-of-ssa-ids `]` `:` memref-type
```
-The `affine.store` op writes an element to a memref, where the index
-for each memref dimension is an affine expression of loop induction
-variables and symbols. The 'affine.store' op stores a new value which is the
-same type as the elements of the memref. An affine expression of loop IVs
-and symbols must be specified for each dimension of the memref. The keyword
-'symbol' can be used to indicate SSA identifiers which are symbolic.
+
+The `affine.store` op writes an element to a memref, where the index for each
+memref dimension is an affine expression of loop induction variables and
+symbols. The 'affine.store' op stores a new value which is the same type as the
+elements of the memref. An affine expression of loop IVs and symbols must be
+specified for each dimension of the memref. The keyword 'symbol' can be used to
+indicate SSA identifiers which are symbolic.
Example:
-```mlir {.mlir}
+```mlir
Example 1:
Syntax:
-``` {.ebnf}
+```
operation ::= `affine.dma_Start` ssa-use `[` multi-dim-affine-map-of-ssa-ids `]`, `[` multi-dim-affine-map-of-ssa-ids `]`, `[` multi-dim-affine-map-of-ssa-ids `]`, ssa-use `:` memref-type
```
Example:
-```mlir {.mlir}
+```mlir
For example, a DmaStartOp operation that transfers 256 elements of a memref
'%src' in memory space 0 at indices [%i + 3, %j] to memref '%dst' in memory
%stride, %num_elt_per_stride : ...
```
+
#### 'affine.dma_wait' operation
Syntax:
-``` {.ebnf}
+```
operation ::= `affine.dma_Start` ssa-use `[` multi-dim-affine-map-of-ssa-ids `]`, `[` multi-dim-affine-map-of-ssa-ids `]`, `[` multi-dim-affine-map-of-ssa-ids `]`, ssa-use `:` memref-type
```
Example:
-```mlir {.mlir}
+```mlir
affine.dma_start %src[%i, %j], %dst[%k, %l], %tag[%index], %num_elements :
memref<2048xf32, 0>, memref<256xf32, 1>, memref<1xi32, 2>
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `affine.min` affine-map dim-and-symbol-use-list
```
Example:
-```mlir {.mlir}
+```mlir
%0 = affine.min (d0)[s0] -> (1000, d0 + 512, s0) (%arg0)[%arg1]
Syntax:
-``` {.ebnf}
+```
operation ::= `"affine.terminator"() : () -> ()`
```
Example:
-```mlir {.mlir}
+```mlir
%bDimX = "gpu.block_dim"() {dimension = "x"} : () -> (index)
```
Example:
-```mlir {.mlir}
+```mlir
%bIdY = "gpu.block_id"() {dimension = "y"} : () -> (index)
```
Example:
-```mlir {.mlir}
+```mlir
%gDimZ = "gpu.grid_dim"() {dimension = "z"} : () -> (index)
```
Example:
-```mlir {.mlir}
+```mlir
module attributes {gpu.container_module} {
// This module creates a separate compilation unit for the GPU compiler.
Example:
-```mlir {.mlir}
+```mlir
%tIdX = "gpu.thread_id"() {dimension = "x"} : () -> (index)
```
Example:
-```mlir {.mlir}
+```mlir
gpu.yield %f0, %f1 : f32, f32
```
-
### `gpu.all_reduce`
The "all_reduce" op reduces the value of every work item across a local
For example, both
-```mlir {.mlir}
+```mlir
%1 = "gpu.all_reduce"(%0) ({}) { op = "add" } : (f32) -> (f32)
%2 = "gpu.all_reduce"(%0) ({
^bb(%lhs : f32, %rhs : f32):
"gpu.yield"(%sum) : (f32) -> ()
}) : (f32) -> (f32)
```
-compute the sum of each work item's %0 value. The first version specifies
-the accumulation as operation, whereas the second version specifies the
-accumulation as code region. The accumulation operation must either be
-`add` or `mul`.
+
+compute the sum of each work item's %0 value. The first version specifies the
+accumulation as operation, whereas the second version specifies the accumulation
+as code region. The accumulation operation must either be `add` or `mul`.
Either none or all work items of a workgroup need to execute this op
in convergence.
The "barrier" op synchronizes all work items of a workgroup. It is used
to coordinate communication between the work items of the workgroup.
-```mlir {.mlir}
+```mlir
gpu.barrier
```
-waits until all work items in the workgroup have reached this point
-and all memory accesses made by these work items prior to the op are
-visible to all work items in the workgroup. Data hazards between work items
-accessing the same memory can be avoided by synchronizing work items
-in-between these accesses.
+
+waits until all work items in the workgroup have reached this point and all
+memory accesses made by these work items prior to the op are visible to all work
+items in the workgroup. Data hazards between work items accessing the same
+memory can be avoided by synchronizing work items in-between these accesses.
Either none or all work items of a workgroup need to execute this op
in convergence.
The LLVM IR dialect defines a single MLIR type, `LLVM::LLVMType`, that can wrap
any existing LLVM IR type. Its syntax is as follows
-``` {.ebnf}
+```
type ::= `!llvm<"` llvm-canonical-type `">
llvm-canonical-type ::= <canonical textual representation defined by LLVM>
```
LLVM functions, such as linkage and calling convention, that may be modeled
differently by the built-in MLIR function.
-```mlir {.mlir}
+```mlir
// The type of @bar is !llvm<"i64 (i64)">
llvm.func @bar(%arg0: !llvm.i64) -> !llvm.i64 {
llvm.return %arg0 : !llvm.i64
Examples:
-```mlir {.mlir}
+```mlir
// Integer addition.
%0 = llvm.add %a, %b : !llvm.i32
Examples:
-```mlir {.mlir}
+```mlir
// Float addition.
%0 = llvm.fadd %a, %b : !llvm.float
Examples:
-```mlir {.mlir}
+```mlir
// Allocate an array of 4 floats on stack
%c4 = llvm.mlir.constant(4) : !llvm.i64
%0 = llvm.alloca %c4 x !llvm.float : (!llvm.i64) -> !llvm<"float*">
Examples:
-```mlir {.mlir}
+```mlir
// Get the value third element of the second element of a structure.
%0 = llvm.extractvalue %s[1, 2] : !llvm<"{i32, {i1, i8, i16}">
Examples:
-```mlir {.mlir}
+```mlir
// Branch without arguments.
^bb0:
llvm.br ^bb0
Examples:
-```mlir {.mlir}
+```mlir
// Direct call without arguments and with one result.
%0 = llvm.call @foo() : () -> (!llvm.float)
Examples:
-```mlir {.mlir}
+```mlir
func @foo() {
// Get the address of a global.
%0 = llvm.mlir.addressof @const : !llvm<"i32*">
Examples:
-```mlir {.mlir}
+```mlir
// Integer constant, internal i32 is mandatory
%0 = llvm.mlir.constant(42 : i32) : !llvm.i32
There are two forms of initialization syntax. Simple constants that can be
represented as MLIR attributes can be given in-line:
-```mlir {.mlir}
+```mlir
llvm.mlir.global @variable(32.0 : f32) : !llvm.float
```
More complex constants that cannot be represented as MLIR attributes can be
given in an initializer region:
-```mlir {.mlir}
+```mlir
// This global is initialized with the equivalent of:
// i32* getelementptr (i32* @g2, i32 2)
llvm.mlir.global constant @int_gep() : !llvm<"i32*"> {
Examples:
-```mlir {.mlir}
+```mlir
// Global values use @-identifiers.
llvm.mlir.global constant @cst(42 : i32) : !llvm.i32
Examples:
-```mlir {.mlir}
+```mlir
// Null pointer to i8 value.
%0 = llvm.mlir.null : !llvm<"i8*">
Example:
-```mlir {.mlir}
+```mlir
// Create a structure with a 32-bit integer followed by a float.
%0 = llvm.mlir.undef : !llvm<"{i32, float}">
```
The SPIR-V dialect reuses standard integer, float, and vector types and defines
the following dialect-specific types:
-``` {.ebnf}
+```
spirv-type ::= array-type
| pointer-type
| runtime-array-type
This corresponds to SPIR-V [array type][ArrayType]. Its syntax is
-``` {.ebnf}
+```
element-type ::= integer-type
| floating-point-type
| vector-type
This corresponds to SPIR-V [image type][ImageType]. Its syntax is
-``` {.ebnf}
+```
dim ::= `1D` | `2D` | `3D` | `Cube` | <and other SPIR-V Dim specifiers...>
depth-info ::= `NoDepth` | `IsDepth` | `DepthUnknown`
For example,
-``` {.mlir}
+```
!spv.image<f32, 1D, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Unknown>
!spv.image<f32, Cube, IsDepth, Arrayed, MultiSampled, NeedSampler, Rgba32f>
```
This corresponds to SPIR-V [pointer type][PointerType]. Its syntax is
-``` {.ebnf}
+```
storage-class ::= `UniformConstant`
| `Uniform`
| `Workgroup`
This corresponds to SPIR-V [runtime array type][RuntimeArrayType]. Its syntax is
-``` {.ebnf}
+```
runtime-array-type ::= `!spv.rtarray<` element-type `>`
```
This corresponds to SPIR-V [struct type][StructType]. Its syntax is
-``` {.ebnf}
+```
struct-member-decoration ::= integer-literal? spirv-decoration*
struct-type ::= `!spv.struct<` spirv-type (`[` struct-member-decoration `]`)?
(`, ` spirv-type (`[` struct-member-decoration `]`)?
For Example,
-``` {.mlir}
+```
!spv.struct<f32>
!spv.struct<f32 [0]>
!spv.struct<f32, !spv.image<f32, 1D, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Unknown>>
Syntax:
-``` {.ebnf}
+```
operation ::= `br` successor
successor ::= bb-id branch-use-list?
branch-use-list ::= `(` ssa-use-list `:` type-list-no-parens `)`
Syntax:
-``` {.ebnf}
+```
operation ::= `cond_br` ssa-use `,` successor `,` successor
```
The following example illustrates a function with a conditional branch operation
that targets the same block:
-```mlir {.mlir}
+```mlir
func @select(i32, i32, i1) -> i32 {
^bb0(%a : i32, %b :i32, %flag : i1) :
// Both targets are the same, operands differ
Syntax:
-``` {.ebnf}
+```
operation ::= `return` (ssa-use-list `:` type-list-no-parens)?
```
Syntax:
-``` {.ebnf}
-operation ::=
+```
+operation ::=
(ssa-id `=`)? `call` symbol-ref-id `(` ssa-use-list? `)` `:` function-type
```
Example:
-```mlir {.mlir}
+```mlir
// Calling the function my_add.
%31 = call @my_add(%0, %1) : (tensor<16xf32>, tensor<16xf32>) -> tensor<16xf32>
```
Syntax:
-``` {.ebnf}
+```
operation ::= `call_indirect` ssa-use `(` ssa-use-list? `)` `:` function-type
```
Example:
-```mlir {.mlir}
+```mlir
%31 = call_indirect %15(%0, %1)
: (tensor<16xf32>, tensor<16xf32>) -> tensor<16xf32>
```
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `dim` ssa-id `,` integer-literal `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Always returns 4, can be constant folded:
%x = dim %A, 0 : tensor<4 x ? x f32>
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `alloc` dim-and-symbol-use-list `:` memref-type
```
Example:
-```mlir {.mlir}
+```mlir
// Allocating memref for a fully static shape.
%A = alloc() : memref<1024x64xf32, #layout_map0, memspace0>
Syntax:
-``` {.ebnf}
+```
operation ::=
ssa-id `=` `alloc_static` `(` integer-literal `)` : memref-type
```
Example:
-```mlir {.mlir}
+```mlir
%A = alloc_static(0x1232a00) : memref<1024 x 64 x f32, #layout_map0, memspace0>
```
Syntax:
-``` {.ebnf}
+```
operation ::= `dealloc` ssa-use `:` memref-type
```
Example:
-```mlir {.mlir}
+```mlir
dealloc %A : memref<128 x f32, #layout, memspace0>
```
Syntax:
-``` {.ebnf}
+```
operation ::= `dma_start` ssa-use`[`ssa-use-list`]` `,`
ssa-use`[`ssa-use-list`]` `,` ssa-use `,`
ssa-use`[`ssa-use-list`]` (`,` ssa-use `,` ssa-use)?
Example:
-```mlir {.mlir}
+```mlir
%size = constant 32 : index
%tag = alloc() : memref<1 x i32, (d0) -> (d0), 4>
%idx = constant 0 : index
Example:
-```mlir {.mlir}
+```mlir
dma_wait %tag[%idx], %size : memref<1 x i32, (d0) -> (d0), 4>
```
Examples:
-```mlir {.mlir}
+```mlir
%3 = extract_element %v[%1, %2] : vector<4x4xi32>
%4 = extract_element %t[%1, %2] : tensor<4x4xi32>
%5 = extract_element %ut[%1, %2] : tensor<*xi32>
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `load` ssa-use `[` ssa-use-list `]` `:` memref-type
```
Example:
-```mlir {.mlir}
+```mlir
%1 = affine.apply (d0, d1) -> (3*d0) (%i, %j)
%2 = affine.apply (d0, d1) -> (d1+1) (%i, %j)
%12 = load %A[%1, %2] : memref<8x?xi32, #layout, memspace0>
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `splat` ssa-use `:` ( vector-type | tensor-type )
```
Example:
-```mlir {.mlir}
+```mlir
%s = load %A[%i] : memref<128xf32>
%v = splat %s : vector<4xf32>
%t = splat %s : tensor<8x16xi32>
```
TODO: This operation is easy to extend to broadcast to dynamically shaped
-tensors in the same way dynamically shaped memrefs are handled. `mlir {.mlir} //
+tensors in the same way dynamically shaped memrefs are handled. `mlir //
Broadcasts %s to a 2-d dynamically shaped tensor, with %m, %n binding // to the
sizes of the two dynamic dimensions. %m = "foo"() : () -> (index) %n = "bar"() :
() -> (index) %t = splat %s [%m, %n] : tensor<?x?xi32>`
Syntax:
-``` {.ebnf}
+```
operation ::= `store` ssa-use `,` ssa-use `[` ssa-use-list `]` `:` memref-type
```
Example:
-```mlir {.mlir}
+```mlir
store %100, %A[%1, 1023] : memref<4x?xf32, #layout, memspace0>
```
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `tensor_load` ssa-use-and-type
```
Example:
-```mlir {.mlir}
+```mlir
// Produces a value of tensor<4x?xf32> type.
%12 = tensor_load %10 : memref<4x?xf32, #layout, memspace0>
```
Syntax:
-``` {.ebnf}
+```
operation ::= `tensor_store` ssa-use `,` ssa-use `:` memref-type
```
Example:
-```mlir {.mlir}
+```mlir
%9 = dim %8, 1 : tensor<4x?xf32>
%10 = alloc(%9) : memref<4x?xf32, #layout, memspace0>
tensor_store %8, %10 : memref<4x?xf32, #layout, memspace0>
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `absf` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar absolute value.
%a = absf %b : f64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `ceilf` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar ceiling value.
%a = ceilf %b : f64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `cos` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar cosine value.
%a = cos %b : f64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `exp` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar natural exponential.
%a = exp %b : f64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `negf` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar negation value.
%a = negf %b : f64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `tanh` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar hyperbolic tangent value.
%a = tanh %b : f64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `addi` ssa-use `,` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar addition.
%a = addi %b, %c : i64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `addf` ssa-use `,` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar addition.
%a = addf %b, %c : f64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `and` ssa-use `,` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar integer bitwise and.
%a = and %b, %c : i64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `cmpi` string-literal `,` ssa-id `,` ssa-id `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Custom form of scalar "signed less than" comparison.
%x = cmpi "slt", %lhs, %rhs : i32
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `constant` attribute-value `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Integer constant
%1 = constant 42 : i32
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `copysign` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar copysign value.
%a = copysign %b %c : f64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `divis` ssa-use `,` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar signed integer division.
%a = divis %b, %c : i64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `diviu` ssa-use `,` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar unsigned integer division.
%a = diviu %b, %c : i64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `memref_cast` ssa-use `:` type `to` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Discard static dimension information.
%3 = memref_cast %2 : memref<4x?xf32> to memref<?x?xf32>
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `mulf` ssa-use `,` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar multiplication.
%a = mulf %b, %c : f64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `or` ssa-use `,` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar integer bitwise or.
%a = or %b, %c : i64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `remis` ssa-use `,` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar signed integer division remainder.
%a = remis %b, %c : i64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `remiu` ssa-use `,` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar unsigned integer division remainder.
%a = remiu %b, %c : i64
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `select` ssa-use `,` ssa-use `,` ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Custom form of scalar selection.
%x = select %cond, %true, %false : i32
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `tensor_cast` ssa-use `:` type `to` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Convert from unknown rank to rank 2 with unknown dimension sizes.
%2 = "std.tensor_cast"(%1) : (tensor<*xf32>) -> tensor<?x?xf32>
%2 = tensor_cast %1 : tensor<*xf32> to tensor<?x?xf32>
Syntax:
-``` {.ebnf}
+```
operation ::= ssa-id `=` `xor` ssa-use, ssa-use `:` type
```
Examples:
-```mlir {.mlir}
+```mlir
// Scalar integer bitwise xor.
%a = xor %b, %c : i64
Depending on the function signature on which this emitter is called, the
generated IR resembles the following, for a 4-D memref of `vector<4xi8>`:
-``` {.mlir}
+```
// CHECK-LABEL: func @t1(%lhs: memref<3x4x5x6xvector<4xi8>>, %rhs: memref<3x4x5x6xvector<4xi8>>, %result: memref<3x4x5x6xvector<4xi8>>) -> () {
// CHECK: affine.for {{.*}} = 0 to 3 {
// CHECK: affine.for {{.*}} = 0 to 4 {
or the following, for a 0-D `memref<f32>`:
-``` {.mlir}
+```
// CHECK-LABEL: func @t3(%lhs: memref<f32>, %rhs: memref<f32>, %result: memref<f32>) -> () {
// CHECK: {{.*}} = load %arg1[] : memref<f32>
// CHECK: {{.*}} = load %arg0[] : memref<f32>
Here's an example of an MLIR module:
-```mlir {.mlir}
+```mlir
// Compute A*B using an implementation of multiply kernel and print the
// result using a TensorFlow op. The dimensions of A and B are partially
// known. The shapes are assumed to match.
This is the EBNF grammar used in this document, presented in yellow boxes.
-``` {.ebnf}
+```
alternation ::= expr0 | expr1 | expr2 // Either expr0 or expr1 or expr2.
sequence ::= expr0 expr1 expr2 // Sequence of expr0 expr1 expr2.
repetition0 ::= expr* // 0 or more occurrences.
Code examples are presented in blue boxes.
-```mlir {.mlir}
+```mlir
// This is an example use of the grammar above:
// This matches things like: ba, bana, boma, banana, banoma, bomana...
example ::= `b` (`an` | `om`)* `a`
The following core grammar productions are used in this document:
-``` {.ebnf}
+```
// TODO: Clarify the split between lexing (tokens) and parsing (grammar).
digit ::= [0-9]
hex_digit ::= [0-9a-fA-F]
Syntax:
-``` {.ebnf}
+```
// Identifiers
bare-id ::= (letter|[_]) (letter|digit|[_$.])*
bare-id-list ::= bare-id (`,` bare-id)*
Example:
-```mlir {.mlir}
+```mlir
// LLVM: %x = call {i16, i1} @llvm.sadd.with.overflow.i16(i16 %a, i16 %b)
%x:2 = "llvm.sadd.with.overflow.i16"(%a, %b) : (i16, i16) -> (i16, i1)
```
Syntax:
-``` {.ebnf}
+```
operation ::= op-result-list? (generic-operation | custom-operation)
trailing-location?
generic-operation ::= string-literal '(' ssa-use-list? ')' attribute-dict?
Example:
-```mlir {.mlir}
+```mlir
// An operation that produces two results.
// The results of %result can be accessed via the <name> `#` <opNo> syntax.
%result:2 = "foo_div"() : () -> (f32, i32)
Example:
-```mlir {.mlir}
+```mlir
// Branch to ^bb1 or ^bb2 depending on the condition %cond.
// Pass value %v to ^bb2, but not to ^bb1.
"cond_br"(%cond)[^bb1, ^bb2(%v : index)] : (i1) -> ()
### Module
-``` {.ebnf}
+```
module ::= `module` symbol-ref-id? (`attributes` attribute-dict)? region
```
or attributes that establish a symbolic connection (e.g. symbols referenced by
name via a string attribute like [SymbolRefAttr](#symbol-reference-attribute)):
-``` {.ebnf}
+```
function ::= `func` function-signature function-attributes? function-body?
function-signature ::= symbol-ref-id `(` argument-list `)`
Examples:
-```mlir {.mlir}
+```mlir
// External function definitions.
func @abort()
func @scribble(i32, i64, memref<? x 128 x f32, #layout_map0>) -> f64
Syntax:
-``` {.ebnf}
+```
block ::= block-label operation+
block-label ::= block-id block-arg-list? `:`
block-id ::= caret-id
Here is a simple example function showing branches, returns, and block
arguments:
-```mlir {.mlir}
+```mlir
func @simple(i64, i1) -> i64 {
^bb0(%a: i64, %cond: i1): // Code dominated by ^bb0 may refer to %a
cond_br %cond, ^bb1, ^bb2
The first block in the region cannot be a successor of any other block. The
syntax for the region is as follows:
-``` {.ebnf}
+```
region ::= `{` block* `}`
```
Example:
-```mlir {.mlir}
+```mlir
func @accelerator_compute(i64, i1) -> i64 {
^bb0(%a: i64, %cond: i1): // Code dominated by ^bb0 may refer to %a
cond_br %cond, ^bb1, ^bb2
may have application-specific semantics. For example, MLIR supports a set of
[dialect types](#dialect-types).
-``` {.ebnf}
+```
type ::= type-alias | dialect-type | standard-type
type-list-no-parens ::= type (`,` type)*
### Type Aliases
-``` {.ebnf}
+```
type-alias-def ::= '!' alias-name '=' 'type' type
type-alias ::= '!' alias-name
```
Example:
-```mlir {.mlir}
+```mlir
!avx_m128 = type vector<4 x f32>
// Using the original type.
Similarly to operations, dialects may define custom extensions to the type
system.
-``` {.ebnf}
+```
dialect-namespace ::= bare-id
opaque-dialect-item ::= dialect-namespace '<' string-literal '>'
Dialect types can be specified in a verbose form, e.g. like this:
-```mlir {.mlir}
+```mlir
// LLVM type that wraps around llvm IR types.
!llvm<"i32*">
Dialect types that are simple enough can use the pretty format, which is a
lighter weight syntax that is equivalent to the above forms:
-```mlir {.mlir}
+```mlir
// Tensor flow string type.
!tf.string
Standard types are a core set of [dialect types](#dialect-types) that are
defined in a builtin dialect and thus available to all users of MLIR.
-``` {.ebnf}
+```
standard-type ::= complex-type
| float-type
| function-type
Syntax:
-``` {.ebnf}
+```
complex-type ::= `complex` `<` type `>`
```
Examples:
-```mlir {.mlir}
+```mlir
complex<f32>
complex<i32>
```
Syntax:
-``` {.ebnf}
+```
// Floating point.
float-type ::= `f16` | `bf16` | `f32` | `f64`
```
Syntax:
-``` {.ebnf}
+```
// MLIR functions can return multiple values.
function-result-type ::= type-list-parens
| non-function-type
Syntax:
-``` {.ebnf}
+```
// Target word-sized integer.
index-type ::= `index`
```
Syntax:
-``` {.ebnf}
+```
// Sized integers like i1, i4, i8, i16, i32.
integer-type ::= `i` [1-9][0-9]*
```
Syntax:
-``` {.ebnf}
+```
memref-type ::= ranked-memref-type | unranked-memref-type
op will be needed for this purpose) and perform a switch and cast to a ranked
memref as a prerequisite to codegen.
-Example
-```mlir {.mlir}
-// With static ranks, we need a function for each
-// possible argument type
-%A = alloc() : memref<16x32xf32>
-%B = alloc() : memref<16x32x64xf32>
-call @helper_2D(%A) : (memref<16x32xf32>)->()
-call @helper_3D(%B) : (memref<16x32x64xf32>)->()
+Example ```mlir // With static ranks, we need a function for each // possible
+argument type %A = alloc() : memref<16x32xf32> %B = alloc() :
+memref<16x32x64xf32> call @helper_2D(%A) : (memref<16x32xf32>)->() call
+@helper_3D(%B) : (memref<16x32x64xf32>)->()
// With unknown rank, the functions can be unified under one unranked type
%A = alloc() : memref<16x32xf32>
Examples of memref static type
-```mlir {.mlir}
+```mlir
// Identity index/layout map
#identity = (d0, d1) -> (d0, d1)
Examples
-```mlir {.mlir}
+```mlir
// Allocates a memref with 2D index space:
// { (i, j) : 0 <= i < 16, 0 <= j < 32 }
%A = alloc() : memref<16x32xf32, #imapA, memspace0>
Layout map examples:
-```mlir {.mlir}
+```mlir
// MxN matrix stored in row major layout in memory:
#layout_map_row_major = (i, j) -> (i, j)
Syntax:
-``` {.ebnf}
+```
none-type ::= `none`
```
Syntax:
-``` {.ebnf}
+```
tensor-type ::= `tensor` `<` dimension-list tensor-memref-element-type `>`
tensor-memref-element-type ::= vector-element-type | vector-type | complex-type
Examples:
-```mlir {.mlir}
+```mlir
// Tensor with unknown rank.
tensor<* x f32>
Syntax:
-``` {.ebnf}
+```
tuple-type ::= `tuple` `<` (type ( `,` type)*)? `>`
```
Examples:
-```mlir {.mlir}
+```mlir
// Empty tuple.
tuple<>
Syntax:
-``` {.ebnf}
+```
vector-type ::= `vector` `<` static-dimension-list vector-element-type `>`
vector-element-type ::= float-type | integer-type
Syntax:
-``` {.ebnf}
+```
attribute-dict ::= `{` `}`
| `{` attribute-entry (`,` attribute-entry)* `}`
attribute-entry ::= dialect-attribute-entry | dependent-attribute-entry
Attribute values are represented by the following forms:
-``` {.ebnf}
+```
attribute-value ::= attribute-alias | dialect-attribute | standard-attribute
```
### Attribute Value Aliases
-``` {.ebnf}
+```
attribute-alias ::= '#' alias-name '=' attribute-value
attribute-alias ::= '#' alias-name
```
Example:
-```mlir {.mlir}
+```mlir
#map = (d0) -> (d0 + 10)
// Using the original attribute.
except that dialect attributes values are distinguished with a leading '#',
while dialect types are distinguished with a leading '!'.
-``` {.ebnf}
+```
dialect-attribute ::= '#' opaque-dialect-item
dialect-attribute ::= '#' pretty-dialect-item
```
Dialect attributes can be specified in a verbose form, e.g. like this:
-```mlir {.mlir}
+```mlir
// Complex attribute
#foo<"something<abcd>">
Dialect attributes that are simple enough can use the pretty format, which is a
lighter weight syntax that is equivalent to the above forms:
-```mlir {.mlir}
+```mlir
// Complex attribute
#foo.something<abcd>
```
[dialect attributes](#dialect-attribute-values) that are defined in a builtin
dialect and thus available to all users of MLIR.
-``` {.ebnf}
+```
standard-attribute ::= affine-map-attribute
| array-attribute
| bool-attribute
Syntax:
-``` {.ebnf}
+```
affine-map-attribute ::= affine-map
```
Syntax:
-``` {.ebnf}
+```
array-attribute ::= `[` (attribute-value (`,` attribute-value)*)? `]`
```
Syntax:
-``` {.ebnf}
+```
bool-attribute ::= bool-literal
```
Syntax:
-``` {.ebnf}
+```
dictionary-attribute ::= `{` (attribute-entry (`,` attribute-entry)*)? `}`
```
Syntax:
-``` {.ebnf}
+```
elements-attribute ::= dense-elements-attribute
| opaque-elements-attribute
| sparse-elements-attribute
Syntax:
-``` {.ebnf}
+```
dense-elements-attribute ::= `dense` `<` attribute-value `>` `:`
( tensor-type | vector-type )
```
Syntax:
-``` {.ebnf}
+```
opaque-elements-attribute ::= `opaque` `<` dialect-namespace `,`
hex-string-literal `>` `:`
( tensor-type | vector-type )
Syntax:
-``` {.ebnf}
+```
sparse-elements-attribute ::= `sparse` `<` attribute-value `,` attribute-value
`>` `:` ( tensor-type | vector-type )
```
Example:
-```mlir {.mlir}
+```mlir
sparse<[[0, 0], [1, 2]], [1, 5]> : tensor<3x4xi32>
// This represents the following tensor:
Syntax:
-``` {.ebnf}
+```
float-attribute ::= (float-literal (`:` float-type)?)
| (hexadecimal-literal `:` float-type)
```
Examples:
-``` {.mlir}
+```
42.0 // float attribute defaults to f64 type
42.0 : f32 // float attribute of f32 type
0x7C00 : f16 // positive infinity
Syntax:
-``` {.ebnf}
+```
integer-attribute ::= integer-literal ( `:` (index-type | integer-type) )?
```
Syntax:
-``` {.ebnf}
+```
integer-set-attribute ::= affine-map
```
Syntax:
-``` {.ebnf}
+```
string-attribute ::= string-literal (`:` type)?
```
Syntax:
-``` {.ebnf}
+```
symbol-ref-attribute ::= symbol-ref-id (`::` symbol-ref-id)*
```
Syntax:
-``` {.ebnf}
+```
type-attribute ::= type
```
#### Unit Attribute
-``` {.ebnf}
+```
unit-attribute ::= `unit`
```
false), but a value of false doesn't really bring any value. The parameter
either is the self/context or it isn't.
-```mlir {.mlir}
+```mlir
// A unit attribute defined with the `unit` value specifier.
func @verbose_form(i1) attributes {dialectName.unitAttr = unit}
simple pass that checks dependencies and emits them as "notes", allowing him to
write tests like this:
-```mlir {.mlir}
+```mlir
// RUN: mlir-opt %s -memref-dependence-check -verify-diagnostics
func @different_memrefs() {
%m.a = alloc() : memref<100xf32>
The [MLIR Tensor Type](LangRef.md#tensor-type) directly captures shape
information, so you can have things like:
-```mlir {.mlir}
+```mlir
%x = tf.Add %x, %y : tensor<128 x 8 x ? x f32>
```
Example:
-```mlir {.mlir}
+```mlir
func foo(...) {
%A = alloc <8x?xf32, #lmap> (%N)
...
Example:
-```mlir {.mlir}
+```mlir
%s = "foo"() : () -> !llvm<"i32*">
```
### Non-affine control flow
-```mlir {.mlir}
+```mlir
// A simple linear search in every row of a matrix
for (i = 0; i < N; i++) {
for (j = 0; j < N; j++) {
The presence of dynamic control flow leads to an inner non-affine function
nested in an outer function that using affine loops.
-```mlir {.mlir}
+```mlir
func @search(%A: memref<?x?xi32, %S: <?xi32>, %key : i32) {
%ni = dim %A, 0 : memref<?x?xi32>
// This loop can be parallelized
}
```
-```mlir {.mlir}
+```mlir
func @outer_nest(%n : index) {
affine.for %i = 0 to %n {
affine.for %j = 0 to %n {
convolution, which uses an integer set `#domain` to represent valid input data
in a dilated convolution.
-```mlir {.mlir}
+```mlir
// Dilation factors S0 and S1 can be constant folded if constant at compile time.
#domain = (d0, d1)[S0,S1,S2,S3]: (d0 % S0 == 0, d1 % S1 == 0, d0 >= 0, d1 >= 0,
S3 - d0 - 1 >= 0, S4 - d1 - 1 >= 0)
polyhedral dimension corresponding to an abstract fused loop for each ML
instruction that appears in that branch. Each leaf node is an ML Instruction.
-```mlir {.mlir}
+```mlir
// A tiled matmul code (128x128x128) represented in schedule tree form
// #map0 = (d0, d1, d2, d3, d4, d5) -> (128*d0 + d3, 128*d1 + d4, 128*d2 + d5)
Syntax:
-``` {.ebnf}
+```
// Affine relation definition at the top of file
affine-rel-def ::= affine-rel-id `=` affine-relation-inline
Example:
-```mlir {.mlir}
+```mlir
// read relation: two elements ( d0 <= r0 <= d0+1 )
##aff_rel9 = (d0) -> (r0) : r0 - d0 >= 0, d0 - r0 + 1 >= 0
Example:
-```mlir {.mlir}
+```mlir
##rel9 ( ) [s0] -> (r0, r1) : 0 <= r0 <= 1023, 0 <= r1 <= s0 - 1
func @cblas_reduce_ffi(%M: memref<1024 x ? x f32, #layout_map0, /*mem=*/0>)
Syntax:
-``` {.ebnf}
+```
[<out-var-list> =]
for %<index-variable-name> = <lower-bound> ... <upper-bound> step <step>
[with <in-var-list>] { <loop-instruction-list> }
Example:
-```mlir {.mlir}
+```mlir
// Return sum of elements in 1-dimensional mref A
func i32 @sum(%A : memref<?xi32>, %N : i32) -> (i32) {
%init = 0
Syntax:
-``` {.ebnf}
+```
<out-var-list> = affine.if (<cond-list>) {...} [else {...}]
```
Example:
-```mlir {.mlir}
+```mlir
// Compute sum of half of the array
func i32 @sum_half(%A : memref<?xi32>, %N : i32) -> (i32) {
%s0 = 0
An example FileCheck test is shown below:
-```mlir {.mlir}
+```mlir
// RUN: mlir-opt %s -cse | FileCheck %s
// CHECK-LABEL: func @simple_constant
self-contained as possible and focus on testing the minimal set of
functionalities needed. Let's see an example:
-```mlir {.mlir}
+```mlir
// RUN: mlir-opt %s -cse | FileCheck %s
// CHECK-LABEL: func @simple_constant() -> (i32, i32)
If we naively remove the unrelated `CHECK` lines in our source file, we may end
up with:
-```mlir {.mlir}
+```mlir
// CHECK-LABEL: func @simple_constant
func @simple_constant() -> (i32, i32) {
// CHECK-NEXT: %result = constant 1 : i32
Utilizing the above, we end up with the example shown in the main
[FileCheck tests](#filecheck-tests) section.
-```mlir {.mlir}
+```mlir
// CHECK-LABEL: func @simple_constant
func @simple_constant() -> (i32, i32) {
/// Here we use a substitution variable as the output of the constant is
An example .mlir test running under `mlir-opt` is shown below:
-```mlir {.mlir}
+```mlir
// RUN: mlir-opt %s -split-input-file -verify-diagnostics
// Expect an error on the same line.
parser and printer is to provide the `type-data` bits. We will define our
`StructType` as having the following form:
-``` {.ebnf}
+```
struct-type ::= `struct` `<` type (`,` type)* `>`
```
Syntax:
- ``` {.ebnf}
+ ```
op ::= `gpu.func` symbol-ref-id `(` argument-list `)` (`->`
function-result-list)?
memory-attribution `kernel`? function-attributes? region
Syntax:
- ``` {.ebnf}
+ ```
operation ::= `gpu.launch` `block` `(` ssa-id-list `)` `in` ssa-reassignment
`threads` `(` ssa-id-list `)` `in` ssa-reassignment
(`args` ssa-reassignment `:` type-list)?
Example:
- ```mlir {.mlir}
+ ```mlir
gpu.launch blocks(%bx, %by, %bz) in (%sz_bx = %0, %sz_by = %1, %sz_bz = %2)
threads(%tx, %ty, %tz) in (%sz_tx = %3, %sz_ty = %4, %sz_tz = %5)
args(%arg0 = %6, %arg1 = 7) : f32, memref<?xf32, 1> {
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fadd-op ::= ssa-id `=` `spv.FAdd` ssa-use, ssa-use
if Operand 2 is 0.
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fdiv-op ::= ssa-id `=` `spv.FDiv` ssa-use, ssa-use
sign of Operand 2.
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fmod-op ::= ssa-id `=` `spv.FMod` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fmul-op ::= `spv.FMul` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fmul-op ::= `spv.FNegate` ssa-use `:` float-scalar-vector-type
sign of Operand 1.
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
frem-op ::= ssa-id `=` `spv.FRemOp` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fsub-op ::= ssa-id `=` `spv.FRemOp` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
iadd-op ::= ssa-id `=` `spv.IAdd` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
imul-op ::= ssa-id `=` `spv.IMul` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
isub-op ::= `spv.ISub` ssa-use, ssa-use
if Operand 2 is 0.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
sdiv-op ::= ssa-id `=` `spv.SDiv` ssa-use, ssa-use
sign of Operand 2.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
smod-op ::= ssa-id `=` `spv.SMod` ssa-use, ssa-use
sign of Operand 1.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
srem-op ::= ssa-id `=` `spv.SRem` ssa-use, ssa-use
if Operand 2 is 0.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
udiv-op ::= ssa-id `=` `spv.UDiv` ssa-use, ssa-use
if Operand 2 is 0.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
umod-op ::= ssa-id `=` `spv.UMod` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
scope ::= `"CrossDevice"` | `"Device"` | `"Workgroup"` | ...
memory-semantics ::= `"None"` | `"Acquire"` | "Release"` | ...
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
bitcount-op ::= ssa-id `=` `spv.BitCount` ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
bitfield-insert-op ::= ssa-id `=` `spv.BitFieldInsert` ssa-use `,` ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
bitfield-extract-s-op ::= ssa-id `=` `spv.BitFieldSExtract` ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
bitfield-extract-u-op ::= ssa-id `=` `spv.BitFieldUExtract` ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
bitreverse-op ::= ssa-id `=` `spv.BitReverse` ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
bitwise-and-op ::= ssa-id `=` `spv.BitwiseAnd` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
bitwise-or-op ::= ssa-id `=` `spv.BitwiseOr` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
bitwise-xor-op ::= ssa-id `=` `spv.BitwiseXor` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
shift-left-logical-op ::= ssa-id `=` `spv.ShiftLeftLogical`
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
shift-right-arithmetic-op ::= ssa-id `=` `spv.ShiftRightArithmetic`
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
shift-right-logical-op ::= ssa-id `=` `spv.ShiftRightLogical`
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
not-op ::= ssa-id `=` `spv.BitNot` ssa-use `:` integer-scalar-vector-type
### Custom assembly form
- ``` {.ebnf}
+ ```
bitcast-op ::= ssa-id `=` `spv.Bitcast` ssa-use
`:` operand-type `to` result-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
convert-f-to-s-op ::= ssa-id `=` `spv.ConvertFToSOp` ssa-use
`:` operand-type `to` result-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
convert-f-to-u-op ::= ssa-id `=` `spv.ConvertFToUOp` ssa-use
`:` operand-type `to` result-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
convert-s-to-f-op ::= ssa-id `=` `spv.ConvertSToFOp` ssa-use
`:` operand-type `to` result-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
convert-u-to-f-op ::= ssa-id `=` `spv.ConvertUToFOp` ssa-use
`:` operand-type `to` result-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
f-convert-op ::= ssa-id `=` `spv.FConvertOp` ssa-use
`:` operand-type `to` result-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
s-convert-op ::= ssa-id `=` `spv.SConvertOp` ssa-use
`:` operand-type `to` result-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
u-convert-op ::= ssa-id `=` `spv.UConvertOp` ssa-use
`:` operand-type `to` result-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
composite-construct-op ::= ssa-id `=` `spv.CompositeConstruct`
(ssa-use (`,` ssa-use)* )? `:` composite-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
composite-extract-op ::= ssa-id `=` `spv.CompositeExtract` ssa-use
`[` integer-literal (',' integer-literal)* `]`
`:` composite-type
### Custom assembly form
- ``` {.ebnf}
+ ```
composite-insert-op ::= ssa-id `=` `spv.CompositeInsert` ssa-use, ssa-use
`[` integer-literal (',' integer-literal)* `]`
`:` object-type `into` composite-type
### Custom assembly form
- ``` {.ebnf}
+ ```
branch-op ::= `spv.Branch` successor
successor ::= bb-id branch-use-list?
branch-use-list ::= `(` ssa-use-list `:` type-list-no-parens `)`
### Custom assembly form
- ``` {.ebnf}
+ ```
branch-conditional-op ::= `spv.BranchConditional` ssa-use
(`[` integer-literal, integer-literal `]`)?
`,` successor `,` successor
### Custom assembly form
- ``` {.ebnf}
+ ```
function-call-op ::= `spv.FunctionCall` function-id `(` ssa-use-list `)`
`:` function-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
return-op ::= `spv.Return`
```
}];
### Custom assembly form
- ``` {.ebnf}
+ ```
unreachable-op ::= `spv.Unreachable`
```
}];
### Custom assembly form
- ``` {.ebnf}
+ ```
return-value-op ::= `spv.ReturnValue` ssa-use `:` spirv-type
```
per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
abs-op ::= ssa-id `=` `spv.GLSL.FAbs` ssa-use `:`
with the same component width. Results are computed per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
abs-op ::= ssa-id `=` `spv.GLSL.SAbs` ssa-use `:`
per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
ceil-op ::= ssa-id `=` `spv.GLSL.Ceil` ssa-use `:`
per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
restricted-float-scalar-type ::= `f16` | `f32`
restricted-float-scalar-vector-type ::=
restricted-float-scalar-type |
computed per component.";
### Custom assembly format
- ``` {.ebnf}
+ ```
restricted-float-scalar-type ::= `f16` | `f32`
restricted-float-scalar-vector-type ::=
restricted-float-scalar-type |
per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
floor-op ::= ssa-id `=` `spv.GLSL.Floor` ssa-use `:`
per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
rsqrt-op ::= ssa-id `=` `spv.GLSL.InverseSqrt` ssa-use `:`
per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
restricted-float-scalar-type ::= `f16` | `f32`
restricted-float-scalar-vector-type ::=
restricted-float-scalar-type |
type. Results are computed per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fmax-op ::= ssa-id `=` `spv.GLSL.FMax` ssa-use `:`
component.
### Custom assembly format
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
smax-op ::= ssa-id `=` `spv.GLSL.SMax` ssa-use `:`
computed per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fmin-op ::= ssa-id `=` `spv.GLSL.FMin` ssa-use `:`
component.
### Custom assembly format
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
smin-op ::= ssa-id `=` `spv.GLSL.SMin` ssa-use `:`
per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
sign-op ::= ssa-id `=` `spv.GLSL.FSign` ssa-use `:`
with the same component width. Results are computed per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
sign-op ::= ssa-id `=` `spv.GLSL.SSign` ssa-use `:`
per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
sqrt-op ::= ssa-id `=` `spv.GLSL.Sqrt` ssa-use `:`
per component.
### Custom assembly format
- ``` {.ebnf}
+ ```
restricted-float-scalar-type ::= `f16` | `f32`
restricted-float-scalar-vector-type ::=
restricted-float-scalar-type |
### Custom assembly form
- ``` {.ebnf}
+ ```
subgroup-ballot-op ::= ssa-id `=` `spv.SubgroupBallotKHR`
ssa-use `:` `vector` `<` 4 `x` `i32` `>`
```
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fordequal-op ::= ssa-id `=` `spv.FOrdEqual` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fordgt-op ::= ssa-id `=` `spv.FOrdGreaterThan` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fordgte-op ::= ssa-id `=` `spv.FOrdGreaterThanEqual` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fordlt-op ::= ssa-id `=` `spv.FOrdLessThan` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fordlte-op ::= ssa-id `=` `spv.FOrdLessThanEqual` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
fordneq-op ::= ssa-id `=` `spv.FOrdNotEqual` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
funordequal-op ::= ssa-id `=` `spv.FUnordEqual` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
funordgt-op ::= ssa-id `=` `spv.FUnordGreaterThan` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
funordgte-op ::= ssa-id `=` `spv.FUnordGreaterThanEqual` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
funordlt-op ::= ssa-id `=` `spv.FUnordLessThan` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
funordlte-op ::= ssa-id `=` `spv.FUnordLessThanEqual` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
float-scalar-vector-type ::= float-type |
`vector<` integer-literal `x` float-type `>`
funordneq-op ::= ssa-id `=` `spv.FUnordNotEqual` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
iequal-op ::= ssa-id `=` `spv.IEqual` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
inot-equal-op ::= ssa-id `=` `spv.INotEqual` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
logical-and ::= `spv.LogicalAnd` ssa-use `,` ssa-use
`:` operand-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
logical-equal ::= `spv.LogicalEqual` ssa-use `,` ssa-use
`:` operand-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
logical-not ::= `spv.LogicalNot` ssa-use `:` operand-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
logical-not-equal ::= `spv.LogicalNotEqual` ssa-use `,` ssa-use
`:` operand-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
logical-or ::= `spv.LogicalOr` ssa-use `,` ssa-use
`:` operand-type
```
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
sgreater-than-op ::= ssa-id `=` `spv.SGreaterThan` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
sgreater-than-equal-op ::= ssa-id `=` `spv.SGreaterThanEqual` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
sless-than-op ::= ssa-id `=` `spv.SLessThan` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
sless-than-equal-op ::= ssa-id `=` `spv.SLessThanEqual` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
scalar-type ::= integer-type | float-type | boolean-type
select-object-type ::= scalar-type
| `vector<` integer-literal `x` scalar-type `>`
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
ugreater-than-op ::= ssa-id `=` `spv.UGreaterThan` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
ugreater-than-equal-op ::= ssa-id `=` `spv.UGreaterThanEqual` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
uless-than-op ::= ssa-id `=` `spv.ULessThan` ssa-use, ssa-use
Results are computed per component.
### Custom assembly form
- ``` {.ebnf}
+ ```
integer-scalar-vector-type ::= integer-type |
`vector<` integer-literal `x` integer-type `>`
uless-than-equal-op ::= ssa-id `=` `spv.ULessThanEqual` ssa-use, ssa-use
### Custom assembly form
- ``` {.ebnf}
+ ```
scope ::= `"Workgroup"` | `"Subgroup"`
non-uniform-ballot-op ::= ssa-id `=` `spv.GroupNonUniformBallot` scope
ssa-use `:` `vector` `<` 4 `x` `integer-type` `>`
- must be an OpConstant when indexing into a structure.
### Custom assembly form
- ``` {.ebnf}
+ ```
access-chain-op ::= ssa-id `=` `spv.AccessChain` ssa-use
`[` ssa-use (',' ssa-use)* `]`
`:` pointer-type
### Custom assembly form
- ``` {.ebnf}
+ ```
scope ::= `"CrossDevice"` | `"Device"` | `"Workgroup"` | ...
memory-semantics ::= `"None"` | `"Acquire"` | "Release"` | ...
### Custom assembly form
- ``` {.ebnf}
+ ```
execution-mode ::= "Invocations" | "SpacingEqual" |
<and other SPIR-V execution modes...>
### Custom assembly form
- ``` {.ebnf}
+ ```
memory-access ::= `"None"` | `"Volatile"` | `"Aligned", ` integer-literal
| `"NonTemporal"`
### Custom assembly form
- ``` {.ebnf}
+ ```
scope ::= `"CrossDevice"` | `"Device"` | `"Workgroup"` | ...
memory-semantics ::= `"None"` | `"Acquire"` | `"Release"` | ...
### Custom assembly form
- ``` {.ebnf}
+ ```
store-op ::= `spv.Store ` storage-class ssa-use `, ` ssa-use `, `
(`[` memory-access `]`)? `:` spirv-element-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
undef-op ::= `spv.undef` `:` spirv-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
variable-op ::= ssa-id `=` `spv.Variable` (`init(` ssa-use `)`)?
(`bind(` integer-literal, integer-literal `)`)?
(`built_in(` string-literal `)`)?
### Custom assembly form
- ``` {.ebnf}
+ ```
spv-address-of-op ::= ssa-id `=` `spv._address_of` symbol-ref-id
`:` spirv-pointer-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
spv-constant-op ::= ssa-id `=` `spv.constant` attribute-value
(`:` spirv-type)?
```
### Custom assembly form
- ``` {.ebnf}
+ ```
execution-model ::= "Vertex" | "TesellationControl" |
<and other SPIR-V execution models...>
### Custom assembly form
- ``` {.ebnf}
+ ```
variable-op ::= `spv.globalVariable` spirv-type symbol-ref-id
(`initializer(` symbol-ref-id `)`)?
(`bind(` integer-literal, integer-literal `)`)?
### Custom assembly form
- ``` {.ebnf}
+ ```
addressing-model ::= `"Logical"` | `"Physical32"` | `"Physical64"`
memory-model ::= `"Simple"` | `"GLSL450"` | `"OpenCL"` | `"VulkanKHR"`
spv-module-op ::= `spv.module` addressing-model memory-model
### Custom assembly form
- ``` {.ebnf}
+ ```
spv-reference-of-op ::= ssa-id `=` `spv._reference_of` symbol-ref-id
`:` spirv-scalar-type
```
### Custom assembly form
- ``` {.ebnf}
+ ```
spv-spec-constant-op ::= `spv.specConstant` symbol-ref-id
`spec_id(` integer `)`
`=` attribute-value (`:` spirv-type)?
More precisely, let's dive deeper into the permutation_map for the following
MLIR:
- ```mlir {.mlir}
+ ```mlir
vector.transfer_read %A[%expr1, %expr2, %expr3, %expr4]
{ permutation_map : (d0,d1,d2,d3) -> (d2,0,d0) } :
memref<?x?x?x?xf32>, vector<3x4x5xf32>
implemented using a warp-shuffle if loop `j` were mapped to `threadIdx.x`.
Syntax
- ``` {.ebnf}
+ ```
operation ::= ssa-id `=` `vector.transfer_read` ssa-use-list
`{` attribute-entry `} :` memref-type `,` vector-type
```
Syntax:
- ``` {.ebnf}
+ ```
operation ::= `vector.transfer_write` ssa-use-list `{` attribute-entry `} :
` vector-type ', ' memref-type '
```
Examples:
- ```mlir {.mlir}
+ ```mlir
// write vector<16x32x64xf32> into the slice
// `%A[%i0, %i1:%i1+32, %i2:%i2+64, %i3:%i3+16]`:
for %i0 = 0 to %0 {
Syntax:
- ``` {.ebnf}
+ ```
operation ::= `vector.type_cast` ssa-use : memref-type to memref-type
```
///
/// Consider the case:
///
-/// ```mlir {.mlir}
+/// ```mlir
/// // Read the slice `%A[%i0, %i1:%i1+256, %i2:%i2+32]` into
/// // vector<32x256xf32> and pad with %f0 to handle the boundary case:
/// %f0 = constant 0.0f : f32
/// resembling the following (while guaranteeing an always full-tile
/// abstraction):
///
-/// ```mlir {.mlir}
+/// ```mlir
/// loop.for %d2 = 0 to %c256 {
/// loop.for %d1 = 0 to %c32 {
/// %s = %A[%i0, %i1 + %d1, %i2 + %d2] : f32
description = existing_info.get('description', None)
if description is None:
- assembly = '\n ``` {.ebnf}\n'\
+ assembly = '\n ```\n'\
' [TODO]\n'\
' ```\n\n'\
' For example:\n\n'\
projects / platform / upstream / llvm.git / commitdiff