Drop Markdown style annotations
authorAlex Zinenko <zinenko@google.com>
Tue, 10 Dec 2019 11:00:29 +0000 (03:00 -0800)
committerA. Unique TensorFlower <gardener@tensorflow.org>
Tue, 10 Dec 2019 11:00:57 +0000 (03:00 -0800)
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

29 files changed:
mlir/g3doc/ConversionToLLVMDialect.md
mlir/g3doc/Diagnostics.md
mlir/g3doc/Dialects/Affine.md
mlir/g3doc/Dialects/GPU.md
mlir/g3doc/Dialects/LLVM.md
mlir/g3doc/Dialects/SPIR-V.md
mlir/g3doc/Dialects/Standard.md
mlir/g3doc/EDSC.md
mlir/g3doc/LangRef.md
mlir/g3doc/MLIRForGraphAlgorithms.md
mlir/g3doc/Rationale.md
mlir/g3doc/TestingGuide.md
mlir/g3doc/Tutorials/Toy/Ch-7.md
mlir/include/mlir/Dialect/GPU/GPUOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVArithmeticOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVAtomicOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVBitOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVCastOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVCompositeOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVControlFlowOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVGLSLOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVGroupOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVLogicalOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVNonUniformOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVOps.td
mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td
mlir/include/mlir/Dialect/VectorOps/VectorOps.td
mlir/lib/Conversion/VectorToLoops/ConvertVectorToLoops.cpp
mlir/utils/spirv/gen_spirv_dialect.py

index 595049ad440ab4267f33222800780ab1bd429416..782889fa3616b53dddab2bb10355b02b67076f97 100644 (file)
@@ -79,7 +79,7 @@ resulting in a struct containing two pointers + offset.
 
 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] }">
@@ -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<f32>
 
@@ -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}">
 
index 0c6ef7a24facf886577c059388c4112f1def0e3f..69a30942c0039b041480750c6edd8e14b8a5138d 100644 (file)
@@ -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}}
index 11cb93c10ea1c51d1b6b5a27deecff093aab4e17..c5dcf6a679027082da0618198b21734bbf8aeacf 100644 (file)
@@ -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<?x?xf32>, %B: memref<?x?xf32>) {
@@ -385,7 +385,7 @@ 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
 ```
@@ -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"() : () -> ()`
 ```
 
index bcb677d7660b2c5beb9434e3335472acf3b6ee5d..f29179ae98ff06606d86e7bd2308210f664d3560 100644 (file)
@@ -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.
index 9791352aa5615cf00acf7ba21ce94f77eaa03f4d..00d0fa02fece6c4de8fd4a17c38b366ae280e7d3 100644 (file)
@@ -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 ::= <canonical textual representation defined by LLVM>
 ```
@@ -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}">
 ```
index 58bd5ee828ffe3cd6b8f8612f05b45ad60158512..1413b181407a790be2252059e2afd2cbd67f42ba 100644 (file)
@@ -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` | <and other SPIR-V Dim specifiers...>
 
 depth-info ::= `NoDepth` | `IsDepth` | `DepthUnknown`
@@ -134,7 +134,7 @@ image-type ::= `!spv.image<` element-type `,` dim `,` depth-info `,`
 
 For example,
 
-``` {.mlir}
+```
 !spv.image<f32, 1D, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Unknown>
 !spv.image<f32, Cube, IsDepth, Arrayed, MultiSampled, NeedSampler, Rgba32f>
 ```
@@ -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<f32>
 !spv.struct<f32 [0]>
 !spv.struct<f32, !spv.image<f32, 1D, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Unknown>>
index 9d53eba328ea27e46b453a1a58619ed956aa01a6..fb0a4ed4fe1a3bbacc50aaeb593f54f5d42b9c9c 100644 (file)
@@ -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<?x?xi32>`
@@ -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<?x?xf32>
 
@@ -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<?x?xf32>
 %2 = tensor_cast %1 : tensor<*xf32> to tensor<?x?xf32>
@@ -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
 
index 264dcaf33e440f1ce5e0cc4b1aeb1079371e4ea3..afceac2dfc1644b7565ac83449c6d673bd398e4c 100644 (file)
@@ -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<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>
index a810330d37e6fb4bb5b6f0f0df1d25490ff9e6f2..cd6d3314c7c8a028840b95e9d130582988d8027d 100644 (file)
@@ -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 <name> `#` <opNo> 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<? x 128 x f32, #layout_map0>) -> 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<f32>
 complex<i32>
 ```
@@ -687,7 +687,7 @@ complex<i32>
 
 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<?x?xf32>`.
 
 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<i32, f32, tensor<i1>, 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<abcd>">
 
@@ -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<abcd>
 ```
@@ -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}
 
index 9130d031b4f1d291532b07c41b23bdbe447778b2..ac26e5beb9b93829945e8f25a1192a390650fea1 100644 (file)
@@ -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>
 ```
 
