From ac4873322f0951a00f8591f2305337ca3d6c7762 Mon Sep 17 00:00:00 2001 From: Alex Zinenko Date: Tue, 10 Dec 2019 03:00:29 -0800 Subject: [PATCH] Drop Markdown style annotations These come from a non-standard extenion that is not available on Github, so it only clutters the documentation source with {.mlir} or {.ebnf} tags. PiperOrigin-RevId: 284733003 --- mlir/g3doc/ConversionToLLVMDialect.md | 22 +-- mlir/g3doc/Diagnostics.md | 12 +- mlir/g3doc/Dialects/Affine.md | 87 ++++++------ mlir/g3doc/Dialects/GPU.md | 35 +++-- mlir/g3doc/Dialects/LLVM.md | 30 ++-- mlir/g3doc/Dialects/SPIR-V.md | 16 +-- mlir/g3doc/Dialects/Standard.md | 152 ++++++++++----------- mlir/g3doc/EDSC.md | 4 +- mlir/g3doc/LangRef.md | 142 ++++++++++--------- mlir/g3doc/MLIRForGraphAlgorithms.md | 4 +- mlir/g3doc/Rationale.md | 28 ++-- mlir/g3doc/TestingGuide.md | 10 +- mlir/g3doc/Tutorials/Toy/Ch-7.md | 2 +- mlir/include/mlir/Dialect/GPU/GPUOps.td | 6 +- .../mlir/Dialect/SPIRV/SPIRVArithmeticOps.td | 30 ++-- mlir/include/mlir/Dialect/SPIRV/SPIRVAtomicOps.td | 2 +- mlir/include/mlir/Dialect/SPIRV/SPIRVBitOps.td | 24 ++-- mlir/include/mlir/Dialect/SPIRV/SPIRVCastOps.td | 16 +-- .../mlir/Dialect/SPIRV/SPIRVCompositeOps.td | 6 +- .../mlir/Dialect/SPIRV/SPIRVControlFlowOps.td | 12 +- mlir/include/mlir/Dialect/SPIRV/SPIRVGLSLOps.td | 32 ++--- mlir/include/mlir/Dialect/SPIRV/SPIRVGroupOps.td | 2 +- mlir/include/mlir/Dialect/SPIRV/SPIRVLogicalOps.td | 56 ++++---- .../mlir/Dialect/SPIRV/SPIRVNonUniformOps.td | 2 +- mlir/include/mlir/Dialect/SPIRV/SPIRVOps.td | 16 +-- .../mlir/Dialect/SPIRV/SPIRVStructureOps.td | 14 +- mlir/include/mlir/Dialect/VectorOps/VectorOps.td | 10 +- .../VectorToLoops/ConvertVectorToLoops.cpp | 4 +- mlir/utils/spirv/gen_spirv_dialect.py | 2 +- 29 files changed, 388 insertions(+), 390 deletions(-) diff --git a/mlir/g3doc/ConversionToLLVMDialect.md b/mlir/g3doc/ConversionToLLVMDialect.md index 595049a..782889f 100644 --- a/mlir/g3doc/ConversionToLLVMDialect.md +++ b/mlir/g3doc/ConversionToLLVMDialect.md @@ -79,7 +79,7 @@ resulting in a struct containing two pointers + offset. Examples: -```mlir {.mlir} +```mlir memref -> !llvm.type<"{ float*, float*, i64 }"> memref<1 x f32> -> !llvm.type<"{ float*, float*, i64, [1 x i64], [1 x i64] }"> memref -> !llvm.type<"{ float*, float*, i64, [1 x i64], [1 x i64] }"> @@ -102,7 +102,7 @@ descriptor pointer. Examples: -```mlir {.mlir} +```mlir // unranked descriptor memref<*xf32> -> !llvm.type<"{i64, i8*}"> ``` @@ -125,7 +125,7 @@ converted using these rules. Examples: -```mlir {.mlir} +```mlir // zero-ary function type with no results. () -> () // is converted to a zero-ary function with `void` result @@ -162,7 +162,7 @@ definition operation uses MLIR syntax. Examples: -```mlir {.mlir} +```mlir // zero-ary function type with no results. func @foo() -> () // gets LLVM type void(). @@ -195,7 +195,7 @@ defines and uses of the values being returned. Example: -```mlir {.mlir} +```mlir func @foo(%arg0: i32, %arg1: i64) -> (i32, i64) { return %arg0, %arg1 : i32, i64 } @@ -249,7 +249,7 @@ are used instead of them in the original terminator operation. Example: -```mlir {.mlir} +```mlir cond_br %0, ^bb1(%1 : i32), ^bb1(%2 : i32) ^bb1(%3 : i32) "use"(%3) : (i32) -> () @@ -257,7 +257,7 @@ Example: leads to a new basic block being inserted, -```mlir {.mlir} +```mlir cond_br %0, ^bb1(%1 : i32), ^dummy ^bb1(%3 : i32): "use"(%3) : (i32) -> () @@ -267,7 +267,7 @@ leads to a new basic block being inserted, 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">) -> () @@ -307,7 +307,7 @@ Examples: An access to a zero-dimensional memref is converted into a plain load: -```mlir {.mlir} +```mlir // before %0 = load %m[] : memref @@ -317,13 +317,13 @@ An access to a zero-dimensional memref is converted into a plain load: 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}"> diff --git a/mlir/g3doc/Diagnostics.md b/mlir/g3doc/Diagnostics.md index 0c6ef7a..69a3094 100644 --- a/mlir/g3doc/Diagnostics.md +++ b/mlir/g3doc/Diagnostics.md @@ -16,7 +16,7 @@ different location types depending on the situational need. ### CallSite Location -``` {.ebnf} +``` callsite-location ::= 'callsite' '(' location 'at' location ')' ``` @@ -26,7 +26,7 @@ location usages. This connects a location of a `callee` with the location of a ### FileLineCol Location -``` {.ebnf} +``` filelinecol-location ::= string-literal ':' integer-literal ':' integer-literal ``` @@ -36,7 +36,7 @@ languages. ### Fused Location -``` {.ebnf} +``` fused-location ::= `fused` fusion-metadata? '[' location (location ',')* ']' fusion-metadata ::= '<' attribute-value '>' ``` @@ -49,7 +49,7 @@ loss of location information. With `fused` locations, this is a non-issue. ### Name Location -``` {.ebnf} +``` name-location ::= string-literal ('(' location ')')? ``` @@ -66,7 +66,7 @@ optional location is used during serialization. ### Unknown Location -``` {.ebnf} +``` unknown-location ::= `unknown` ``` @@ -311,7 +311,7 @@ annotate your source file with expected diagnostics in the form of: 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}} diff --git a/mlir/g3doc/Dialects/Affine.md b/mlir/g3doc/Dialects/Affine.md index 11cb93c..c5dcf6a 100644 --- a/mlir/g3doc/Dialects/Affine.md +++ b/mlir/g3doc/Dialects/Affine.md @@ -19,7 +19,7 @@ brackets. 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) @@ -36,7 +36,7 @@ use the same parenthesized vs square bracket list to distinguish the two. Syntax: -``` {.ebnf} +``` // Uses of SSA values that are passed to dimensional identifiers. dim-use-list ::= `(` ssa-use-list? `)` @@ -51,7 +51,7 @@ SSA values bound to dimensions and symbols must always have 'index' type. 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> @@ -79,7 +79,7 @@ other dimensions and symbols). Syntax: -``` {.ebnf} +``` affine-expr ::= `(` affine-expr `)` | affine-expr `+` affine-expr | affine-expr `-` affine-expr @@ -131,7 +131,7 @@ i^2)$$, $$(i \mod j, i/j)$$ are not affine functions of $$(i, j)$$. Syntax: -``` {.ebnf} +``` affine-map-inline ::= dim-and-symbol-id-lists `->` multi-dim-affine-expr ``` @@ -158,7 +158,7 @@ the representation closed with respect to several operations of interest. Syntax: -``` {.ebnf} +``` affine-map-id ::= `#` suffix-id // Definitions of affine maps are at the top of the file. @@ -175,7 +175,7 @@ name. Examples: -```mlir {.mlir} +```mlir // Affine map out-of-line definition and usage example. #affine_map42 = (d0, d1)[s0] -> (d0, d0 + d1 + s0 floordiv 2) @@ -195,7 +195,7 @@ Semi-affine maps are thus a strict superset of affine maps. Syntax of semi-affine expressions: -``` {.ebnf} +``` semi-affine-expr ::= `(` semi-affine-expr `)` | semi-affine-expr `+` semi-affine-expr | semi-affine-expr `-` semi-affine-expr @@ -216,7 +216,7 @@ as that for [affine expressions](#affine-expressions). Syntax of semi-affine maps: -``` {.ebnf} +``` semi-affine-map-inline ::= dim-and-symbol-id-lists `->` multi-dim-semi-affine-expr ``` @@ -225,7 +225,7 @@ Semi-affine maps may be defined inline at the point of use, or may be hoisted to 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. @@ -247,7 +247,7 @@ while its symbols are enclosed in square brackets. Syntax of affine constraints: -``` {.ebnf} +``` affine-constraint ::= affine-expr `>=` `0` | affine-expr `==` `0` affine-constraint-conjunction ::= affine-constraint (`,` affine-constraint)* @@ -257,7 +257,7 @@ Integer sets may be defined inline at the point of use, or may be hoisted to the 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 @@ -278,7 +278,7 @@ dimensions. 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) @@ -298,7 +298,7 @@ affine.if #set42(%i, %j)[%M, %N] { Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `affine.apply` affine-map dim-and-symbol-use-list ``` @@ -311,7 +311,7 @@ value. The input operands and result must all have 'index' type. Example: -```mlir {.mlir} +```mlir #map10 = (d0, d1) -> (d0 floordiv 8 + d1 floordiv 128) ... %1 = affine.apply #map10 (%s, %t) @@ -324,7 +324,7 @@ Example: Syntax: -``` {.ebnf} +``` operation ::= `affine.for` ssa-id `=` lower-bound `to` upper-bound (`step` integer-literal)? `{` op* `}` @@ -365,7 +365,7 @@ nullary mapping function that returns the constant value (e.g. `()->(-42)()`). Example showing reverse iteration of the inner loop: -```mlir {.mlir} +```mlir #map57 = (d0)[s0] -> (s0 - d0 - 1) func @simple_example(%A: memref, %B: memref) { @@ -385,7 +385,7 @@ func @simple_example(%A: memref, %B: memref) { Syntax: -``` {.ebnf} +``` operation ::= `affine.if` if-op-cond `{` op* `}` (`else` `{` op* `}`)? if-op-cond ::= integer-set dim-and-symbol-use-list ``` @@ -409,7 +409,7 @@ blocks must not have any arguments. 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) { @@ -431,19 +431,20 @@ 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: @@ -460,19 +461,20 @@ Example: 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: @@ -489,7 +491,7 @@ Example: 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 ``` @@ -515,7 +517,7 @@ specified. The value of 'num_elements' must be a multiple of 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 @@ -536,11 +538,12 @@ space 1 at indices [%k + 7, %l], would be specified as follows: %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 ``` @@ -553,7 +556,7 @@ associated with the DMA operation. For example: 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> @@ -567,7 +570,7 @@ Example: Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `affine.min` affine-map dim-and-symbol-use-list ``` @@ -580,7 +583,7 @@ returns one value. The input operands and result must all have 'index' type. Example: -```mlir {.mlir} +```mlir %0 = affine.min (d0)[s0] -> (1000, d0 + 512, s0) (%arg0)[%arg1] @@ -590,7 +593,7 @@ Example: Syntax: -``` {.ebnf} +``` operation ::= `"affine.terminator"() : () -> ()` ``` diff --git a/mlir/g3doc/Dialects/GPU.md b/mlir/g3doc/Dialects/GPU.md index bcb677d..f29179a 100644 --- a/mlir/g3doc/Dialects/GPU.md +++ b/mlir/g3doc/Dialects/GPU.md @@ -43,7 +43,7 @@ x, y, or z `dimension`. Example: -```mlir {.mlir} +```mlir %bDimX = "gpu.block_dim"() {dimension = "x"} : () -> (index) ``` @@ -54,7 +54,7 @@ the x, y, or z `dimension`. Example: -```mlir {.mlir} +```mlir %bIdY = "gpu.block_id"() {dimension = "y"} : () -> (index) ``` @@ -65,7 +65,7 @@ Returns the number of thread blocks in the grid along the x, y, or z Example: -```mlir {.mlir} +```mlir %gDimZ = "gpu.grid_dim"() {dimension = "z"} : () -> (index) ``` @@ -93,7 +93,7 @@ A custom syntax for this operation is currently not available. Example: -```mlir {.mlir} +```mlir module attributes {gpu.container_module} { // This module creates a separate compilation unit for the GPU compiler. @@ -140,7 +140,7 @@ along the x, y, or z `dimension`. Example: -```mlir {.mlir} +```mlir %tIdX = "gpu.thread_id"() {dimension = "x"} : () -> (index) ``` @@ -151,11 +151,10 @@ returns values to the immediately enclosing gpu op. 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 @@ -163,7 +162,7 @@ workgroup. The result is equal for all work items of a workgroup. 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): @@ -171,10 +170,10 @@ For example, both "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. @@ -184,14 +183,14 @@ 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. diff --git a/mlir/g3doc/Dialects/LLVM.md b/mlir/g3doc/Dialects/LLVM.md index 9791352..00d0fa0 100644 --- a/mlir/g3doc/Dialects/LLVM.md +++ b/mlir/g3doc/Dialects/LLVM.md @@ -25,7 +25,7 @@ must exist in the dialect's context. 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 ::= ``` @@ -60,7 +60,7 @@ type. LLVM function operation is intended to capture additional properties of 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 @@ -100,7 +100,7 @@ same type. Examples: -```mlir {.mlir} +```mlir // Integer addition. %0 = llvm.add %a, %b : !llvm.i32 @@ -121,7 +121,7 @@ the same type. Examples: -```mlir {.mlir} +```mlir // Float addition. %0 = llvm.fadd %a, %b : !llvm.float @@ -146,7 +146,7 @@ non-pointer arguments of LLVM IR's `getelementptr`. 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*"> @@ -179,7 +179,7 @@ they are modeled as array attributes. 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}"> @@ -209,7 +209,7 @@ arguments. Examples: -```mlir {.mlir} +```mlir // Branch without arguments. ^bb0: llvm.br ^bb0 @@ -248,7 +248,7 @@ wrapped LLVM IR function type. Examples: -```mlir {.mlir} +```mlir // Direct call without arguments and with one result. %0 = llvm.call @foo() : () -> (!llvm.float) @@ -292,7 +292,7 @@ referenced. If the global value is a constant, storing into it is not allowed. Examples: -```mlir {.mlir} +```mlir func @foo() { // Get the address of a global. %0 = llvm.mlir.addressof @const : !llvm<"i32*"> @@ -319,7 +319,7 @@ to the attribute type converted to LLVM IR. Examples: -```mlir {.mlir} +```mlir // Integer constant, internal i32 is mandatory %0 = llvm.mlir.constant(42 : i32) : !llvm.i32 @@ -343,7 +343,7 @@ both cases. 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 ``` @@ -354,7 +354,7 @@ types must be compatible. 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*"> { @@ -374,7 +374,7 @@ other @-identifiers in it. Examples: -```mlir {.mlir} +```mlir // Global values use @-identifiers. llvm.mlir.global constant @cst(42 : i32) : !llvm.i32 @@ -406,7 +406,7 @@ type. Examples: -```mlir {.mlir} +```mlir // Null pointer to i8 value. %0 = llvm.mlir.null : !llvm<"i8*"> @@ -423,7 +423,7 @@ dialect type wrapping an LLVM IR structure type. Example: -```mlir {.mlir} +```mlir // Create a structure with a 32-bit integer followed by a float. %0 = llvm.mlir.undef : !llvm<"{i32, float}"> ``` diff --git a/mlir/g3doc/Dialects/SPIR-V.md b/mlir/g3doc/Dialects/SPIR-V.md index 58bd5ee..1413b18 100644 --- a/mlir/g3doc/Dialects/SPIR-V.md +++ b/mlir/g3doc/Dialects/SPIR-V.md @@ -84,7 +84,7 @@ instructions are represented in the SPIR-V dialect. Notably, 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 @@ -94,7 +94,7 @@ spirv-type ::= array-type This corresponds to SPIR-V [array type][ArrayType]. Its syntax is -``` {.ebnf} +``` element-type ::= integer-type | floating-point-type | vector-type @@ -114,7 +114,7 @@ For example, This corresponds to SPIR-V [image type][ImageType]. Its syntax is -``` {.ebnf} +``` dim ::= `1D` | `2D` | `3D` | `Cube` | depth-info ::= `NoDepth` | `IsDepth` | `DepthUnknown` @@ -134,7 +134,7 @@ image-type ::= `!spv.image<` element-type `,` dim `,` depth-info `,` For example, -``` {.mlir} +``` !spv.image !spv.image ``` @@ -143,7 +143,7 @@ For example, This corresponds to SPIR-V [pointer type][PointerType]. Its syntax is -``` {.ebnf} +``` storage-class ::= `UniformConstant` | `Uniform` | `Workgroup` @@ -163,7 +163,7 @@ For example, This corresponds to SPIR-V [runtime array type][RuntimeArrayType]. Its syntax is -``` {.ebnf} +``` runtime-array-type ::= `!spv.rtarray<` element-type `>` ``` @@ -178,7 +178,7 @@ For example, 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 `]`)? @@ -186,7 +186,7 @@ struct-type ::= `!spv.struct<` spirv-type (`[` struct-member-decoration `]`)? For Example, -``` {.mlir} +``` !spv.struct !spv.struct !spv.struct> diff --git a/mlir/g3doc/Dialects/Standard.md b/mlir/g3doc/Dialects/Standard.md index 9d53eba3..fb0a4ed 100644 --- a/mlir/g3doc/Dialects/Standard.md +++ b/mlir/g3doc/Dialects/Standard.md @@ -21,7 +21,7 @@ list of successors, i.e. other blocks to which the control flow will proceed. Syntax: -``` {.ebnf} +``` operation ::= `br` successor successor ::= bb-id branch-use-list? branch-use-list ::= `(` ssa-use-list `:` type-list-no-parens `)` @@ -37,7 +37,7 @@ The MLIR branch operation is not allowed to target the entry block for a region. Syntax: -``` {.ebnf} +``` operation ::= `cond_br` ssa-use `,` successor `,` successor ``` @@ -53,7 +53,7 @@ allowed to be the same. 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 @@ -68,7 +68,7 @@ func @select(i32, i32, i1) -> i32 { Syntax: -``` {.ebnf} +``` operation ::= `return` (ssa-use-list `:` type-list-no-parens)? ``` @@ -83,8 +83,8 @@ single function to return. Syntax: -``` {.ebnf} -operation ::= +``` +operation ::= (ssa-id `=`)? `call` symbol-ref-id `(` ssa-use-list? `)` `:` function-type ``` @@ -94,7 +94,7 @@ encoded as a function attribute named "callee". Example: -```mlir {.mlir} +```mlir // Calling the function my_add. %31 = call @my_add(%0, %1) : (tensor<16xf32>, tensor<16xf32>) -> tensor<16xf32> ``` @@ -103,7 +103,7 @@ Example: Syntax: -``` {.ebnf} +``` operation ::= `call_indirect` ssa-use `(` ssa-use-list? `)` `:` function-type ``` @@ -117,7 +117,7 @@ Function values can be created with the Example: -```mlir {.mlir} +```mlir %31 = call_indirect %15(%0, %1) : (tensor<16xf32>, tensor<16xf32>) -> tensor<16xf32> ``` @@ -126,7 +126,7 @@ Example: Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `dim` ssa-id `,` integer-literal `:` type ``` @@ -139,7 +139,7 @@ The `dim` operation is represented with a single integer attribute named Examples: -```mlir {.mlir} +```mlir // Always returns 4, can be constant folded: %x = dim %A, 0 : tensor<4 x ? x f32> @@ -157,7 +157,7 @@ Examples: Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `alloc` dim-and-symbol-use-list `:` memref-type ``` @@ -172,7 +172,7 @@ destroyed by the `dealloc` operation. Example: -```mlir {.mlir} +```mlir // Allocating memref for a fully static shape. %A = alloc() : memref<1024x64xf32, #layout_map0, memspace0> @@ -186,7 +186,7 @@ Example: Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `alloc_static` `(` integer-literal `)` : memref-type ``` @@ -198,7 +198,7 @@ require dynamic symbols in their layout function (use the Example: -```mlir {.mlir} +```mlir %A = alloc_static(0x1232a00) : memref<1024 x 64 x f32, #layout_map0, memspace0> ``` @@ -209,7 +209,7 @@ has been performed. Syntax: -``` {.ebnf} +``` operation ::= `dealloc` ssa-use `:` memref-type ``` @@ -219,7 +219,7 @@ allocation. It is paired with an [`alloc`](#alloc-operation) or Example: -```mlir {.mlir} +```mlir dealloc %A : memref<128 x f32, #layout, memspace0> ``` @@ -227,7 +227,7 @@ 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)? @@ -257,7 +257,7 @@ specified as shown below. Example: -```mlir {.mlir} +```mlir %size = constant 32 : index %tag = alloc() : memref<1 x i32, (d0) -> (d0), 4> %idx = constant 0 : index @@ -283,7 +283,7 @@ load/store indices. Example: -```mlir {.mlir} +```mlir dma_wait %tag[%idx], %size : memref<1 x i32, (d0) -> (d0), 4> ``` @@ -304,7 +304,7 @@ then 3 indices are required for the extract. The indices should all be of 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> @@ -314,7 +314,7 @@ Examples: Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `load` ssa-use `[` ssa-use-list `]` `:` memref-type ``` @@ -333,7 +333,7 @@ values or the recursively result of such an `affine.apply` operation. 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> @@ -356,7 +356,7 @@ in these contexts. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `splat` ssa-use `:` ( vector-type | tensor-type ) ``` @@ -366,14 +366,14 @@ it has to be statically shaped. 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` @@ -382,7 +382,7 @@ sizes of the two dynamic dimensions. %m = "foo"() : () -> (index) %n = "bar"() : Syntax: -``` {.ebnf} +``` operation ::= `store` ssa-use `,` ssa-use `[` ssa-use-list `]` `:` memref-type ``` @@ -400,7 +400,7 @@ of such an `affine.apply` operation. Example: -```mlir {.mlir} +```mlir store %100, %A[%1, 1023] : memref<4x?xf32, #layout, memspace0> ``` @@ -417,7 +417,7 @@ in these contexts. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `tensor_load` ssa-use-and-type ``` @@ -427,7 +427,7 @@ operand. Example: -```mlir {.mlir} +```mlir // Produces a value of tensor<4x?xf32> type. %12 = tensor_load %10 : memref<4x?xf32, #layout, memspace0> ``` @@ -436,7 +436,7 @@ Example: Syntax: -``` {.ebnf} +``` operation ::= `tensor_store` ssa-use `,` ssa-use `:` memref-type ``` @@ -446,7 +446,7 @@ element types of these must match, and are specified by the 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> @@ -458,13 +458,13 @@ 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 @@ -484,13 +484,13 @@ attributes. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `ceilf` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar ceiling value. %a = ceilf %b : f64 @@ -510,13 +510,13 @@ has no standard attributes. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `cos` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar cosine value. %a = cos %b : f64 @@ -536,13 +536,13 @@ attributes. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `exp` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar natural exponential. %a = exp %b : f64 @@ -561,13 +561,13 @@ tensor of floats. It has no standard attributes. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `negf` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar negation value. %a = negf %b : f64 @@ -587,13 +587,13 @@ has no standard attributes. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `tanh` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar hyperbolic tangent value. %a = tanh %b : f64 @@ -618,13 +618,13 @@ section. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `addi` ssa-use `,` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar addition. %a = addi %b, %c : i64 @@ -644,13 +644,13 @@ attributes. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `addf` ssa-use `,` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar addition. %a = addf %b, %c : f64 @@ -676,13 +676,13 @@ Bitwise integer and. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `and` ssa-use `,` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar integer bitwise and. %a = and %b, %c : i64 @@ -702,13 +702,13 @@ attributes. 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 @@ -772,7 +772,7 @@ positives Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `constant` attribute-value `:` type ``` @@ -786,7 +786,7 @@ The type specifies the result type of the operation. Examples: -```mlir {.mlir} +```mlir // Integer constant %1 = constant 42 : i32 @@ -809,13 +809,13 @@ function simplifies this Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `copysign` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar copysign value. %a = copysign %b %c : f64 @@ -841,13 +841,13 @@ value divided by -1) is TBD; do NOT assume any specific behavior. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `divis` ssa-use `,` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar signed integer division. %a = divis %b, %c : i64 @@ -874,13 +874,13 @@ behavior. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `diviu` ssa-use `,` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar unsigned integer division. %a = diviu %b, %c : i64 @@ -900,13 +900,13 @@ standard attributes. 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 @@ -932,13 +932,13 @@ if converting to a mismatching static rank. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `mulf` ssa-use `,` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar multiplication. %a = mulf %b, %c : f64 @@ -964,13 +964,13 @@ Bitwise integer or. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `or` ssa-use `,` ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar integer bitwise or. %a = or %b, %c : i64 @@ -996,13 +996,13 @@ behavior. 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 @@ -1028,13 +1028,13 @@ behavior. 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 @@ -1054,13 +1054,13 @@ standard attributes. 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 @@ -1088,13 +1088,13 @@ implement `min` and `max` with signed or unsigned comparison semantics. 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 %2 = tensor_cast %1 : tensor<*xf32> to tensor @@ -1119,13 +1119,13 @@ Bitwise integer xor. Syntax: -``` {.ebnf} +``` operation ::= ssa-id `=` `xor` ssa-use, ssa-use `:` type ``` Examples: -```mlir {.mlir} +```mlir // Scalar integer bitwise xor. %a = xor %b, %c : i64 diff --git a/mlir/g3doc/EDSC.md b/mlir/g3doc/EDSC.md index 264dcaf..afceac2 100644 --- a/mlir/g3doc/EDSC.md +++ b/mlir/g3doc/EDSC.md @@ -97,7 +97,7 @@ def AddOp : Op<"x.add">, 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 { @@ -111,7 +111,7 @@ generated IR resembles the following, for a 4-D memref of `vector<4xi8>`: or the following, for a 0-D `memref`: -``` {.mlir} +``` // CHECK-LABEL: func @t3(%lhs: memref, %rhs: memref, %result: memref) -> () { // CHECK: {{.*}} = load %arg1[] : memref // CHECK: {{.*}} = load %arg0[] : memref diff --git a/mlir/g3doc/LangRef.md b/mlir/g3doc/LangRef.md index a810330..cd6d331 100644 --- a/mlir/g3doc/LangRef.md +++ b/mlir/g3doc/LangRef.md @@ -42,7 +42,7 @@ even arbitrary user-defined high-level operations including the 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. @@ -115,7 +115,7 @@ This document describes the grammar using 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. @@ -127,7 +127,7 @@ literal ::= `abcd` // Matches the literal `abcd`. 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` @@ -137,7 +137,7 @@ 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] @@ -158,7 +158,7 @@ starting with a `//` and going until the end of the line. Syntax: -``` {.ebnf} +``` // Identifiers bare-id ::= (letter|[_]) (letter|digit|[_$.])* bare-id-list ::= bare-id (`,` bare-id)* @@ -226,7 +226,7 @@ with an "llvm." name. 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) ``` @@ -238,7 +238,7 @@ GPUs), and are required to align with the LLVM definition of these intrinsics. Syntax: -``` {.ebnf} +``` operation ::= op-result-list? (generic-operation | custom-operation) trailing-location? generic-operation ::= string-literal '(' ssa-use-list? ')' attribute-dict? @@ -270,7 +270,7 @@ types of the results and operands. Example: -```mlir {.mlir} +```mlir // An operation that produces two results. // The results of %result can be accessed via the `#` syntax. %result:2 = "foo_div"() : () -> (f32, i32) @@ -295,7 +295,7 @@ also have a list of successors ([blocks](#blocks) and their arguments). 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) -> () @@ -303,7 +303,7 @@ Example: ### Module -``` {.ebnf} +``` module ::= `module` symbol-ref-id? (`attributes` attribute-dict)? region ``` @@ -321,7 +321,7 @@ outside of the function, and all external references must use function arguments 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 `)` @@ -352,7 +352,7 @@ function arguments, results, or the function itself. Examples: -```mlir {.mlir} +```mlir // External function definitions. func @abort() func @scribble(i32, i64, memref) -> f64 @@ -377,7 +377,7 @@ func @example_fn_attr() attributes {dialectName.attrName = false} Syntax: -``` {.ebnf} +``` block ::= block-label operation+ block-label ::= block-id block-arg-list? `:` block-id ::= caret-id @@ -402,7 +402,7 @@ provided for these block arguments by branches that go to the block. 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 @@ -447,7 +447,7 @@ attributes. 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* `}` ``` @@ -468,7 +468,7 @@ would have been legal to use them as operands to the enclosing operation. 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 @@ -548,7 +548,7 @@ MLIR has an open type system (i.e. there is no fixed list of types), and types 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)* @@ -564,7 +564,7 @@ ssa-use-and-type-list ::= ssa-use-and-type (`,` ssa-use-and-type)* ### Type Aliases -``` {.ebnf} +``` type-alias-def ::= '!' alias-name '=' 'type' type type-alias ::= '!' alias-name ``` @@ -576,7 +576,7 @@ names are reserved for [dialect types](#dialect-types). Example: -```mlir {.mlir} +```mlir !avx_m128 = type vector<4 x f32> // Using the original type. @@ -591,7 +591,7 @@ Example: Similarly to operations, dialects may define custom extensions to the type system. -``` {.ebnf} +``` dialect-namespace ::= bare-id opaque-dialect-item ::= dialect-namespace '<' string-literal '>' @@ -613,7 +613,7 @@ dialect-type ::= '!' pretty-dialect-item 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*"> @@ -630,7 +630,7 @@ Dialect types can be specified in a verbose form, e.g. like this: 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 @@ -651,7 +651,7 @@ See [here](DefiningAttributesAndTypes.md) to learn how to define dialect types. 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 @@ -668,7 +668,7 @@ standard-type ::= complex-type Syntax: -``` {.ebnf} +``` complex-type ::= `complex` `<` type `>` ``` @@ -678,7 +678,7 @@ type. The element must be a floating point or integer scalar type. Examples: -```mlir {.mlir} +```mlir complex complex ``` @@ -687,7 +687,7 @@ complex Syntax: -``` {.ebnf} +``` // Floating point. float-type ::= `f16` | `bf16` | `f32` | `f64` ``` @@ -699,7 +699,7 @@ above. Syntax: -``` {.ebnf} +``` // MLIR functions can return multiple values. function-result-type ::= type-list-parens | non-function-type @@ -721,7 +721,7 @@ Function types are also used to indicate the arguments and results of Syntax: -``` {.ebnf} +``` // Target word-sized integer. index-type ::= `index` ``` @@ -739,7 +739,7 @@ sizes, dimensionalities and subscripts. Syntax: -``` {.ebnf} +``` // Sized integers like i1, i4, i8, i16, i32. integer-type ::= `i` [1-9][0-9]* ``` @@ -759,7 +759,7 @@ TODO: Need to decide on a representation for quantized integers Syntax: -``` {.ebnf} +``` memref-type ::= ranked-memref-type | unranked-memref-type @@ -811,14 +811,10 @@ exposed to codegen but one may query the rank of an unranked memref (a special 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> @@ -852,7 +848,7 @@ and index maps. Examples of memref static type -```mlir {.mlir} +```mlir // Identity index/layout map #identity = (d0, d1) -> (d0, d1) @@ -920,7 +916,7 @@ multidimensional index space defined by the memref's dimension list. 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> @@ -959,7 +955,7 @@ maps, `memref`. Layout map examples: -```mlir {.mlir} +```mlir // MxN matrix stored in row major layout in memory: #layout_map_row_major = (i, j) -> (i, j) @@ -1024,7 +1020,7 @@ normalized memref descriptor when lowering to LLVM. Syntax: -``` {.ebnf} +``` none-type ::= `none` ``` @@ -1035,7 +1031,7 @@ where its value does not have a defined dynamic representation. Syntax: -``` {.ebnf} +``` tensor-type ::= `tensor` `<` dimension-list tensor-memref-element-type `>` tensor-memref-element-type ::= vector-element-type | vector-type | complex-type @@ -1066,7 +1062,7 @@ types, such tensors should be optimized away before lowering tensors to vectors. Examples: -```mlir {.mlir} +```mlir // Tensor with unknown rank. tensor<* x f32> @@ -1093,7 +1089,7 @@ tensor<0xf32> Syntax: -``` {.ebnf} +``` tuple-type ::= `tuple` `<` (type ( `,` type)*)? `>` ``` @@ -1106,7 +1102,7 @@ no standard operations for operating on `tuple` types Examples: -```mlir {.mlir} +```mlir // Empty tuple. tuple<> @@ -1121,7 +1117,7 @@ tuple, i5> Syntax: -``` {.ebnf} +``` vector-type ::= `vector` `<` static-dimension-list vector-element-type `>` vector-element-type ::= float-type | integer-type @@ -1143,7 +1139,7 @@ shape `(0, 42)` and zero shapes are not allowed. Syntax: -``` {.ebnf} +``` attribute-dict ::= `{` `}` | `{` attribute-entry (`,` attribute-entry)* `}` attribute-entry ::= dialect-attribute-entry | dependent-attribute-entry @@ -1170,13 +1166,13 @@ dialect and not the function argument. 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 ``` @@ -1189,7 +1185,7 @@ These aliases *must* be defined before their uses. Alias names may not contain a Example: -```mlir {.mlir} +```mlir #map = (d0) -> (d0 + 10) // Using the original attribute. @@ -1206,14 +1202,14 @@ syntactic structure of these values is identical to custom dialect type values, 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"> @@ -1224,7 +1220,7 @@ Dialect attributes can be specified in a verbose form, e.g. like this: 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 ``` @@ -1244,7 +1240,7 @@ Standard attributes are a core set of [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 @@ -1263,7 +1259,7 @@ standard-attribute ::= affine-map-attribute Syntax: -``` {.ebnf} +``` affine-map-attribute ::= affine-map ``` @@ -1273,7 +1269,7 @@ An affine-map attribute is an attribute that represents a affine-map object. Syntax: -``` {.ebnf} +``` array-attribute ::= `[` (attribute-value (`,` attribute-value)*)? `]` ``` @@ -1284,7 +1280,7 @@ values. Syntax: -``` {.ebnf} +``` bool-attribute ::= bool-literal ``` @@ -1295,7 +1291,7 @@ value, true or false. Syntax: -``` {.ebnf} +``` dictionary-attribute ::= `{` (attribute-entry (`,` attribute-entry)*)? `}` ``` @@ -1307,7 +1303,7 @@ unique within the collection. Syntax: -``` {.ebnf} +``` elements-attribute ::= dense-elements-attribute | opaque-elements-attribute | sparse-elements-attribute @@ -1320,7 +1316,7 @@ An elements attribute is a literal attribute that represents a constant Syntax: -``` {.ebnf} +``` dense-elements-attribute ::= `dense` `<` attribute-value `>` `:` ( tensor-type | vector-type ) ``` @@ -1334,7 +1330,7 @@ floating point type. Syntax: -``` {.ebnf} +``` opaque-elements-attribute ::= `opaque` `<` dialect-namespace `,` hex-string-literal `>` `:` ( tensor-type | vector-type ) @@ -1351,7 +1347,7 @@ Note: The parsed string literal must be in hexadecimal form. Syntax: -``` {.ebnf} +``` sparse-elements-attribute ::= `sparse` `<` attribute-value `,` attribute-value `>` `:` ( tensor-type | vector-type ) ``` @@ -1368,7 +1364,7 @@ corresponding values for the indices. Example: -```mlir {.mlir} +```mlir sparse<[[0, 0], [1, 2]], [1, 5]> : tensor<3x4xi32> // This represents the following tensor: @@ -1381,7 +1377,7 @@ Example: Syntax: -``` {.ebnf} +``` float-attribute ::= (float-literal (`:` float-type)?) | (hexadecimal-literal `:` float-type) ``` @@ -1396,7 +1392,7 @@ attribute. Examples: -``` {.mlir} +``` 42.0 // float attribute defaults to f64 type 42.0 : f32 // float attribute of f32 type 0x7C00 : f16 // positive infinity @@ -1408,7 +1404,7 @@ Examples: Syntax: -``` {.ebnf} +``` integer-attribute ::= integer-literal ( `:` (index-type | integer-type) )? ``` @@ -1420,7 +1416,7 @@ is not specified, is a 64-bit integer. Syntax: -``` {.ebnf} +``` integer-set-attribute ::= affine-map ``` @@ -1430,7 +1426,7 @@ An integer-set attribute is an attribute that represents an integer-set object. Syntax: -``` {.ebnf} +``` string-attribute ::= string-literal (`:` type)? ``` @@ -1440,7 +1436,7 @@ A string attribute is an attribute that represents a string literal value. Syntax: -``` {.ebnf} +``` symbol-ref-attribute ::= symbol-ref-id (`::` symbol-ref-id)* ``` @@ -1468,7 +1464,7 @@ reference, we can always opaquely reason about a symbols usage characteristics. Syntax: -``` {.ebnf} +``` type-attribute ::= type ``` @@ -1476,7 +1472,7 @@ A type attribute is an attribute that represents a [type object](#type-system). #### Unit Attribute -``` {.ebnf} +``` unit-attribute ::= `unit` ``` @@ -1490,7 +1486,7 @@ could be represented as a [boolean attribute](#boolean-attribute)(true or 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} diff --git a/mlir/g3doc/MLIRForGraphAlgorithms.md b/mlir/g3doc/MLIRForGraphAlgorithms.md index 9130d03..ac26e5b 100644 --- a/mlir/g3doc/MLIRForGraphAlgorithms.md +++ b/mlir/g3doc/MLIRForGraphAlgorithms.md @@ -207,7 +207,7 @@ test the dependence analysis infra in the code generator, Andy Davis wrote a 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> @@ -241,7 +241,7 @@ with ShapeRefiner. 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> ``` diff --git a/mlir/g3doc/Rationale.md b/mlir/g3doc/Rationale.md index efccf07..66cf800 100644 --- a/mlir/g3doc/Rationale.md +++ b/mlir/g3doc/Rationale.md @@ -134,7 +134,7 @@ unknown dimension can be queried using the "dim" builtin as shown below. Example: -```mlir {.mlir} +```mlir func foo(...) { %A = alloc <8x?xf32, #lmap> (%N) ... @@ -532,7 +532,7 @@ the existing LLVM type parsing infrastructure. Example: -```mlir {.mlir} +```mlir %s = "foo"() : () -> !llvm<"i32*"> ``` @@ -588,7 +588,7 @@ represents computation. ### 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++) { @@ -604,7 +604,7 @@ for (i = 0; i < N; i++) { 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, %key : i32) { %ni = dim %A, 0 : memref // This loop can be parallelized @@ -662,7 +662,7 @@ for (i = 0; i < N; i++)        } ``` -```mlir {.mlir} +```mlir func @outer_nest(%n : index) { affine.for %i = 0 to %n { affine.for %j = 0 to %n { @@ -689,7 +689,7 @@ The following example illustrates a reference implementation of a 2D 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) @@ -805,7 +805,7 @@ which is called a schedule tree. Each non-leaf node of the tree is an abstract 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) @@ -859,7 +859,7 @@ constraints on the identifiers. Syntax: -``` {.ebnf} +``` // Affine relation definition at the top of file affine-rel-def ::= affine-rel-id `=` affine-relation-inline @@ -890,7 +890,7 @@ dimensions. Example: -```mlir {.mlir} +```mlir // read relation: two elements ( d0 <= r0 <= d0+1 ) ##aff_rel9 = (d0) -> (r0) : r0 - d0 >= 0, d0 - r0 + 1 >= 0 @@ -966,7 +966,7 @@ TODO: Design this, and update to use function attribute syntax. 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>) @@ -1027,7 +1027,7 @@ The abandoned design of supporting escaping scalars is as follows: Syntax: -``` {.ebnf} +``` [ =] for % = ... step [with ] { } @@ -1040,7 +1040,7 @@ is a list of instructions that may also include a yield instruction. Example: -```mlir {.mlir} +```mlir // Return sum of elements in 1-dimensional mref A func i32 @sum(%A : memref, %N : i32) -> (i32) { %init = 0 @@ -1057,7 +1057,7 @@ func i32 @sum(%A : memref, %N : i32) -> (i32) { Syntax: -``` {.ebnf} +``` = affine.if () {...} [else {...}] ``` @@ -1070,7 +1070,7 @@ this situation. Example: -```mlir {.mlir} +```mlir // Compute sum of half of the array func i32 @sum_half(%A : memref, %N : i32) -> (i32) { %s0 = 0 diff --git a/mlir/g3doc/TestingGuide.md b/mlir/g3doc/TestingGuide.md index 5a8c039..723b78b 100644 --- a/mlir/g3doc/TestingGuide.md +++ b/mlir/g3doc/TestingGuide.md @@ -27,7 +27,7 @@ different aspects of the IR - such as the output of a transformation pass. An example FileCheck test is shown below: -```mlir {.mlir} +```mlir // RUN: mlir-opt %s -cse | FileCheck %s // CHECK-LABEL: func @simple_constant @@ -52,7 +52,7 @@ brittle tests that are essentially `diff` tests. FileCheck tests should be as 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) @@ -89,7 +89,7 @@ IR output. 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 @@ -111,7 +111,7 @@ as well as named 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 @@ -140,7 +140,7 @@ accessible via the `verify-diagnostics` flag in mlir-opt. 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. diff --git a/mlir/g3doc/Tutorials/Toy/Ch-7.md b/mlir/g3doc/Tutorials/Toy/Ch-7.md index a55c583..398983a 100644 --- a/mlir/g3doc/Tutorials/Toy/Ch-7.md +++ b/mlir/g3doc/Tutorials/Toy/Ch-7.md @@ -254,7 +254,7 @@ form available under certain circumstances. The responsibility of our `Toy` 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)* `>` ``` diff --git a/mlir/include/mlir/Dialect/GPU/GPUOps.td b/mlir/include/mlir/Dialect/GPU/GPUOps.td index a1f4be6..dfb7ea9 100644 --- a/mlir/include/mlir/Dialect/GPU/GPUOps.td +++ b/mlir/include/mlir/Dialect/GPU/GPUOps.td @@ -68,7 +68,7 @@ def GPU_GPUFuncOp : GPU_Op<"func", [FunctionLike, IsolatedFromAbove, Symbol]> { Syntax: - ``` {.ebnf} + ``` op ::= `gpu.func` symbol-ref-id `(` argument-list `)` (`->` function-result-list)? memory-attribution `kernel`? function-attributes? region @@ -213,7 +213,7 @@ def GPU_LaunchOp : GPU_Op<"launch", [IsolatedFromAbove]>, Syntax: - ``` {.ebnf} + ``` operation ::= `gpu.launch` `block` `(` ssa-id-list `)` `in` ssa-reassignment `threads` `(` ssa-id-list `)` `in` ssa-reassignment (`args` ssa-reassignment `:` type-list)? @@ -223,7 +223,7 @@ def GPU_LaunchOp : GPU_Op<"launch", [IsolatedFromAbove]>, 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 { diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVArithmeticOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVArithmeticOps.td index 38b6391..f15d274 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVArithmeticOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVArithmeticOps.td @@ -53,7 +53,7 @@ def SPV_FAddOp : SPV_ArithmeticBinaryOp<"FAdd", SPV_Float, [Commutative]> { 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 @@ -83,7 +83,7 @@ def SPV_FDivOp : SPV_ArithmeticBinaryOp<"FDiv", SPV_Float, []> { 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 @@ -118,7 +118,7 @@ def SPV_FModOp : SPV_ArithmeticBinaryOp<"FMod", SPV_Float, []> { 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 @@ -148,7 +148,7 @@ def SPV_FMulOp : SPV_ArithmeticBinaryOp<"FMul", SPV_Float, [Commutative]> { ### Custom assembly form - ``` {.ebnf} + ``` float-scalar-vector-type ::= float-type | `vector<` integer-literal `x` float-type `>` fmul-op ::= `spv.FMul` ssa-use, ssa-use @@ -178,7 +178,7 @@ def SPV_FNegateOp : SPV_ArithmeticUnaryOp<"FNegate", SPV_Float, []> { ### 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 @@ -212,7 +212,7 @@ def SPV_FRemOp : SPV_ArithmeticBinaryOp<"FRem", SPV_Float, []> { 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 @@ -242,7 +242,7 @@ def SPV_FSubOp : SPV_ArithmeticBinaryOp<"FSub", SPV_Float, []> { 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 @@ -277,7 +277,7 @@ def SPV_IAddOp : SPV_ArithmeticBinaryOp<"IAdd", SPV_Integer, [Commutative]> { 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 @@ -315,7 +315,7 @@ def SPV_IMulOp : SPV_ArithmeticBinaryOp<"IMul", SPV_Integer, [Commutative]> { 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 @@ -353,7 +353,7 @@ def SPV_ISubOp : SPV_ArithmeticBinaryOp<"ISub", SPV_Integer, []> { 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 @@ -388,7 +388,7 @@ def SPV_SDivOp : SPV_ArithmeticBinaryOp<"SDiv", SPV_Integer, []> { 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 @@ -426,7 +426,7 @@ def SPV_SModOp : SPV_ArithmeticBinaryOp<"SMod", SPV_Integer, []> { 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 @@ -463,7 +463,7 @@ def SPV_SRemOp : SPV_ArithmeticBinaryOp<"SRem", SPV_Integer, []> { 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 @@ -495,7 +495,7 @@ def SPV_UDivOp : SPV_ArithmeticBinaryOp<"UDiv", SPV_Integer, []> { 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 @@ -527,7 +527,7 @@ def SPV_UModOp : SPV_ArithmeticBinaryOp<"UMod", SPV_Integer> { 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 diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVAtomicOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVAtomicOps.td index 7042bf2..d8e62bb 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVAtomicOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVAtomicOps.td @@ -35,7 +35,7 @@ def SPV_AtomicCompareExchangeWeakOp : SPV_Op<"AtomicCompareExchangeWeak", []> { ### Custom assembly form - ``` {.ebnf} + ``` scope ::= `"CrossDevice"` | `"Device"` | `"Workgroup"` | ... memory-semantics ::= `"None"` | `"Acquire"` | "Release"` | ... diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVBitOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVBitOps.td index e927f31..d76a1e3 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVBitOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVBitOps.td @@ -83,7 +83,7 @@ def SPV_BitCountOp : SPV_BitUnaryOp<"BitCount", []> { ### Custom assembly form - ``` {.ebnf} + ``` integer-scalar-vector-type ::= integer-type | `vector<` integer-literal `x` integer-type `>` bitcount-op ::= ssa-id `=` `spv.BitCount` ssa-use @@ -132,7 +132,7 @@ def SPV_BitFieldInsertOp : SPV_Op<"BitFieldInsert", [NoSideEffect]> { ### 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 @@ -190,7 +190,7 @@ def SPV_BitFieldSExtractOp : SPV_BitFieldExtractOp<"BitFieldSExtract", []> { ### 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 @@ -219,7 +219,7 @@ def SPV_BitFieldUExtractOp : SPV_BitFieldExtractOp<"BitFieldUExtract", []> { ### 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 @@ -253,7 +253,7 @@ def SPV_BitReverseOp : SPV_BitUnaryOp<"BitReverse", []> { ### Custom assembly form - ``` {.ebnf} + ``` integer-scalar-vector-type ::= integer-type | `vector<` integer-literal `x` integer-type `>` bitreverse-op ::= ssa-id `=` `spv.BitReverse` ssa-use @@ -287,7 +287,7 @@ def SPV_BitwiseAndOp : SPV_BitBinaryOp<"BitwiseAnd", [Commutative]> { ### 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 @@ -321,7 +321,7 @@ def SPV_BitwiseOrOp : SPV_BitBinaryOp<"BitwiseOr", [Commutative]> { ### 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 @@ -355,7 +355,7 @@ def SPV_BitwiseXorOp : SPV_BitBinaryOp<"BitwiseXor", [Commutative]> { ### 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 @@ -397,7 +397,7 @@ def SPV_ShiftLeftLogicalOp : SPV_ShiftOp<"ShiftLeftLogical", []> { ### Custom assembly form - ``` {.ebnf} + ``` integer-scalar-vector-type ::= integer-type | `vector<` integer-literal `x` integer-type `>` shift-left-logical-op ::= ssa-id `=` `spv.ShiftLeftLogical` @@ -438,7 +438,7 @@ def SPV_ShiftRightArithmeticOp : SPV_ShiftOp<"ShiftRightArithmetic", []> { ### Custom assembly form - ``` {.ebnf} + ``` integer-scalar-vector-type ::= integer-type | `vector<` integer-literal `x` integer-type `>` shift-right-arithmetic-op ::= ssa-id `=` `spv.ShiftRightArithmetic` @@ -480,7 +480,7 @@ def SPV_ShiftRightLogicalOp : SPV_ShiftOp<"ShiftRightLogical", []> { ### Custom assembly form - ``` {.ebnf} + ``` integer-scalar-vector-type ::= integer-type | `vector<` integer-literal `x` integer-type `>` shift-right-logical-op ::= ssa-id `=` `spv.ShiftRightLogical` @@ -514,7 +514,7 @@ def SPV_NotOp : SPV_BitUnaryOp<"Not", []> { ### 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 diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVCastOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVCastOps.td index 2faffb6..e4fe526 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVCastOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVCastOps.td @@ -72,7 +72,7 @@ def SPV_BitcastOp : SPV_Op<"Bitcast", [NoSideEffect]> { ### Custom assembly form - ``` {.ebnf} + ``` bitcast-op ::= ssa-id `=` `spv.Bitcast` ssa-use `:` operand-type `to` result-type ``` @@ -118,7 +118,7 @@ def SPV_ConvertFToSOp : SPV_CastOp<"ConvertFToS", SPV_Integer, SPV_Float, []> { ### Custom assembly form - ``` {.ebnf} + ``` convert-f-to-s-op ::= ssa-id `=` `spv.ConvertFToSOp` ssa-use `:` operand-type `to` result-type ``` @@ -151,7 +151,7 @@ def SPV_ConvertFToUOp : SPV_CastOp<"ConvertFToU", SPV_Integer, SPV_Float, []> { ### Custom assembly form - ``` {.ebnf} + ``` convert-f-to-u-op ::= ssa-id `=` `spv.ConvertFToUOp` ssa-use `:` operand-type `to` result-type ``` @@ -182,7 +182,7 @@ def SPV_ConvertSToFOp : SPV_CastOp<"ConvertSToF", SPV_Float, SPV_Integer, []> { ### Custom assembly form - ``` {.ebnf} + ``` convert-s-to-f-op ::= ssa-id `=` `spv.ConvertSToFOp` ssa-use `:` operand-type `to` result-type ``` @@ -213,7 +213,7 @@ def SPV_ConvertUToFOp : SPV_CastOp<"ConvertUToF", SPV_Float, SPV_Integer, []> { ### Custom assembly form - ``` {.ebnf} + ``` convert-u-to-f-op ::= ssa-id `=` `spv.ConvertUToFOp` ssa-use `:` operand-type `to` result-type ``` @@ -246,7 +246,7 @@ def SPV_FConvertOp : SPV_CastOp<"FConvert", SPV_Float, SPV_Float, []> { ### Custom assembly form - ``` {.ebnf} + ``` f-convert-op ::= ssa-id `=` `spv.FConvertOp` ssa-use `:` operand-type `to` result-type ``` @@ -280,7 +280,7 @@ def SPV_SConvertOp : SPV_CastOp<"SConvert", SPV_Integer, SPV_Integer, []> { ### Custom assembly form - ``` {.ebnf} + ``` s-convert-op ::= ssa-id `=` `spv.SConvertOp` ssa-use `:` operand-type `to` result-type ``` @@ -315,7 +315,7 @@ def SPV_UConvertOp : SPV_CastOp<"UConvert", SPV_Integer, SPV_Integer, []> { ### Custom assembly form - ``` {.ebnf} + ``` u-convert-op ::= ssa-id `=` `spv.UConvertOp` ssa-use `:` operand-type `to` result-type ``` diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVCompositeOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVCompositeOps.td index 6392a1b..3a21ddc 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVCompositeOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVCompositeOps.td @@ -54,7 +54,7 @@ def SPV_CompositeConstructOp : SPV_Op<"CompositeConstruct", [NoSideEffect]> { ### Custom assembly form - ``` {.ebnf} + ``` composite-construct-op ::= ssa-id `=` `spv.CompositeConstruct` (ssa-use (`,` ssa-use)* )? `:` composite-type ``` @@ -93,7 +93,7 @@ def SPV_CompositeExtractOp : SPV_Op<"CompositeExtract", [NoSideEffect]> { ### Custom assembly form - ``` {.ebnf} + ``` composite-extract-op ::= ssa-id `=` `spv.CompositeExtract` ssa-use `[` integer-literal (',' integer-literal)* `]` `:` composite-type @@ -143,7 +143,7 @@ def SPV_CompositeInsertOp : SPV_Op<"CompositeInsert", [NoSideEffect]> { ### Custom assembly form - ``` {.ebnf} + ``` composite-insert-op ::= ssa-id `=` `spv.CompositeInsert` ssa-use, ssa-use `[` integer-literal (',' integer-literal)* `]` `:` object-type `into` composite-type diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVControlFlowOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVControlFlowOps.td index ab9d51b..464b670 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVControlFlowOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVControlFlowOps.td @@ -36,7 +36,7 @@ def SPV_BranchOp : SPV_Op<"Branch", [InFunctionScope, Terminator]> { ### 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 `)` @@ -106,7 +106,7 @@ def SPV_BranchConditionalOp : SPV_Op<"BranchConditional", ### Custom assembly form - ``` {.ebnf} + ``` branch-conditional-op ::= `spv.BranchConditional` ssa-use (`[` integer-literal, integer-literal `]`)? `,` successor `,` successor @@ -231,7 +231,7 @@ def SPV_FunctionCallOp : SPV_Op<"FunctionCall", [ ### Custom assembly form - ``` {.ebnf} + ``` function-call-op ::= `spv.FunctionCall` function-id `(` ssa-use-list `)` `:` function-type ``` @@ -355,7 +355,7 @@ def SPV_ReturnOp : SPV_Op<"Return", [InFunctionScope, Terminator]> { ### Custom assembly form - ``` {.ebnf} + ``` return-op ::= `spv.Return` ``` }]; @@ -378,7 +378,7 @@ def SPV_UnreachableOp : SPV_Op<"Unreachable", [InFunctionScope, Terminator]> { ### Custom assembly form - ``` {.ebnf} + ``` unreachable-op ::= `spv.Unreachable` ``` }]; @@ -405,7 +405,7 @@ def SPV_ReturnValueOp : SPV_Op<"ReturnValue", [InFunctionScope, Terminator]> { ### Custom assembly form - ``` {.ebnf} + ``` return-value-op ::= `spv.ReturnValue` ssa-use `:` spirv-type ``` diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVGLSLOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVGLSLOps.td index 2a1e8f3..a031fac 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVGLSLOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVGLSLOps.td @@ -100,7 +100,7 @@ def SPV_GLSLFAbsOp : SPV_GLSLUnaryArithmeticOp<"FAbs", 4, SPV_Float> { 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 `:` @@ -129,7 +129,7 @@ def SPV_GLSLSAbsOp : SPV_GLSLUnaryArithmeticOp<"SAbs", 5, SPV_Integer> { 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 `:` @@ -160,7 +160,7 @@ def SPV_GLSLCeilOp : SPV_GLSLUnaryArithmeticOp<"Ceil", 9, SPV_Float> { 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 `:` @@ -190,7 +190,7 @@ def SPV_GLSLCosOp : SPV_GLSLUnaryArithmeticOp<"Cos", 14, SPV_Float16or32> { per component. ### Custom assembly format - ``` {.ebnf} + ``` restricted-float-scalar-type ::= `f16` | `f32` restricted-float-scalar-vector-type ::= restricted-float-scalar-type | @@ -222,7 +222,7 @@ def SPV_GLSLExpOp : SPV_GLSLUnaryArithmeticOp<"Exp", 27, SPV_Float16or32> { computed per component."; ### Custom assembly format - ``` {.ebnf} + ``` restricted-float-scalar-type ::= `f16` | `f32` restricted-float-scalar-vector-type ::= restricted-float-scalar-type | @@ -255,7 +255,7 @@ def SPV_GLSLFloorOp : SPV_GLSLUnaryArithmeticOp<"Floor", 8, SPV_Float> { 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 `:` @@ -285,7 +285,7 @@ def SPV_GLSLInverseSqrtOp : SPV_GLSLUnaryArithmeticOp<"InverseSqrt", 32, SPV_Flo 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 `:` @@ -316,7 +316,7 @@ def SPV_GLSLLogOp : SPV_GLSLUnaryArithmeticOp<"Log", 28, SPV_Float16or32> { per component. ### Custom assembly format - ``` {.ebnf} + ``` restricted-float-scalar-type ::= `f16` | `f32` restricted-float-scalar-vector-type ::= restricted-float-scalar-type | @@ -349,7 +349,7 @@ def SPV_GLSLFMaxOp : SPV_GLSLBinaryArithmeticOp<"FMax", 40, SPV_Float> { 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 `:` @@ -379,7 +379,7 @@ def SPV_GLSLSMaxOp : SPV_GLSLBinaryArithmeticOp<"SMax", 42, SPV_Integer> { 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 `:` @@ -410,7 +410,7 @@ def SPV_GLSLFMinOp : SPV_GLSLBinaryArithmeticOp<"FMin", 37, SPV_Float> { 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 `:` @@ -440,7 +440,7 @@ def SPV_GLSLSMinOp : SPV_GLSLBinaryArithmeticOp<"SMin", 39, SPV_Integer> { 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 `:` @@ -470,7 +470,7 @@ def SPV_GLSLFSignOp : SPV_GLSLUnaryArithmeticOp<"FSign", 6, SPV_Float> { 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 `:` @@ -499,7 +499,7 @@ def SPV_GLSLSSignOp : SPV_GLSLUnaryArithmeticOp<"SSign", 7, SPV_Integer> { 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 `:` @@ -529,7 +529,7 @@ def SPV_GLSLSqrtOp : SPV_GLSLUnaryArithmeticOp<"Sqrt", 31, SPV_Float> { 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 `:` @@ -559,7 +559,7 @@ def SPV_GLSLTanhOp : SPV_GLSLUnaryArithmeticOp<"Tanh", 21, SPV_Float16or32> { per component. ### Custom assembly format - ``` {.ebnf} + ``` restricted-float-scalar-type ::= `f16` | `f32` restricted-float-scalar-vector-type ::= restricted-float-scalar-type | diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVGroupOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVGroupOps.td index 5f60e6b..c0388fe 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVGroupOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVGroupOps.td @@ -46,7 +46,7 @@ def SPV_SubgroupBallotKHROp : SPV_Op<"SubgroupBallotKHR", []> { ### Custom assembly form - ``` {.ebnf} + ``` subgroup-ballot-op ::= ssa-id `=` `spv.SubgroupBallotKHR` ssa-use `:` `vector` `<` 4 `x` `i32` `>` ``` diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVLogicalOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVLogicalOps.td index be9825a..0c4b290 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVLogicalOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVLogicalOps.td @@ -62,7 +62,7 @@ def SPV_FOrdEqualOp : SPV_LogicalBinaryOp<"FOrdEqual", SPV_Float, [Commutative]> ### 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 @@ -96,7 +96,7 @@ def SPV_FOrdGreaterThanOp : SPV_LogicalBinaryOp<"FOrdGreaterThan", SPV_Float, [] ### 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 @@ -130,7 +130,7 @@ def SPV_FOrdGreaterThanEqualOp : SPV_LogicalBinaryOp<"FOrdGreaterThanEqual", SPV ### 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 @@ -164,7 +164,7 @@ def SPV_FOrdLessThanOp : SPV_LogicalBinaryOp<"FOrdLessThan", SPV_Float, []> { ### 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 @@ -198,7 +198,7 @@ def SPV_FOrdLessThanEqualOp : SPV_LogicalBinaryOp<"FOrdLessThanEqual", SPV_Float ### 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 @@ -229,7 +229,7 @@ def SPV_FOrdNotEqualOp : SPV_LogicalBinaryOp<"FOrdNotEqual", SPV_Float, [Commuta ### 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 @@ -260,7 +260,7 @@ def SPV_FUnordEqualOp : SPV_LogicalBinaryOp<"FUnordEqual", SPV_Float, [Commutati ### 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 @@ -294,7 +294,7 @@ def SPV_FUnordGreaterThanOp : SPV_LogicalBinaryOp<"FUnordGreaterThan", SPV_Float ### 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 @@ -328,7 +328,7 @@ def SPV_FUnordGreaterThanEqualOp : SPV_LogicalBinaryOp<"FUnordGreaterThanEqual", ### 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 @@ -362,7 +362,7 @@ def SPV_FUnordLessThanOp : SPV_LogicalBinaryOp<"FUnordLessThan", SPV_Float, []> ### 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 @@ -396,7 +396,7 @@ def SPV_FUnordLessThanEqualOp : SPV_LogicalBinaryOp<"FUnordLessThanEqual", SPV_F ### 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 @@ -427,7 +427,7 @@ def SPV_FUnordNotEqualOp : SPV_LogicalBinaryOp<"FUnordNotEqual", SPV_Float, [Com ### 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 @@ -457,7 +457,7 @@ def SPV_IEqualOp : SPV_LogicalBinaryOp<"IEqual", SPV_Integer, [Commutative]> { 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 @@ -488,7 +488,7 @@ def SPV_INotEqualOp : SPV_LogicalBinaryOp<"INotEqual", SPV_Integer, [Commutative 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 @@ -523,7 +523,7 @@ def SPV_LogicalAndOp : SPV_LogicalBinaryOp<"LogicalAnd", SPV_Bool, [Commutative] ### Custom assembly form - ``` {.ebnf} + ``` logical-and ::= `spv.LogicalAnd` ssa-use `,` ssa-use `:` operand-type ``` @@ -556,7 +556,7 @@ def SPV_LogicalEqualOp : SPV_LogicalBinaryOp<"LogicalEqual", SPV_Bool, [Commutat ### Custom assembly form - ``` {.ebnf} + ``` logical-equal ::= `spv.LogicalEqual` ssa-use `,` ssa-use `:` operand-type ``` @@ -586,7 +586,7 @@ def SPV_LogicalNotOp : SPV_LogicalUnaryOp<"LogicalNot", SPV_Bool, []> { ### Custom assembly form - ``` {.ebnf} + ``` logical-not ::= `spv.LogicalNot` ssa-use `:` operand-type ``` @@ -620,7 +620,7 @@ def SPV_LogicalNotEqualOp : SPV_LogicalBinaryOp<"LogicalNotEqual", SPV_Bool, [Co ### Custom assembly form - ``` {.ebnf} + ``` logical-not-equal ::= `spv.LogicalNotEqual` ssa-use `,` ssa-use `:` operand-type ``` @@ -653,7 +653,7 @@ def SPV_LogicalOrOp : SPV_LogicalBinaryOp<"LogicalOr", SPV_Bool, [Commutative]> ### Custom assembly form - ``` {.ebnf} + ``` logical-or ::= `spv.LogicalOr` ssa-use `,` ssa-use `:` operand-type ``` @@ -684,7 +684,7 @@ def SPV_SGreaterThanOp : SPV_LogicalBinaryOp<"SGreaterThan", SPV_Integer, []> { 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 @@ -718,7 +718,7 @@ def SPV_SGreaterThanEqualOp : SPV_LogicalBinaryOp<"SGreaterThanEqual", SPV_Integ 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 @@ -751,7 +751,7 @@ def SPV_SLessThanOp : SPV_LogicalBinaryOp<"SLessThan", SPV_Integer, []> { 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 @@ -785,7 +785,7 @@ def SPV_SLessThanEqualOp : SPV_LogicalBinaryOp<"SLessThanEqual", SPV_Integer, [] 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 @@ -827,7 +827,7 @@ def SPV_SelectOp : SPV_Op<"Select", [NoSideEffect]> { ### Custom assembly form - ``` {.ebnf} + ``` scalar-type ::= integer-type | float-type | boolean-type select-object-type ::= scalar-type | `vector<` integer-literal `x` scalar-type `>` @@ -879,7 +879,7 @@ def SPV_UGreaterThanOp : SPV_LogicalBinaryOp<"UGreaterThan", SPV_Integer, []> { 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 @@ -913,7 +913,7 @@ def SPV_UGreaterThanEqualOp : SPV_LogicalBinaryOp<"UGreaterThanEqual", SPV_Integ 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 @@ -946,7 +946,7 @@ def SPV_ULessThanOp : SPV_LogicalBinaryOp<"ULessThan", SPV_Integer, []> { 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 @@ -981,7 +981,7 @@ def SPV_ULessThanEqualOp : 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 diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVNonUniformOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVNonUniformOps.td index a37f5b5..1b3174c 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVNonUniformOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVNonUniformOps.td @@ -49,7 +49,7 @@ def SPV_GroupNonUniformBallotOp : SPV_Op<"GroupNonUniformBallot", []> { ### Custom assembly form - ``` {.ebnf} + ``` scope ::= `"Workgroup"` | `"Subgroup"` non-uniform-ballot-op ::= ssa-id `=` `spv.GroupNonUniformBallot` scope ssa-use `:` `vector` `<` 4 `x` `integer-type` `>` diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVOps.td index a5bc740..91ea8d7 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVOps.td @@ -76,7 +76,7 @@ def SPV_AccessChainOp : SPV_Op<"AccessChain", [NoSideEffect]> { - 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 @@ -144,7 +144,7 @@ def SPV_ControlBarrierOp : SPV_Op<"ControlBarrier", []> { ### Custom assembly form - ``` {.ebnf} + ``` scope ::= `"CrossDevice"` | `"Device"` | `"Workgroup"` | ... memory-semantics ::= `"None"` | `"Acquire"` | "Release"` | ... @@ -190,7 +190,7 @@ def SPV_ExecutionModeOp : SPV_Op<"ExecutionMode", [InModuleScope]> { ### Custom assembly form - ``` {.ebnf} + ``` execution-mode ::= "Invocations" | "SpacingEqual" | @@ -243,7 +243,7 @@ def SPV_LoadOp : SPV_Op<"Load", []> { ### Custom assembly form - ``` {.ebnf} + ``` memory-access ::= `"None"` | `"Volatile"` | `"Aligned", ` integer-literal | `"NonTemporal"` @@ -298,7 +298,7 @@ def SPV_MemoryBarrierOp : SPV_Op<"MemoryBarrier", []> { ### Custom assembly form - ``` {.ebnf} + ``` scope ::= `"CrossDevice"` | `"Device"` | `"Workgroup"` | ... memory-semantics ::= `"None"` | `"Acquire"` | `"Release"` | ... @@ -343,7 +343,7 @@ def SPV_StoreOp : SPV_Op<"Store", []> { ### Custom assembly form - ``` {.ebnf} + ``` store-op ::= `spv.Store ` storage-class ssa-use `, ` ssa-use `, ` (`[` memory-access `]`)? `:` spirv-element-type ``` @@ -391,7 +391,7 @@ def SPV_UndefOp : SPV_Op<"undef", []> { ### Custom assembly form - ``` {.ebnf} + ``` undef-op ::= `spv.undef` `:` spirv-type ``` @@ -439,7 +439,7 @@ def SPV_VariableOp : SPV_Op<"Variable", []> { ### Custom assembly form - ``` {.ebnf} + ``` variable-op ::= ssa-id `=` `spv.Variable` (`init(` ssa-use `)`)? (`bind(` integer-literal, integer-literal `)`)? (`built_in(` string-literal `)`)? diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td index d57a7f4..a494e15 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td @@ -40,7 +40,7 @@ def SPV_AddressOfOp : SPV_Op<"_address_of", [InFunctionScope, NoSideEffect]> { ### Custom assembly form - ``` {.ebnf} + ``` spv-address-of-op ::= ssa-id `=` `spv._address_of` symbol-ref-id `:` spirv-pointer-type ``` @@ -89,7 +89,7 @@ def SPV_ConstantOp : SPV_Op<"constant", [NoSideEffect]> { ### Custom assembly form - ``` {.ebnf} + ``` spv-constant-op ::= ssa-id `=` `spv.constant` attribute-value (`:` spirv-type)? ``` @@ -160,7 +160,7 @@ def SPV_EntryPointOp : SPV_Op<"EntryPoint", [InModuleScope]> { ### Custom assembly form - ``` {.ebnf} + ``` execution-model ::= "Vertex" | "TesellationControl" | @@ -217,7 +217,7 @@ def SPV_GlobalVariableOp : SPV_Op<"globalVariable", [InModuleScope, Symbol]> { ### Custom assembly form - ``` {.ebnf} + ``` variable-op ::= `spv.globalVariable` spirv-type symbol-ref-id (`initializer(` symbol-ref-id `)`)? (`bind(` integer-literal, integer-literal `)`)? @@ -293,7 +293,7 @@ def SPV_ModuleOp : SPV_Op<"module", ### Custom assembly form - ``` {.ebnf} + ``` addressing-model ::= `"Logical"` | `"Physical32"` | `"Physical64"` memory-model ::= `"Simple"` | `"GLSL450"` | `"OpenCL"` | `"VulkanKHR"` spv-module-op ::= `spv.module` addressing-model memory-model @@ -393,7 +393,7 @@ def SPV_ReferenceOfOp : SPV_Op<"_reference_of", [NoSideEffect]> { ### Custom assembly form - ``` {.ebnf} + ``` spv-reference-of-op ::= ssa-id `=` `spv._reference_of` symbol-ref-id `:` spirv-scalar-type ``` @@ -434,7 +434,7 @@ def SPV_SpecConstantOp : SPV_Op<"specConstant", [InModuleScope, Symbol]> { ### Custom assembly form - ``` {.ebnf} + ``` spv-spec-constant-op ::= `spv.specConstant` symbol-ref-id `spec_id(` integer `)` `=` attribute-value (`:` spirv-type)? diff --git a/mlir/include/mlir/Dialect/VectorOps/VectorOps.td b/mlir/include/mlir/Dialect/VectorOps/VectorOps.td index 1e84010..f6e1ae5 100644 --- a/mlir/include/mlir/Dialect/VectorOps/VectorOps.td +++ b/mlir/include/mlir/Dialect/VectorOps/VectorOps.td @@ -484,7 +484,7 @@ def Vector_TransferReadOp : 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, vector<3x4x5xf32> @@ -544,7 +544,7 @@ def Vector_TransferReadOp : 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 ``` @@ -615,14 +615,14 @@ def Vector_TransferWriteOp : 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 { @@ -661,7 +661,7 @@ def Vector_TypeCastOp : Syntax: - ``` {.ebnf} + ``` operation ::= `vector.type_cast` ssa-use : memref-type to memref-type ``` diff --git a/mlir/lib/Conversion/VectorToLoops/ConvertVectorToLoops.cpp b/mlir/lib/Conversion/VectorToLoops/ConvertVectorToLoops.cpp index 43ad91c..681e6e9 100644 --- a/mlir/lib/Conversion/VectorToLoops/ConvertVectorToLoops.cpp +++ b/mlir/lib/Conversion/VectorToLoops/ConvertVectorToLoops.cpp @@ -51,7 +51,7 @@ using vector_type_cast = edsc::intrinsics::ValueBuilder; /// /// 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 @@ -68,7 +68,7 @@ using vector_type_cast = edsc::intrinsics::ValueBuilder; /// 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 diff --git a/mlir/utils/spirv/gen_spirv_dialect.py b/mlir/utils/spirv/gen_spirv_dialect.py index d1530f7..bf4886d 100755 --- a/mlir/utils/spirv/gen_spirv_dialect.py +++ b/mlir/utils/spirv/gen_spirv_dialect.py @@ -472,7 +472,7 @@ def get_op_definition(instruction, doc, existing_info): 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'\ -- 2.7.4