index efccf0752d779a608b5ef23e974b376a80ce953d..66cf800621d48b7d5dba21e8e37b8a95cf0dc0ce 100644 (file)
@@ -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<?x?xi32, %S: <?xi32>, %key : i32) {
   %ni = dim %A, 0 : memref<?x?xi32>
   // 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}
+```
 [<out-var-list> =]
 for %<index-variable-name> = <lower-bound> ... <upper-bound> step <step>
    [with <in-var-list>] { <loop-instruction-list> }
@@ -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<?xi32>, %N : i32) -> (i32) {
    %init = 0
@@ -1057,7 +1057,7 @@ func i32 @sum(%A : memref<?xi32>, %N : i32) -> (i32) {
 
 Syntax:
 
-``` {.ebnf}
+```
 <out-var-list> = affine.if (<cond-list>) {...} [else {...}]
 ```
 
@@ -1070,7 +1070,7 @@ this situation.
 
 Example:
 
-```mlir {.mlir}
+```mlir
 // Compute sum of half of the array
 func i32 @sum_half(%A : memref<?xi32>, %N : i32) -> (i32) {
    %s0 = 0
index 5a8c039fdd636a0be9038833f2b673d222400249..723b78bf0f58236c22e7548b5b8de66e2c2dbb47 100644 (file)
@@ -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.
index a55c5836a663e6812eb56dda1a22d06b5be06407..398983ac469cae9c201a24bebc09ea2a295c6bd2 100644 (file)
@@ -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)* `>`
 ```
 
index a1f4be60e218aa61581f3a83a40df2a195d1c783..dfb7ea91dc2d4159bfdc84f94aa9b1baa46028f0 100644 (file)
@@ -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<?xf32, 1> {
index 38b6391c5e8f71a15f7395b6e0741f43c016f5f4..f15d274922a7e762aa3fa6650aa9fb416468e138 100644 (file)
@@ -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
index 7042bf2cd3ee6d7a7e3e98e79ad54401041412ea..d8e62bb13b8bd77bed6a9277e3f2fbf34a1e522c 100644 (file)
@@ -35,7 +35,7 @@ def SPV_AtomicCompareExchangeWeakOp : SPV_Op<"AtomicCompareExchangeWeak", []> {
 
     ### Custom assembly form
 
-    ``` {.ebnf}
+    ```
     scope ::= `"CrossDevice"` | `"Device"` | `"Workgroup"` | ...
 
     memory-semantics ::= `"None"` | `"Acquire"` | "Release"` | ...
index e927f315a01804c4d4b7311f4241aa315ed121c0..d76a1e3854bd7328c49b9aaebcb1fb1b406c969c 100644 (file)
@@ -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
index 2faffb6280b28a5faea515d0db618d83edb07a98..e4fe526e420062e04466ae6bb843d6c1be3b2a18 100644 (file)
@@ -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
     ```
index 6392a1b52e82ebe1d32f1ba42c2ce7d8c4bdae85..3a21ddc871cf580802386b0e820b84d0a5a49906 100644 (file)
@@ -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
index ab9d51b05e7a34a55d571363620d1b050b655926..464b670dae9a48c511c1d0d0cfdeea55238d3d23 100644 (file)
@@ -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
     ```
 
index 2a1e8f328071a3668f13af7856d7c1c9fdb8f308..a031facdf5a59d36d405eabeadc90f8dafbbf59e 100644 (file)
@@ -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 |
index 5f60e6b013563a31de0d58208cbf5a0b47980d24..c0388fe4e23aad2d4647c0c2d74209e1fe8fb75c 100644 (file)
@@ -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` `>`
     ```
index be9825a89ab097ff96b46794383a3b7fb8f42511..0c4b2902a128a72e1d2d3e6ca097b75754e8bfa0 100644 (file)
@@ -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
index a37f5b576fd63d13867261d35586b03fdaa67fe8..1b3174c9e9fd3d88031c0c1e12d88373e22cd1d7 100644 (file)
@@ -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` `>`
index a5bc7402ece17ba7a24c337fcdc09c6a8d4916fd..91ea8d7d67613bdc0bac1a64a52cf44e59e7cbc1 100644 (file)
@@ -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" |
                        <and other SPIR-V execution modes...>
 
@@ -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 `)`)?
index d57a7f41d81554d2618b7b9b9f38e0aefee4aed6..a494e1517ba960fcd86ab881153f2e54a3b48585 100644 (file)
@@ -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" |
                         <and other SPIR-V execution models...>
 
@@ -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)?
index 1e8401005f41063b18c651edf67d595cbe5cf9c8..f6e1ae508e77505c02b1b4ca9a4b223dfd3ab8d0 100644 (file)
@@ -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<?x?x?x?xf32>, 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
     ```
 
index 43ad91ce87804214d586e68d0deee6cc545affea..681e6e964f25373ea95f244ff7ba264827816a99 100644 (file)
@@ -51,7 +51,7 @@ using vector_type_cast = edsc::intrinsics::ValueBuilder<vector::TypeCastOp>;
 ///
 /// 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<vector::TypeCastOp>;
 /// 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
index d1530f77d5ad18596e106664aa5a26f5e2ce40e4..bf4886dfd511246fa502c01e2afb1f8a8d032a1d 100755 (executable)
@@ -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'\