[mlir][NFC] Update dialect/op documentation to be consistent

[rriddle]

Summary

This revision performs a lot of different cleanups on operation documentation to ensure that they are consistent, e.g. using mlir code blocks, formatting, etc.

This revision also includes the auto-generated documentation into the hand-written documentation for the dialects that have a specific top-level dialect file. This updates the documentation for all dialects aside from SPIRV and STD. These dialects will be updated in a followup.

[rriddle]

Fix a few missed mlir code blocks.

[rriddle]

Rebase

[rriddle]

Rebase

[bondhugula]

Thanks very much for fixing all these!

mlir/docs/Dialects/Affine.md

-# Affine Dialect
+# `affine` Dialect
 
 This dialect provides a powerful abstraction for affine operations and analyses.
 
 [TOC]
 
 ## Polyhedral Structures
 
 MLIR uses techniques from polyhedral compilation to make dependence analysis and

 }
 ```
 
 `d0` and `d1` correspond to dimensional identifiers of the set, while `s0` and
 `s1` are symbol identifiers.
 
 ## Operations
 
-#### 'affine.apply' operation
+[include "Dialects/AffineOps.md"]
 
-Syntax:
-
-```
-operation ::= ssa-id `=` `affine.apply` affine-map-attribute dim-and-symbol-use-list
-```
-
-The `affine.apply` operation applies an
-[affine mapping](#affine-expressions) to a list of SSA values,
-yielding a single SSA value. The number of dimension and symbol arguments to
-affine.apply must be equal to the respective number of dimensional and symbolic
-inputs to the affine mapping; the `affine.apply` operation always returns one
-value. The input operands and result must all have 'index' type.
-
-Example:
-
-```mlir
-#map10 = affine_map<(d0, d1) -> (d0 floordiv 8 + d1 floordiv 128)>
-...
-%1 = affine.apply #map10 (%s, %t)
-
-// Inline example.
-%2 = affine.apply affine_map<(i)[s0] -> (i+s0)> (%42)[%n]
-```
-
-#### 'affine.for' operation
-
-Syntax:
-
-```
-operation   ::= `affine.for` ssa-id `=` lower-bound `to` upper-bound
-                      (`step` integer-literal)? `{` op* `}`
-
-lower-bound ::= `max`? affine-map-attribute dim-and-symbol-use-list | shorthand-bound
-upper-bound ::= `min`? affine-map-attribute dim-and-symbol-use-list | shorthand-bound
-shorthand-bound ::= ssa-id | `-`? integer-literal
-```
-
-The `affine.for` operation represents an affine loop nest. It has one region
-containing its body. This region must contain one block that terminates with
-[`affine.terminator`](#affineterminator-operation). *Note:* when `affine.for` is
-printed in custom format, the terminator is omitted. The block has one argument
-of [`index`](../LangRef.md#index-type) type that represents the induction
-variable of the loop.
-
-The `affine.for` operation executes its body a number of times iterating from a
-lower bound to an upper bound by a stride. The stride, represented by `step`, is
-a positive constant integer which defaults to "1" if not present. The lower and
-upper bounds specify a half-open range: the range includes the lower bound but
-does not include the!upper bound.
-
-The lower and upper bounds of a `affine.for` operation are represented as an
-application of an affine mapping to a list of SSA values passed to the map. The
-[same restrictions](#restrictions-on-dimensions-and-symbols) hold for these SSA
-values as for all bindings of SSA values to dimensions and symbols.
-
-The affine mappings for the bounds may return multiple results, in which case
-the `max`/`min` keywords are required (for the lower/upper bound respectively),
-and the bound is the maximum/minimum of the returned values. There is no
-semantic ambiguity, but MLIR syntax requires the use of these keywords to make
-things more obvious to human readers.
-
-Many upper and lower bounds are simple, so MLIR accepts two custom form
-syntaxes: the form that accepts a single 'ssa-id' (e.g. `%N`) is shorthand for
-applying that SSA value to a function that maps a single symbol to itself, e.g.,
-`()[s]->(s)()[%N]`. The integer literal form (e.g. `-42`) is shorthand for a
-nullary mapping function that returns the constant value (e.g. `()->(-42)()`).
-
-Example showing reverse iteration of the inner loop:
-
-```mlir
-#map57 = affine_map<(d0)[s0] -> (s0 - d0 - 1)>
-
-func @simple_example(%A: memref<?x?xf32>, %B: memref<?x?xf32>) {
-  %N = dim %A, 0 : memref<?x?xf32>
-  affine.for %i = 0 to %N step 1 {
-    affine.for %j = 0 to %N {   // implicitly steps by 1
-      %0 = affine.apply #map57(%j)[%N]
-      %tmp = call @F1(%A, %i, %0) : (memref<?x?xf32>, index, index)->(f32)
-      call @F2(%tmp, %B, %i, %0) : (f32, memref<?x?xf32>, index, index)->()
-    }
-  }
-  return
-}
-```
-
-#### 'affine.if' operation
-
-Syntax:
-
-```
-operation    ::= `affine.if` if-op-cond `{` op* `}` (`else` `{` op* `}`)?
-if-op-cond ::= integer-set-attr dim-and-symbol-use-list
-```
-
-The `affine.if` operation restricts execution to a subset of the loop iteration
-space defined by an integer set (a conjunction of affine constraints). A single
-`affine.if` may end with an optional `else` clause.
-
-The condition of the `affine.if` is represented by an
-[integer set](#integer-sets) (a conjunction of affine constraints),
-and the SSA values bound to the dimensions and symbols in the integer set. The
-[same restrictions](#restrictions-on-dimensions-and-symbols) hold for these SSA
-values as for all bindings of SSA values to dimensions and symbols.
-
-The `affine.if` operation contains two regions for the "then" and "else"
-clauses. The latter may be empty (i.e. contain no blocks), meaning the absence
-of the else clause. When non-empty, both regions must contain exactly one block
-terminating with [`affine.terminator`](#affineterminator-operation). *Note:*
-when `affine.if` is printed in custom format, the terminator is omitted. These
-blocks must not have any arguments.
-
-Example:
-
-```mlir
-#set = affine_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) {
-  affine.for %i = 0 to %N {
-     affine.for %j = 0 to %N {
-       %0 = affine.apply #map42(%j)
-       %tmp = call @S1(%X, %i, %0)
-       affine.if #set(%i, %j)[%N] {
-          %1 = affine.apply #map43(%i, %j)
-          call @S2(%tmp, %A, %i, %1)
-       }
-    }
-  }
-  return
-}
-```
-
-#### 'affine.load' operation
+### 'affine.load' operation
 
 Syntax:
 
 ```
 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
```mlir
 
   Example 2: Uses 'symbol' keyword for symbols '%n' and '%m'.
 
     %1 = affine.load %0[%i0 + symbol(%n), %i1 + symbol(%m)]
       : memref<100x100xf32>
 
 ```
 
-#### 'affine.store' operation
+### 'affine.store' operation
 
 Syntax:
 
 ```
 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
```mlir
 
     Example 2: Uses 'symbol' keyword for symbols '%n' and '%m'.
 
       affine.store %v0, %0[%i0 + symbol(%n), %i1 + symbol(%m)]
         : memref<100x100xf32>
 
 ```
 
-#### 'affine.dma_start' operation
+### 'affine.dma_start' operation
 
 Syntax:
 
 ```
 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
 ```
 
 The `affine.dma_start` op starts a non-blocking DMA operation that transfers

 %num_elements are transferred. Either both or no stride arguments should be
 specified. The value of 'num_elements' must be a multiple of
 'number_of_elements_per_stride'.
 
 
 Example:
 
 ```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
 space 1 at indices [%k + 7, %l], would be specified as follows:
 
   %num_elements = constant 256
   %idx = constant 0 : index
   %tag = alloc() : memref<1xi32, 4>
   affine.dma_start %src[%i + 3, %j], %dst[%k + 7, %l], %tag[%idx],
     %num_elements :
       memref<40x128xf32, 0>, memref<2x1024xf32, 1>, memref<1xi32, 2>
 
   If %stride and %num_elt_per_stride are specified, the DMA is expected to
   transfer %num_elt_per_stride elements every %stride elements apart from
   memory space 0 until %num_elements are transferred.
 
   affine.dma_start %src[%i, %j], %dst[%k, %l], %tag[%idx], %num_elements,
     %stride, %num_elt_per_stride : ...
-
 ```
 
-#### 'affine.dma_wait' operation
+### 'affine.dma_wait' operation
 
 Syntax:
 
 ```
 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
 ```
 
 The `affine.dma_start` op blocks until the completion of a DMA operation
 associated with the tag element '%tag[%index]'. %tag is a memref, and %index
 has to be an index with the same restrictions as any load/store index.
 In particular, index for each memref dimension must be an affine expression of
 loop induction variables and symbols. %num_elements is the number of elements
 associated with the DMA operation. For example:
 
 Example:
 
 ```mlir
-
-  affine.dma_start %src[%i, %j], %dst[%k, %l], %tag[%index], %num_elements :
-    memref<2048xf32, 0>, memref<256xf32, 1>, memref<1xi32, 2>
-  ...
-  ...
-  affine.dma_wait %tag[%index], %num_elements : memref<1xi32, 2>
-
-```
-
-#### 'affine.min' operation
-
-Syntax:
-
-```
-operation ::= ssa-id `=` `affine.min` affine-map-attribute dim-and-symbol-use-list
 ```
-
-The `affine.min` operation applies an
-[affine mapping](#affine-expressions) to a list of SSA values, and returns the
-minimum value of all result expressions. The number of dimension and symbol
-arguments to affine.min must be equal to the respective number of dimensional
-and symbolic inputs to the affine mapping; the `affine.min` operation always
-returns one value. The input operands and result must all have 'index' type.
-
-Example:
-
-```mlir
-
-%0 = affine.min affine_map<(d0)[s0] -> (1000, d0 + 512, s0)> (%arg0)[%arg1]
-
-```
-
-#### `affine.terminator` operation
-
-Syntax:
-
-```
-operation ::= `"affine.terminator"() : () -> ()`
-```
-
-Affine terminator is a special terminator operation for blocks inside affine
-loops ([`affine.for`](#affinefor-operation)) and branches
-([`affine.if`](#affineif-operation)). It unconditionally transmits the control
-flow to the successor of the operation enclosing the region.
-
-*Rationale*: bodies of affine operations are [blocks](../LangRef.md#blocks) that
-must have terminators. Loops and branches represent structured control flow and
-should not accept arbitrary branches as terminators.
-
-This operation does _not_ have a custom syntax. However, affine control
-operations omit the terminator in their custom syntax for brevity.

mlir/docs/Dialects/GPU.md

-# GPU Dialect
+# `gpu` Dialect
 
 Note: this dialect is more likely to change than others in the near future; use
 with caution.
 
 This dialect provides middle-level abstractions for launching GPU kernels
 following a programming model similar to that of CUDA or OpenCL. It provides
 abstractions for kernel invocations (and may eventually provide those for device
 management) that are not present at the lower level (e.g., as LLVM IR intrinsics

 incentive to use the buffers for communicating between different kernels or
 launches of the same kernel, which should be done through function arguments
 instead; we chose not to use `alloca`-style approach that would require more
 complex lifetime analysis following the principles of MLIR that promote
 structure and representing analysis results in the IR.
 
 ## Operations
 
-### `gpu.block_dim`
-
-Returns the number of threads in the thread block (aka the block size) along the
-x, y, or z `dimension`.
-
-Example:
-
-```mlir
-  %bDimX = "gpu.block_dim"() {dimension = "x"} : () -> (index)
-```
-
-### `gpu.block_id`
-
-Returns the block id, i.e. the index of the current block within the grid along
-the x, y, or z `dimension`.
-
-Example:
-
-```mlir
-  %bIdY = "gpu.block_id"() {dimension = "y"} : () -> (index)
-```
-
-### `gpu.grid_dim`
-
-Returns the number of thread blocks in the grid along the x, y, or z
-`dimension`.
-
-Example:
-
-```mlir
-  %gDimZ = "gpu.grid_dim"() {dimension = "z"} : () -> (index)
-```
-
-### `gpu.thread_id`
-
-Returns the thread id, i.e. the index of the current thread within the block
-along the x, y, or z `dimension`.
-
-Example:
-
-```mlir
-  %tIdX = "gpu.thread_id"() {dimension = "x"} : () -> (index)
-```
-
-### `gpu.yield`
-
-Is a special terminator operation for blocks inside regions in gpu ops. It
-returns values to the immediately enclosing gpu op.
-
-Example:
-
-```mlir
-gpu.yield %f0, %f1 : f32, f32
-```
-
-### `gpu.all_reduce`
-
-The "all_reduce" op reduces the value of every work item across a local
-workgroup. The result is equal for all work items of a workgroup.
-
-For example, both
-
-```mlir
-%1 = "gpu.all_reduce"(%0) ({}) { op = "add" } : (f32) -> (f32)
-%2 = "gpu.all_reduce"(%0) ({
-^bb(%lhs : f32, %rhs : f32):
-  %sum = addf %lhs, %rhs : f32
-  "gpu.yield"(%sum) : (f32) -> ()
-}) : (f32) -> (f32)
-```
-
-compute the sum of each work item's %0 value. The first version specifies the
-accumulation as operation, whereas the second version specifies the accumulation
-as code region. The accumulation operation must either be `add` or `mul`.
-
-Either none or all work items of a workgroup need to execute this op
-in convergence.
-
-### `gpu.barrier`
-
-The "barrier" op synchronizes all work items of a workgroup. It is used
-to coordinate communication between the work items of the workgroup.
-
-```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.
-
-Either none or all work items of a workgroup need to execute this op
-in convergence.
+[include "Dialects/GPUOps.md"]

mlir/docs/Dialects/LLVM.md

-# LLVM IR Dialect
+# `llvm` Dialect
 
 This dialect wraps the LLVM IR types and instructions into MLIR types and
 operations. It provides several additional operations that are necessary to
 cover for the differences in the IR structure (e.g., MLIR does not have `phi`
 operations and LLVM IR does not have a `constant` operation).
 
 In this document, we use "LLVM IR" to designate the
 [intermediate representation of LLVM](https://llvm.org/docs/LangRef.html) and

mlir/docs/Dialects/Linalg.md

-# Linalg Dialect
+# `linalg` Dialect
 
 [TOC]
 
 ## Rationale
 
 <img width="90" align="left" alt="MLIR Codegen Flow" src="https://user-images.githubusercontent.com/10148468/73613629-c5586580-45c5-11ea-94b7-074aeea94c7b.png">
 
 Linalg is designed to solve the High-level Hierarchical Optimization

 1. Is this general enough for the community's needs? If not how should this be
 extended, if at all?
 ...
 
 These key questions (and much more) should be really thought of in the general
 context of MLIR in which different levels of IR interoperate seamlessly. In
 practice, it is not necessary (or beneficial) to try and solve all problems in the
 same IR.
+
+## Operations
+
+[include "Dialects/LinalgOps.md"]

mlir/include/mlir/Dialect/Affine/IR/AffineOps.td

 
 // Require regions to have affine terminator.
 def ImplicitAffineTerminator
     : SingleBlockImplicitTerminator<"AffineTerminatorOp">;
 
 def AffineApplyOp : Affine_Op<"apply", [NoSideEffect]> {
   let summary = "affine apply operation";
   let description = [{
-    The affine.apply operation applies an affine mapping to a list of SSA
-    values, yielding a single SSA value. The number of dimension and symbol
-    arguments to affine.apply must be equal to the respective number of
-    dimensional and symbolic inputs to the affine mapping; the affine mapping
-    has to be one-dimensional, and so the affine.apply operation always returns
-    one value. The input operands and result must all have ‘index’ type.
+    The affine.apply operation applies an [affine mapping](#affine-expressions)
+    to a list of SSA values, yielding a single SSA value. The number of
+    dimension and symbol arguments to `affine.apply` must be equal to the
+    respective number of dimensional and symbolic inputs to the affine mapping;
+    the affine mapping has to be one-dimensional, and so the `affine.apply`
+    operation always returns one value. The input operands and result must all
+    have ‘index’ type.
 
     Example:
 
     ```mlir
-      #map10 = affine_map<(d0, d1) -> (d0 floordiv 8 + d1 floordiv 128)>
-      ...
-      %1 = affine.apply #map10 (%s, %t)
 
-      // Inline example.
-      %2 = affine.apply affine_map<(i)[s0] -> (i+s0)> (%42)[%n]
     ```
   }];
   let arguments = (ins AffineMapAttr:$map, Variadic<Index>:$mapOperands);
   let results = (outs Index);
 
   // TODO: The auto-generated builders should check to see if the return type
   // has a constant builder. That way we wouldn't need to explicitly specify the
   // result types here.
def AffineApplyOp : Affine_Op<"apply", [NoSideEffect]> {
   let hasFolder = 1;
 }
 
 def AffineForOp : Affine_Op<"for",
     [ImplicitAffineTerminator, RecursiveSideEffects,
      DeclareOpInterfaceMethods<LoopLikeOpInterface>]> {
   let summary = "for operation";
   let description = [{
-    The "affine.for" operation represents an affine loop nest, defining an SSA
-    value for its induction variable. It has one region capturing the loop body.
-    The induction variable is represented as a argument of this region. This SSA
-    value always has type index, which is the size of the machine word. The
-    stride, represented by step, is a positive constant integer which defaults
-    to "1" if not present. The lower and upper bounds specify a half-open range:
-    the range includes the lower bound but does not include the upper bound.
-
-    The body region must contain exactly one block that terminates with
-    "affine.terminator".  Calling AffineForOp::build will create such region
-    and insert the terminator, so will the parsing even in cases if it is absent
-    from the custom format.
+    Syntax:
 
-    The lower and upper bounds of a for operation are represented as an
+    ```
+    operation   ::= `affine.for` ssa-id `=` lower-bound `to` upper-bound
+                    (`step` integer-literal)? `{` op* `}`
+
+    lower-bound ::= `max`? affine-map-attribute dim-and-symbol-use-list | shorthand-bound
+    upper-bound ::= `min`? affine-map-attribute dim-and-symbol-use-list | shorthand-bound
+    shorthand-bound ::= ssa-id | `-`? integer-literal
+    ```
+
+    The `affine.for` operation represents an affine loop nest. It has one region
+    containing its body. This region must contain one block that terminates with
+    [`affine.terminator`](#affineterminator-operation). *Note:* when
+    `affine.for` is printed in custom format, the terminator is omitted. The
+    block has one argument of [`index`](../LangRef.md#index-type) type that
+    represents the induction variable of the loop.
+
+    The `affine.for` operation executes its body a number of times iterating
+    from a lower bound to an upper bound by a stride. The stride, represented by
+    `step`, is a positive constant integer which defaults to "1" if not present.
+    The lower and upper bounds specify a half-open range: the range includes the
+    lower bound but does not include the!upper bound.
+
+    The lower and upper bounds of a `affine.for` operation are represented as an
     application of an affine mapping to a list of SSA values passed to the map.
-    The same restrictions hold for these SSA values as for all bindings of SSA
-    values to dimensions and symbols. The affine mappings for the bounds may
-    return multiple results, in which case the max/min keywords are required
-    (for the lower/upper bound respectively), and the bound is the
-    maximum/minimum of the returned values.
+    The [same restrictions](#restrictions-on-dimensions-and-symbols) hold for
+    these SSA values as for all bindings of SSA values to dimensions and
+    symbols.
+
+    The affine mappings for the bounds may return multiple results, in which
+    case the `max`/`min` keywords are required (for the lower/upper bound
+    respectively), and the bound is the maximum/minimum of the returned values.
+    There is no semantic ambiguity, but MLIR syntax requires the use of these
+    keywords to make things more obvious to human readers.
+
+    Many upper and lower bounds are simple, so MLIR accepts two custom form
+    syntaxes: the form that accepts a single 'ssa-id' (e.g. `%N`) is shorthand
+    for applying that SSA value to a function that maps a single symbol to
+    itself, e.g., `()[s]->(s)()[%N]`. The integer literal form (e.g. `-42`) is
+    shorthand for a nullary mapping function that returns the constant value
+    (e.g. `()->(-42)()`).
 
-    Example:
+    Example showing reverse iteration of the inner loop:
 
-      affine.for %i = 1 to 10 {
-        ...
-      }
+    ```mlir
+    #map57 = affine_map<(d0)[s0] -> (s0 - d0 - 1)>
 
+    func @simple_example(%A: memref<?x?xf32>, %B: memref<?x?xf32>) {
+      %N = dim %A, 0 : memref<?x?xf32>
+      affine.for %i = 0 to %N step 1 {
+        affine.for %j = 0 to %N {   // implicitly steps by 1
+          %0 = affine.apply #map57(%j)[%N]
+          %tmp = call @F1(%A, %i, %0) : (memref<?x?xf32>, index, index)->(f32)
+          call @F2(%tmp, %B, %i, %0) : (f32, memref<?x?xf32>, index, index)->()
+        }
+      }
+      return
+    }
+    ```
   }];
   let arguments = (ins Variadic<AnyType>);
   let regions = (region SizedRegion<1>:$region);
 
   let skipDefaultBuilders = 1;
   let builders = [
     OpBuilder<"Builder *builder, OperationState &result, "
               "int64_t lowerBound, int64_t upperBound, int64_t step = 1">,
def AffineForOp : Affine_Op<"for",
   let hasCanonicalizer = 1;
   let hasFolder = 1;
 }
 
 def AffineIfOp : Affine_Op<"if",
                            [ImplicitAffineTerminator, RecursiveSideEffects]> {
   let summary = "if-then-else operation";
   let description = [{
-    The "if" operation represents an if-then-else construct for conditionally
-    executing two regions of code. The operands to an if operation are an
-    IntegerSet condition and a set of symbol/dimension operands to the
-    condition set. The operation produces no results. For example:
+    Syntax:
 
-       affine.if #set(%i)  {
-         ...
-       } else {
-         ...
-       }
+    ```
+    operation  ::= `affine.if` if-op-cond `{` op* `}` (`else` `{` op* `}`)?
+    if-op-cond ::= integer-set-attr dim-and-symbol-use-list
+    ```
 
-    The 'else' blocks to the if operation are optional, and may be omitted. For
-    example:
+    The `affine.if` operation restricts execution to a subset of the loop
+    iteration space defined by an integer set (a conjunction of affine
+    constraints). A single `affine.if` may end with an optional `else` clause.
+
+    The condition of the `affine.if` is represented by an
+    [integer set](#integer-sets) (a conjunction of affine constraints),
+    and the SSA values bound to the dimensions and symbols in the integer set.
+    The [same restrictions](#restrictions-on-dimensions-and-symbols) hold for
+    these SSA values as for all bindings of SSA values to dimensions and
+    symbols.
+
+    The `affine.if` operation contains two regions for the "then" and "else"
+    clauses. The latter may be empty (i.e. contain no blocks), meaning the
+    absence of the else clause. When non-empty, both regions must contain
+    exactly one block terminating with
+    [`affine.terminator`](#affineterminator-operation). *Note:* when `affine.if`
+    is printed in custom format, the terminator is omitted. These blocks must
+    not have any arguments.
 
-       affine.if #set(%i)  {
-         ...
+    Example:
+
+    ```mlir
+    #set = affine_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) {
+      affine.for %i = 0 to %N {
+         affine.for %j = 0 to %N {
+           %0 = affine.apply #map42(%j)
+           %tmp = call @S1(%X, %i, %0)
+           affine.if #set(%i, %j)[%N] {
+              %1 = affine.apply #map43(%i, %j)
+              call @S2(%tmp, %A, %i, %1)
+           }
         }
+      }
+      return
+    }
+    ```
   }];
   let arguments = (ins Variadic<AnyType>);
   let regions = (region SizedRegion<1>:$thenRegion, AnyRegion:$elseRegion);
 
   let skipDefaultBuilders = 1;
 
   let builders = [
     OpBuilder<"Builder *builder, OperationState &result, "
class AffineMinMaxOpBase<string mnemonic, list<OpTrait> traits = []> :
   let parser = [{ return ::parseAffineMinMaxOp<$cppClass>(parser, result); }];
   let hasFolder = 1;
   let hasCanonicalizer = 1;
 }
 
 def AffineMinOp : AffineMinMaxOpBase<"min", [NoSideEffect]> {
   let summary = "min operation";
   let description = [{
-    The "min" operation computes the minimum value result from a multi-result
-    affine map.
+    Syntax:
+
+    ```
+    operation ::= ssa-id `=` `affine.min` affine-map-attribute dim-and-symbol-use-list
+    ```
+
+    The `affine.min` operation applies an [affine mapping](#affine-expressions)
+    to a list of SSA values, and returns the minimum value of all result
+    expressions. The number of dimension and symbol arguments to `affine.min`
+    must be equal to the respective number of dimensional and symbolic inputs to
+    the affine mapping; the `affine.min` operation always returns one value. The
+    input operands and result must all have 'index' type.
 
     Example:
 
-       %0 = affine.min (d0) -> (1000, d0 + 512) (%i0) : index
+    ```mlir
+    %0 = affine.min affine_map<(d0)[s0] -> (1000, d0 + 512, s0)> (%arg0)[%arg1]
+    ```
   }];
 }
 
 def AffineMaxOp : AffineMinMaxOpBase<"max", [NoSideEffect]> {
   let summary = "max operation";
   let description = [{
     The "max" operation computes the maximum value result from a multi-result
     affine map.
 
     Example:
 
-       %0 = affine.max (d0) -> (1000, d0 + 512) (%i0) : index
+    ```mlir
+    ```
   }];
 }
 
 def AffineParallelOp : Affine_Op<"parallel", [ImplicitAffineTerminator]> {
   let summary = "multi-index parallel band operation";
   let description = [{
     The "affine.parallel" operation represents a hyper-rectangular affine
     parallel band, defining multiple SSA values for its induction variables. It
let description = [{
     Note: Calling AffineParallelOp::build will create the required region and
     block, and insert the required terminator. Parsing will also create the
     required region, block, and terminator, even when they are missing from the
     textual representation.
 
     Example:
 
     ```mlir
-      affine.parallel (%i, %j) = (0, 0) to (10, 10) step (1, 1) {
-        ...
-      }
     ```
   }];
 
   let arguments = (ins
      AffineMapAttr:$lowerBoundsMap,
      AffineMapAttr:$upperBoundsMap,
      I64ArrayAttr:$steps,
      Variadic<Index>:$mapOperands);

 def AffinePrefetchOp : Affine_Op<"prefetch"> {
   let summary = "affine prefetch operation";
   let description = [{
     The "affine.prefetch" op prefetches data from a memref location described
     with an affine subscript similar to affine.load, and has three attributes:
     a read/write specifier, a locality hint, and a cache type specifier as shown
     below:
 
-      affine.prefetch %0[%i, %j + 5], read, locality<3>, data
-                          : memref<400x400xi32>
+    ```mlir
+    affine.prefetch %0[%i, %j + 5], read, locality<3>, data : memref<400x400xi32>
+    ```
 
     The read/write specifier is either 'read' or 'write', the locality hint
     specifier ranges from locality<0> (no locality) to locality<3> (extremely
     local keep in cache). The cache type specifier is either 'data' or 'instr'
     and specifies whether the prefetch is performed on data cache or on
     instruction cache.
   }];
 
def AffinePrefetchOp : Affine_Op<"prefetch"> {
   let hasCanonicalizer = 1;
   let hasFolder = 1;
 }
 
 def AffineTerminatorOp :
     Affine_Op<"terminator", [NoSideEffect, Terminator]> {
   let summary = "affine terminator operation";
   let description = [{
+    Syntax:
+
+    ```
+    operation ::= `"affine.terminator"() : () -> ()`
+    ```
+
     Affine terminator is a special terminator operation for blocks inside affine
-    loops and branches. It unconditionally transmits the control flow to the
-    successor of the operation enclosing the region.
+    loops ([`affine.for`](#affinefor-operation)) and branches
+    ([`affine.if`](#affineif-operation)). It unconditionally transmits the
+    control flow to the successor of the operation enclosing the region.
+
+    *Rationale*: bodies of affine operations are [blocks](../LangRef.md#blocks)
+    that must have terminators. Loops and branches represent structured control
+    flow and should not accept arbitrary branches as terminators.
 
     This operation does _not_ have a custom syntax. However, affine control
     operations omit the terminator in their custom syntax for brevity.
   }];
 
   // No custom parsing/printing form.
   let parser = ?;
   let printer = ?;
 
   // Fully specified by traits.
   let verifier = ?;
 }
 
 #endif // AFFINE_OPS

mlir/include/mlir/Dialect/Affine/IR/CMakeLists.txt

 add_mlir_dialect(AffineOps affine)
-add_mlir_doc(AffineOps -gen-dialect-doc AffineDialect Dialects/)
+add_mlir_doc(AffineOps -gen-op-doc AffineOps Dialects/)

mlir/include/mlir/Dialect/GPU/CMakeLists.txt

-add_mlir_dialect(GPUOps gpu GPUOps)
-add_mlir_doc(GPUOps -gen-dialect-doc GPUDialect Dialects/)
+add_mlir_dialect(GPUOps gpu)
+add_mlir_doc(GPUOps -gen-op-doc GPUOps Dialects/)
 
 set(LLVM_TARGET_DEFINITIONS ParallelLoopMapperAttr.td)
 mlir_tablegen(ParallelLoopMapperAttr.h.inc -gen-struct-attr-decls)
 mlir_tablegen(ParallelLoopMapperAttr.cpp.inc -gen-struct-attr-defs)
 add_public_tablegen_target(MLIRParallelLoopMapperAttrGen)
 
 set(LLVM_TARGET_DEFINITIONS ParallelLoopMapperAttr.td)
 mlir_tablegen(ParallelLoopMapperEnums.h.inc -gen-enum-decls)
 mlir_tablegen(ParallelLoopMapperEnums.cpp.inc -gen-enum-defs)
 add_public_tablegen_target(MLIRParallelLoopMapperEnumsGen)

mlir/include/mlir/Dialect/GPU/GPUOps.td

class GPU_Op<string mnemonic, list<OpTrait> traits = []> :
     Op<GPU_Dialect, mnemonic, traits>;
 
 class GPU_IndexOp<string mnemonic, list<OpTrait> traits = []> :
     GPU_Op<mnemonic, !listconcat(traits, [NoSideEffect])>,
     Arguments<(ins StrAttr:$dimension)>, Results<(outs Index)> {
   let verifier = [{ return ::verifyIndexOp(*this); }];
 }
 
-def GPU_BlockDimOp : GPU_IndexOp<"block_dim">;
-def GPU_BlockIdOp : GPU_IndexOp<"block_id">;
-def GPU_GridDimOp : GPU_IndexOp<"grid_dim">;
-def GPU_ThreadIdOp : GPU_IndexOp<"thread_id">;
+def GPU_BlockDimOp : GPU_IndexOp<"block_dim"> {
+  let description = [{
+    Returns the number of threads in the thread block (aka the block size) along
+    the x, y, or z `dimension`.
+
+    Example:
+
+    ```mlir
+    %bDimX = "gpu.block_dim"() {dimension = "x"} : () -> (index)
+    ```
+  }];
+}
+def GPU_BlockIdOp : GPU_IndexOp<"block_id"> {
+  let description = [{
+    Returns the block id, i.e. the index of the current block within the grid
+    along the x, y, or z `dimension`.
+
+    Example:
+
+    ```mlir
+    %bIdY = "gpu.block_id"() {dimension = "y"} : () -> (index)
+    ```
+  }];
+}
+def GPU_GridDimOp : GPU_IndexOp<"grid_dim"> {
+  let description = [{
+    Returns the number of thread blocks in the grid along the x, y, or z
+    `dimension`.
+
+    Example:
+
+    ```mlir
+    %gDimZ = "gpu.grid_dim"() {dimension = "z"} : () -> (index)
+    ```
+  }];
+}
+def GPU_ThreadIdOp : GPU_IndexOp<"thread_id"> {
+  let description = [{
+    Returns the thread id, i.e. the index of the current thread within the block
+    along the x, y, or z `dimension`.
+
+    Example:
+
+    ```mlir
+    %tIdX = "gpu.thread_id"() {dimension = "x"} : () -> (index)
+    ```
+  }];
+}
 
 def GPU_GPUFuncOp : GPU_Op<"func", [FunctionLike, IsolatedFromAbove, Symbol]> {
   let summary = "Function executable on a GPU";
 
   let description = [{
     Defines a function that can be executed on a GPU. This supports memory
     attribution and its body has a particular execution model.
 
def GPU_TerminatorOp : GPU_Op<"terminator", [HasParent<"LaunchOp">,
   let parser = [{ return success(); }];
   let printer = [{ p << getOperationName(); }];
 }
 
 def GPU_YieldOp : GPU_Op<"yield", [NoSideEffect, Terminator]>,
     Arguments<(ins Variadic<AnyType>:$values)> {
   let summary = "GPU yield operation";
   let description = [{
-    "gpu.yield" is a special terminator operation for blocks inside regions
+    gpu.yield` is a special terminator operation for blocks inside regions
     in gpu ops. It returns values to the immediately enclosing gpu op.
 
     Example:
 
-       ```gpu.yield %f0, %f1 : f32, f32
-       ```
+    ```mlir
+    gpu.yield %f0, %f1 : f32, f32
   }];
 }
 
 // add, mul mirror the XLA ComparisonDirection enum.
 def GPU_AllReduceOpAdd : StrEnumAttrCase<"add">;
 def GPU_AllReduceOpAnd : StrEnumAttrCase<"and">;
 def GPU_AllReduceOpMax : StrEnumAttrCase<"max">;
 def GPU_AllReduceOpMin : StrEnumAttrCase<"min">;

 
 def GPU_AllReduceOp : GPU_Op<"all_reduce",
     [SameOperandsAndResultType, IsolatedFromAbove]>,
     Arguments<(ins AnyType:$value,
                OptionalAttr<GPU_AllReduceOperationAttr>:$op)>,
     Results<(outs AnyType)> {
   let summary = "Reduce values among workgroup.";
   let description = [{
-    The "all_reduce" op reduces the value of every work item across a local
+    The `all_reduce` op reduces the value of every work item across a local
     workgroup. The result is equal for all work items of a workgroup.
 
     For example, both
-    ```
-      %1 = "gpu.all_reduce"(%0) ({}) { op = "add" } : (f32) -> (f32)
-      %2 = "gpu.all_reduce"(%0) ({
-      ^bb(%lhs : f32, %rhs : f32):
-        %sum = addf %lhs, %rhs : f32
-        "gpu.yield"(%sum) : (f32) -> ()
-      }) : (f32) -> (f32)
+
+    ```mlir
     ```
+
     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 be one of:
     `add`, `and`, `max`, `min`, `mul`, `or`, `xor`.
 
     Either none or all work items of a workgroup need to execute this op
     in convergence.
   }];
def GPU_ShuffleOp : GPU_Op<"shuffle", [NoSideEffect]>,
     Arguments<(ins AnyType:$value, I32:$offset, I32:$width,
                GPU_ShuffleModeAttr:$mode)>,
     Results<(outs AnyType:$result, I1:$valid)> {
   let summary = "Shuffles values within a subgroup.";
   let description = [{
     The "shuffle" op moves values to a different invocation within the same
     subgroup.
 
-    For example
-    ```
-      %1, %2 = gpu.shuffle %0, %offset, %width xor : f32
+    Example:
+
+    ```mlir
     ```
-    for lane k returns the value from lane `k ^ offset` and `true` if that lane
+
+    For lane k returns the value from lane `k ^ offset` and `true` if that lane
     is smaller than %width. Otherwise it returns an unspecified value and
     `false`. A lane is the index of an invocation relative to its subgroup.
 
     The width specifies the number of invocations that participate in the
     shuffle. The width needs to be the same for all invocations that participate
     in the shuffle. Exactly the first `width` invocations of a subgroup need to
     execute this op in convergence.
   }];
   let verifier = [{ return ::verifyShuffleOp(*this); }];
   let printer = [{ printShuffleOp(p, *this); }];
   let parser = [{ return parseShuffleOp(parser, result); }];
 }
 
 def GPU_BarrierOp : GPU_Op<"barrier"> {
   let summary = "Synchronizes all work items of a workgroup.";
   let description = [{
     The "barrier" op synchronizes all work items of a workgroup. It is used
     to coordinate communication between the work items of the workgroup.
 
-    ```
-      gpu.barrier
+    ```mlir
     ```
+
     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.

mlir/include/mlir/Dialect/Linalg/IR/CMakeLists.txt

 add_mlir_dialect(LinalgOps linalg)
-add_mlir_doc(LinalgDoc -gen-dialect-doc LinalgDialect Dialects/)
+add_mlir_doc(LinalgDoc -gen-op-doc LinalgOps Dialects/)
+
 set(LLVM_TARGET_DEFINITIONS LinalgStructuredOps.td)
 mlir_tablegen(LinalgStructuredOps.h.inc -gen-op-decls)
 mlir_tablegen(LinalgStructuredOps.cpp.inc -gen-op-defs)
 add_public_tablegen_target(MLIRLinalgStructuredOpsIncGen)
 
 set(LLVM_TARGET_DEFINITIONS LinalgStructuredOpsInterface.td)
 mlir_tablegen(LinalgStructuredOpsInterfaces.h.inc -gen-op-interface-decls)
 mlir_tablegen(LinalgStructuredOpsInterfaces.cpp.inc -gen-op-interface-defs)
 add_public_tablegen_target(MLIRLinalgStructuredOpsInterfaceIncGen)

mlir/include/mlir/Dialect/Linalg/IR/LinalgOps.td

def Linalg_RangeOp :
   let description = [{
     The `linalg.range` op creates a `!linalg.range` from 3 values of type
     `index` that represent the min, max and step values of the `range`. This
     type does not pass function boundaries at the moment.
 
     Example:
 
     ```mlir
-      %3 = linalg.range %0:%1:%2 : !linalg.range
     ````
   }];
   let builders = [OpBuilder<
     "Builder *builder, OperationState &result, Value min, Value max, "
     "Value step",
     [{
       auto rangeType = RangeType::get(builder->getContext());
       build(builder, result, rangeType, min, max, step);
let description = [{
     relationship between source and target memref ranks. The verification rule
     is that the reassociation maps are applied to the memref with the larger
     rank to obtain the memref with the smaller rank. In the case of a dimension
     expansion, the reassociation maps can be interpreted as inverse maps.
 
     Examples:
 
     ```mlir
-       // Dimension collapse (i, j) -> i' and k -> k'
-       %1 = linalg.reshape %0 [(i, j, k) -> (i, j), (i, j, k) -> (k)] :
-         memref<?x?x?xf32, stride_spec> into memref<?x?xf32, stride_spec_2>
     ```
 
     ```mlir
-       // Dimension expansion i -> (i', j') and (k) -> (k')
-       %1 = linalg.reshape %0 [(i, j, k) -> (i, j), (i, j, k) -> (k)] :
-         memref<?x?xf32, stride_spec> into memref<?x?x?xf32, stride_spec_2>
     ```
   }];
 
   let builders = [
     // Builder for a contracting reshape whose result type is computed from
     // `view` and `reassociation`.
     OpBuilder<"Builder *b, OperationState &result, Value view, "
     "ArrayRef<ArrayRef<AffineExpr>> reassociation, "
let description = [{
     it to be within bounds but there are tradeoffs involved now that `std.view`
     is a standard op.
 
     Examples:
 
       1. rank-preserving `slice`:
 
       ```mlir
-        %4 = linalg.slice %0[%1, %2] : memref<?x?xf32, stride_spec>,
-          !linalg.range, !linalg.range, memref<?x?xf32, stride_spec>
-       ```
 
       2. rank-reducing `slice` (from 2-D to 1-D):
 
       ```mlir
-        %4 = linalg.slice %0[%1, %2] : memref<?x?xf32, stride_spec>,
-          index, !linalg.range, memref<?x?xf32, stride_spec>
       ```
 
       3. rank-reducing `slice` (from 2-D to 0-D):
 
       ```mlir
-        %4 = linalg.slice %0[%1, %2] : memref<?x?xf32, stride_spec>,
-          index, index, memref<?x?xf32, stride_spec>
       ```
   }];
 
   let builders = [OpBuilder<
     "Builder *b, OperationState &result, Value base, "
     "ValueRange indexings">];
 
   let extraClassDeclaration = [{
def Linalg_TransposeOp : Linalg_Op<"transpose", [NoSideEffect]>,
   let description = [{
     The `linalg.transpose` op produces a strided memref whose sizes and strides
     are a permutation of the original `view`. This is a pure metadata
     transformation.
 
     Example:
 
     ```mlir
-       %1 = linalg.transpose %0 (i, j) -> (j, i) : memref<?x?xf32, stride_spec>
     ```
   }];
 
   let builders = [OpBuilder<
     "Builder *b, OperationState &result, Value view, "
     "AffineMapAttr permutation, ArrayRef<NamedAttribute> attrs = {}">];
 
   let verifier = [{
def Linalg_YieldOp : Linalg_Op<"yield", [NoSideEffect, Terminator]>,
   let description = [{
     `linalg.yield` is a special terminator operation for blocks inside regions
     in `linalg` generic ops. It returns values to the immediately enclosing
     `linalg` generic op.
 
     Example:
 
     ```mlir
-       linalg.yield %f0, %f1 : f32, f32
     ```
   }];
 }
 
 #endif // LINALG_OPS

mlir/include/mlir/Dialect/Linalg/IR/LinalgStructuredOps.td

 //===----------------------------------------------------------------------===//
 // At the moment these are not declarative and require a bunch of C++ code.
 // In the future, these should be migrated to a declarative specification.
 def CopyOp : LinalgStructured_Op<"copy", [NInputs<1>, NOutputs<1>]> {
   let description = [{
     Copies the data in the input view into the output view.
 
     Usage:
-      ```mlir
-      linalg.copy(%arg0, %arg1) : memref<?xf32, stride_specification>,
-                                  memref<?xf32, stride_specification>
-      ```
+
 
     One possible lowering to loop form is:
-      ```mlir
-      %0 = linalg.dim %arg0, 0 : index
-      loop.for %i0 = %c0 to %0 step %c1 {
-        %1 = load %arg0[%i0] : memref<?xf32, stride_specification>
-        store %1, %arg1[%i0] : memref<?xf32, stride_specification>
-      }
-      ```
+
 
     Optionally, can take `input_permutation` and `output_permutation` attributes
     to reorder the dimensions of the input and output views.
 
     Usage:
-      ```mlir
-      linalg.copy(%arg0, %arg1) {inputPermutation : (i, j, k) -> (i, k, j),
-                                 outputPermutation : (i, j, k) -> (k, j, i)} :
-        memref<?x?x?xf32, stride_specification>,
-        memref<?x?x?xf32, stride_specification>
-     ```
+
 
     One possible lowering to loop form is:
-      ```mlir
-      %0 = linalg.dim %arg0, 0
-      %1 = linalg.dim %arg0, 1
-      %2 = linalg.dim %arg0, 2
-      loop.for %i0 = %c0 to %{{.*}} step %c1 {
-        loop.for %i1 = %c0 to %{{.*}} step %c1 {
-          loop.for %i2 = %c0 to %{{.*}} step %c1 {
-            %3 = load %arg0[%i0, %i2, %i1] :
-                    memref<?x?x?xf32, stride_specification>
-            store %3, %arg1[%i2, %i1, %i0] :
-                    memref<?x?x?xf32, stride_specification>
-      ```
+
 
     The views are expected to be compatible for correctness but this is not
     enforced at the moment.
   }];
 
   let arguments = (ins
     AnyStridedMemRef:$input,
     AnyStridedMemRef:$output,

 
 /// Index-free GenericOp.
 def GenericOp : GenericOpBase<"generic"> {
   let description = [{
     Generic Linalg op form where the key properties of the computation are
     specified as attributes. In pretty form, a linalg.generic op is written as:
 
       ```mlir
-        linalg.generic #trait_attribute %A, %B, %C {other-attributes} :
-          memref<?x?xf32, stride_specification>,
-          memref<?x?xf32, stride_specification>,
-          memref<?x?xf32, stride_specification>
       ```
 
     Where #trait_attributes is an alias of a dictionary attribute containing:
       - args_in: an I64Attr representing the number of input (readonly) views
       - args_out: an I64Attr representing the number of output (readwrite) views
       - doc [optional]: a documentation string
       - fun: a FlatSymbolRefAttr that must resolve to an existing function
         symbol. To support inplace updates in a generic fashion, the signature
Where #trait_attributes is an alias of a dictionary attribute containing:
       - iterator_types: an ArrayAttr specifying the type of the enclosing loops.
         Each element of the list represents and iterator of one of the following
         types:
           parallel, reduction, window
 
     Example:
     Defining a #matmul_trait attribute in MLIR can be done as follows:
       ```mlir
-        func @fma(%a: f32, %b: f32, %c: f32) -> f32 {
-          %d = mulf %a, %b: f32
-          %e = addf %c, %d: f32
-          return %e: f32
-        }
-        #matmul_accesses = [
-          (m, n, k) -> (m, k),
-          (m, n, k) -> (k, n),
-          (m, n, k) -> (m, n)
-        ]
-        #matmul_trait = {
-          doc = "C(m, n) += A(m, k) * B(k, n)",
-          fun = @fma,
-          indexing_maps = #matmul_accesses,
-          library_call = "linalg_matmul",
-          n_views = [2, 1],
-          iterator_types = ["parallel", "parallel", "reduction"]
-        }
       ```
 
     And can be reused in multiple places as:
       ```mlir
-        linalg.generic #matmul_trait %A, %B, %C [other-attributes] :
-          memref<?x?xf32, stride_specification>,
-          memref<?x?xf32, stride_specification>,
-          memref<?x?xf32, stride_specification>
       ```
 
     This may lower to either:
       ```mlir
-        call @linalg_matmul(%A, %B, %C) :
-          (memref<?x?xf32, stride_specification>,
-           memref<?x?xf32, stride_specification>,
-           memref<?x?xf32, stride_specification>)
-          -> ()
       ```
 
     or IR resembling:
     ```mlir
     loop.for %m = %c0 to %M step %c1 {
       loop.for %n = %c0 to %N step %c1 {
         loop.for %k = %c0 to %K step %c1 {
           %a = load %A[%m, %k] : memref<?x?xf32, stride_specification>
           %b = load %B[%k, %n] : memref<?x?xf32, stride_specification>
           %c = load %C[%m, %n] : memref<?x?xf32, stride_specification>
           %d = call @func_of_elements(%a, %b, %c)
                  : (f32, f32, f32) -> (f32)
           store %d, %C[%m, %n] : memref<?x?x?xf32, stride_specification>
         }
       }
     }
     ```
 
     To allow progressive lowering from the value world (a.k.a tensor values) to
     the buffer world (a.k.a memref values), a `linalg.generic` op accepts
     mixing input and output ranked tensor values with input and output memrefs.
 
     ```mlir
-      %C = linalg.generic #trait_attribute %A, %B {other-attributes} :
-        tensor<?x?xf32>,
-        memref<?x?xf32, stride_specification>
-        -> (tensor<?x?xf32>)
     ```
 
     In this case, the number of outputs (args_out) must match the sum of (1) the
     number of output buffer operands and (2) the number of tensor return values.
     The semantics is that the `linalg.indexed_generic` op produces (i.e.
     allocates and fills) its tensor return values.
 
     Tensor values must be legalized by a buffer allocation pass before most

 /// the enclosing loop induction variables)
 def IndexedGenericOp : GenericOpBase<"indexed_generic"> {
   let description = [{
     Indexed Generic Linalg op form where the key properties of the computation
     are specified as attributes. In pretty form, a linalg.indexed_generic op is
     written as:
 
       ```mlir
-        linalg.indexed_generic #trait_attribute %A, %B, %C {other-attributes} :
-          memref<?x?xf32, stride_specification>,
-          memref<?x?xf32, stride_specification>,
-          memref<?x?xf32, stride_specification>
       ```
 
     Where #trait_attributes is an alias of a dictionary attribute containing:
       - args_in: an I64Attr representing the number of input (readonly) views
       - args_out: an I64Attr representing the number of output (readwrite) views
       - doc [optional]: a documentation string
       - fun: a FlatSymbolRefAttr that must resolve to an existing function
         symbol. To support inplace updates in a generic fashion, the signature
Where #trait_attributes is an alias of a dictionary attribute containing:
         a library call, linalg.indexed_generic will always lower to loops.
       - iterator_types: an ArrayAttr they type of the enclosing loops; Each
         element of the list represents and iterator of one of the following
         types:
           parallel, reduction, window
 
     Example:
     Defining a #matmul_trait attribute in MLIR can be done as follows:
-      ```mlir
-        func @fma(%offset_m: index, %offset_n: index, %offset_k: index,
-                  %a: f32, %b: f32, %c: f32)
-          -> f32
-        {
-          "some_optional_condition"(%offset_m, %offset_n, %offset_k)
-          %d = mulf %a, %b: f32
-          %e = addf %c, %d: f32
-          return %e: f32
-        }
-        #matmul_accesses = [
-          (m, n, k) -> (m, k),
-          (m, n, k) -> (k, n),
-          (m, n, k) -> (m, n)
-        ]
-        #matmul_trait = {
-          doc = "C(m, n) += A(m, k) * B(k, n)",
-          fun = @fma,
-          indexing_maps = #matmul_accesses,
-          library_call = "linalg_matmul",
-          n_views = [2, 1],
-          iterator_types = ["parallel", "parallel", "reduction"]
-        }
-      ```
+
 
     And can be reused in multiple places as:
-      ```mlir
-        linalg.indexed_generic #matmul_trait %A, %B, %C [other-attributes] :
-          memref<?x?xf32, stride_specification>,
-          memref<?x?xf32, stride_specification>,
-          memref<?x?xf32, stride_specification>
-      ```
+
 
     This may lower to either:
-      ```mlir
-        call @linalg_matmul(%offset_m, %offset_n, %offset_k, %A, %B, %C) :
-          (memref<?x?xf32, stride_specification>,
-           memref<?x?xf32, stride_specification>,
-           memref<?x?xf32, stride_specification>)
-          -> ()
-      ```
+
 
     or IR resembling:
+
     ```mlir
     loop.for %m = %c0 to %M step %c1 {
       loop.for %n = %c0 to %N step %c1 {
         loop.for %k = %c0 to %K step %c1 {
           %a = load %A[%m, %k] : memref<?x?xf32, stride_specification>
           %b = load %B[%k, %n] : memref<?x?xf32, stride_specification>
           %c = load %C[%m, %n] : memref<?x?xf32, stride_specification>
           %d = call @func_of_elements_and_indices(%m, %n, %k, %a, %b, %c)
                  : (index, index, index, f32, f32, f32) -> (f32)
           store %d, %C[%m, %n] : memref<?x?x?xf32, stride_specification>
         }
       }
     }
     ```
 
     To allow progressive lowering from the value world (a.k.a tensor values) to
     the buffer world (a.k.a memref values), a `linalg.indexed_generic` op
     accepts mixing input and output ranked tensor values with input and output
     memrefs.
 
     ```mlir
-      %C = linalg.indexed_generic #trait_attribute %A, %B {other-attributes}
-      : tensor<?x?xf32>,
-        memref<?x?xf32, stride_specification>
-        -> (tensor<?x?xf32>)
     ```
 
     In this case, the number of outputs (args_out) must match the sum of (1) the
     number of output buffer operands and (2) the number of tensor return values.
     The semantics is that the `linalg.indexed_generic` op produces (i.e.
     allocates and fills) its return values.
 
     Tensor values must be legalized by a buffer allocation pass before most

mlir/include/mlir/Dialect/LoopOps/LoopOps.td

let description = [{
     the lower bound but does not include the upper bound.
 
     The body region must contain exactly one block that terminates with
     "loop.yield".  Calling ForOp::build will create such a region and insert
     the terminator implicitly if none is defined, so will the parsing even in
     cases when it is absent from the custom format. For example:
 
     ```mlir
-       loop.for %iv = %lb to %ub step %step {
-         ... // body
-       }
     ```
 
-    "loop.for" can also operate on loop-carried variables and returns the final
+    `loop.for` can also operate on loop-carried variables and returns the final
     values after loop termination. The initial values of the variables are
     passed as additional SSA operands to the "loop.for" following the 3 loop
     control SSA values mentioned above (lower bound, upper bound and step). The
     operation region has equivalent arguments for each variable representing
     the value of the variable at the current iteration.
 
     The region must terminate with a "loop.yield" that passes all the current
     iteration variables to the next iteration, or to the "loop.for" result, if
func @conditional_reduce(%buffer: memref<1024xf32>, %lb: index,
           loop.yield %new_sum : f32
         } else {
           loop.yield %sum_iter : f32
         }
         loop.yield %sum_next : f32
       }
       return %sum : f32
     }
     ```
   }];
   let arguments = (ins Index:$lowerBound,
                        Index:$upperBound,
                        Index:$step,
                        Variadic<AnyType>:$initArgs);
   let results = (outs Variadic<AnyType>:$results);
   let regions = (region SizedRegion<1>:$region);
 
let extraClassDeclaration = [{
     }
   }];
 }
 
 def IfOp : Loop_Op<"if",
       [SingleBlockImplicitTerminator<"YieldOp">, RecursiveSideEffects]> {
   let summary = "if-then-else operation";
   let description = [{
-    The "loop.if" operation represents an if-then-else construct for
+    The `loop.if` operation represents an if-then-else construct for
     conditionally executing two regions of code. The operand to an if operation
     is a boolean value. For example:
 
     ```mlir
-       loop.if %b  {
-         ...
-       } else {
-         ...
-       }
     ```
 
-    "loop.if" may also return results that are defined in its regions. The
+    `loop.if` may also return results that are defined in its regions. The
     values defined are determined by which execution path is taken.
-    For example:
+
+    Example:
+
     ```mlir
-       %x, %y = loop.if %b -> (f32, f32) {
-         %x_true = ...
-         %y_true = ...
-         loop.yield %x_true, %y_true : f32, f32
-       } else {
-         %x_false = ...
-         %y_false = ...
-         loop.yield %x_false, %y_false : f32, f32
-       }
     ```
 
-    "loop.if" regions are always terminated with "loop.yield". If "loop.if"
+    `loop.if` regions are always terminated with "loop.yield". If "loop.if"
     defines no values, the "loop.yield" can be left out, and will be inserted
     implicitly. Otherwise, it must be explicit.
     Also, if "loop.if" defines one or more values, the 'else' block cannot be
     omitted.
 
-    For example:
+    Example:
+
     ```mlir
-       loop.if %b  {
-         ...
-       }
     ```
   }];
   let arguments = (ins I1:$condition);
   let results = (outs Variadic<AnyType>:$results);
   let regions = (region SizedRegion<1>:$thenRegion, AnyRegion:$elseRegion);
 
   let skipDefaultBuilders = 1;
   let builders = [
let description = [{
     initial values, respectively. The operation defines a variadic number of
     SSA values for its induction variables. It has one region capturing the
     loop body. The induction variables are represented as an argument of this
     region. These SSA values always have type index, which is the size of the
     machine word. The steps are values of type index, required to be positive.
     The lower and upper bounds specify a half-open range: the range includes
     the lower bound but does not include the upper bound. The initial values
     have the same types as results of "loop.parallel". If there are no results,
-    the keyword `init` can be omitted. 
+    the keyword `init` can be omitted.
 
     Semantically we require that the iteration space can be iterated in any
     order, and the loop body can be executed in parallel. If there are data
     races, the behavior is undefined.
 
     The parallel loop operation supports reduction of values produced by
     individual iterations into a single result. This is modeled using the
     loop.reduce operation (see loop.reduce for details). Each result of a
     loop.parallel operation is associated with an initial value operand and
     reduce operation that is an immediate child. Reductions are matched to
     result and initial values in order of their appearance in the body.
     Consequently, we require that the body region has the same number of
     results and initial values as it has reduce operations.
 
     The body region must contain exactly one block that terminates with
     "loop.yield" without operands. Parsing ParallelOp will create such a region
     and insert the terminator when it is absent from the custom format.
-    For example:
+
+    Example:
 
     ```mlir
-       loop.parallel (%iv) = (%lb) to (%ub) step (%step) -> f32 {
-         %zero = constant 0.0 : f32
-         loop.reduce(%zero) : f32 {
-           ^bb0(%lhs : f32, %rhs: f32):
-             %res = addf %lhs, %rhs : f32
-             loop.reduce.return %res : f32
-         }
-       }
     ```
-   }];
 
   let arguments = (ins Variadic<Index>:$lowerBound,
                        Variadic<Index>:$upperBound,
                        Variadic<Index>:$step,
                        Variadic<AnyType>:$initVals);
   let results = (outs Variadic<AnyType>:$results);
   let regions = (region SizedRegion<1>:$region);
 
let description = [{
     reduction may be non-deterministic unless the operation is associative and
     commutative.
 
     The result of the reduce operation's body must have the same type as the
     operands and associated result value of the parallel loop operation.
     Example:
 
     ```mlir
-      %operand = constant 1.0 : f32
-      loop.reduce(%operand) : f32 {
-        ^bb0(%lhs : f32, %rhs: f32):
-          %res = addf %lhs, %rhs : f32
-          loop.reduce.return %res : f32
-      }
     ```
-
   }];
 
   let skipDefaultBuilders = 1;
   let builders = [
     OpBuilder<"Builder *builder, OperationState &result, "
               "Value operand">
   ];
 
   let arguments = (ins AnyType:$operand);
   let regions = (region SizedRegion<1>:$reductionOperator);
 }
 
 def ReduceReturnOp :
     Loop_Op<"reduce.return", [HasParent<"ReduceOp">, NoSideEffect,
                               Terminator]> {
   let summary = "terminator for reduce operation";
   let description = [{
     "loop.reduce.return" is a special terminator operation for the block inside
     "loop.reduce". It terminates the region. It should have the same type as
     the operand of "loop.reduce". Example for the custom format:
 
     ```mlir
-      loop.reduce.return %res : f32
     ```
   }];
 
   let arguments = (ins AnyType:$result);
   let assemblyFormat = "$result attr-dict `:` type($result)";
 }
 
 def YieldOp : Loop_Op<"yield", [NoSideEffect, Terminator]> {

mlir/include/mlir/Dialect/Quant/QuantOps.td

 // high-precision ops with quantization specifications for all the inputs
 // and outputs. Some quantization specifications can be undetermined and
 // derived from other ports by the target specification of the kernel.
 def quant_QuantizeRegionOp : quant_Op<"region", [
     NoSideEffect,
     IsolatedFromAbove,
     SingleBlockImplicitTerminator<"ReturnOp">]> {
   let summary = [{
-    The `region operation wraps high-precision ops as a logical low-precision
+    The `region` operation wraps high-precision ops as a logical low-precision
     quantized kernel.
   }];
 
   let arguments = (ins Variadic<AnyType>:$inputs,
                     TypeArrayAttr:$input_specs,
                     TypeArrayAttr:$output_specs,
                     StrAttr:$logical_kernel);
   let results = (outs Variadic<AnyType>:$outputs);

 }
 
 //===----------------------------------------------------------------------===//
 // Training integration and instrumentation ops
 //===----------------------------------------------------------------------===//
 
 def quant_ConstFakeQuant : quant_Op<"const_fake_quant",
                                     [SameOperandsAndResultType, NoSideEffect]> {
-  let summary =
-      "Simulates the effect of uniform quantization with const range.";
+  let summary = [{
+    Simulates the effect of uniform quantization with const range.
+  }];
 
   let description = [{
     Given a const min, max, num_bits and narrow_range attribute, applies the
     same uniform quantization simulation as is done by the TensorFlow
     fake_quant_with_min_max_args op. See the fakeQuantAttrsToType() utility
     method and the quant-convert-simulated-quantization pass for further details.
   }];
 
def quant_ConstFakeQuant : quant_Op<"const_fake_quant",
 
   let results = (outs
     F32Tensor:$outputs
   );
 }
 
 def quant_ConstFakeQuantPerAxis : quant_Op<"const_fake_quant_per_axis",
                                     [SameOperandsAndResultType, NoSideEffect]> {
-  let summary =
-      "Simulates the effect of per axis uniform quantization with const range.";
+  let summary = [{
+    Simulates the effect of per axis uniform quantization with const range.
+  }];
 
   let description = [{
     Given a const min, max, num_bits and narrow_range attribute, applies the
     same per axis uniform quantization simulation as is done by the TensorFlow
     fake_quant_with_min_max_vars_per_channel op. See the fakeQuantAttrsToType()
     utility method and the quant-convert-simulated-quantization pass for further
     details.
   }];
def quant_ConstFakeQuantPerAxis : quant_Op<"const_fake_quant_per_axis",
   );
 
   let results = (outs
     F32Tensor:$outputs
   );
 }
 
 def quant_StatisticsRefOp : quant_Op<"stats_ref", [SameOperandsAndResultType]> {
-  let summary =
-      "Indicates that statistics are resolved by reference.";
+  let summary = "Indicates that statistics are resolved by reference.";
 
   let description = [{
     This op acts as an identity that, when encountered at runtime, should result
     in statistics being collected about about the value of its operand/result.
     Such statistics will be stored with the provided key, allowing this node
     to later be converted to a 'stats' op if statistics with that key have been
     encountered.
   }];
 
   let arguments = (ins
     quant_RealValueType:$arg,
     StrAttr:$statsKey
   );
   let results = (outs quant_RealValueType);
 }
 
 def quant_StatisticsOp : quant_Op<"stats", [SameOperandsAndResultType]> {
-  let summary =
-      "Identity op which associates statistics with the value.";
+  let summary = "Identity op which associates statistics with the value.";
 
   let description = [{
     Associates statistics about the runtime ranges of values observed for
     evaluations of this node.
 
     Statistics about the entire type are reported in the 'layerStats' attribute
     and those for each axis, in the (optional) `axisStats` attribute. The
     interpretation of each is determined by the last dimension of its shape.
     Currently, only dim=2 is supported, which is interpreted as [min, max].
 
     `layerStats` must be a rank 1 tensor: [2]
     `axisStats` must be a rank 2 tensor: [N, 2], where N=the slice size
       splitted by the `axis` dimension. For example:
-      <?x?x3x2>, axis=3 => N=2
-      <?x?x3x2>, axis=2 => N=6
+
+    ```
+    ```
   }];
 
   let arguments = (ins
     quant_RealValueType:$arg,
     ElementsAttr:$layerStats,
     OptionalAttr<ElementsAttr>:$axisStats,
     OptionalAttr<I64Attr>:$axis);
   let results = (outs quant_RealValueType);
if (axisStats()) {
                            "where N = the slice size defined by the axis dim");
       }
     }
     return success();
   }];
 }
 
 def quant_CoupledRefOp : quant_Op<"coupled_ref", [SameOperandsAndResultType]> {
-  let summary =
-      "Indicates that one point of the computation is coupled to another.";
+  let summary = [{
+    Indicates that one point of the computation is coupled to another.
+  }];
 
   let description = [{
     Ordinarily, relationships between ops for the purposes of determining
     compatible quantized types is explicit based on the use-def chain. However,
     in some situations, a use may be separated from its def by arbitrary
     external connections. In such a case, during analysis, all coupled_ref
     nodes in a module which share a coupledKey will be considered to be
     directly connected as via an identity op for the purpose of type inference.

mlir/include/mlir/Dialect/Shape/IR/CMakeLists.txt

-set(LLVM_TARGET_DEFINITIONS ShapeOps.td)
-mlir_tablegen(ShapeOps.h.inc -gen-op-decls)
-mlir_tablegen(ShapeOps.cpp.inc -gen-op-defs)
-mlir_tablegen(ShapeOpsDialect.h.inc -gen-dialect-decls)
-add_public_tablegen_target(MLIRShapeOpsIncGen)
+add_mlir_dialect(ShapeOps shape)
+add_mlir_doc(ShapeOps -gen-dialect-doc ShapeDialect Dialects/)

mlir/include/mlir/Dialect/Vector/VectorOps.td

*) "parallel": Batch dimensions are iterator type "parallel", and
                    lhs or rhs (but not both). The lhs and rhs free dimensions
                    are unrelated to each other and do not co-iterate, which
                    should be expressed in their indexing maps.
 
     An indexing map attribute list must be specified with an entry for lhs, rhs
     and acc arguments. An indexing map attribute specifies a mapping from each
     iterator in the iterator type list, to each dimension of an N-D vector.
 
-    Examples:
+    Example:
+
     ```mlir
-      // Simple dot product (K = 0).
-      #contraction_accesses = [
-       affine_map<(i) -> (i)>,
-       affine_map<(i) -> (i)>,
-       affine_map<(i) -> ()>
-      ]
-      #contraction_trait = {
-        indexing_maps = #contraction_accesses,
-        iterator_types = ["reduction"]
-      }
-      %3 = vector.contract #contraction_trait %0, %1, %2
-        : vector<10xf32>, vector<10xf32> into f32
 
-      // 2D vector contraction with one contracting dimension (matmul, K = 2).
-      #contraction_accesses = [
-        affine_map<(i, j, k) -> (i, k)>,
-        affine_map<(i, j, k) -> (k, j)>,
-        affine_map<(i, j, k) -> (i, j)>
-      ]
-      #contraction_trait = {
-        indexing_maps = #contraction_accesses,
-        iterator_types = ["parallel", "parallel", "reduction"]
-      }
 
-      %3 = vector.contract #contraction_trait %0, %1, %2
-        : vector<4x3xf32>, vector<3x7xf32> into vector<4x7xf32>
 
-      // 4D to 3D vector contraction with two contracting dimensions and
-      // one batch dimension (K = 3).
-      #contraction_accesses = [
-        affine_map<(b0, f0, f1, c0, c1) -> (c0, b0, c1, f0)>,
-        affine_map<(b0, f0, f1, c0, c1) -> (b0, c1, c0, f1)>,
-        affine_map<(b0, f0, f1, c0, c1) -> (b0, f0, f1)>
-      ]
-      #contraction_trait = {
-        indexing_maps = #contraction_accesses,
-        iterator_types = ["parallel", "parallel", "parallel",
-                          "reduction", "reduction"]
-      }
 
-      %4 = vector.contract #contraction_trait %0, %1, %2
-          : vector<7x8x16x15xf32>, vector<8x16x7x5xf32> into vector<8x15x5xf32>
 
-      // 4D vector contraction with two contracting dimensions and optional
-      // vector mask arguments.
-      %lhs_mask = vector.constant_mask [7, 8, 16, 15] : vector<7x8x16x15xi1>
-      %rhs_mask = vector.constant_mask [8, 16, 7, 5] : vector<8x16x7x5xi1>
 
-      %5 = vector.contract #contraction_trait %0, %1, %2, %lhs_mask, %rhs_mask
-         : vector<7x8x16x15xf32>, vector<8x16x7x5xf32> into vector<8x15x8x5xf32>
     ```
   }];
   let builders = [OpBuilder<
     "Builder *builder, OperationState &result, Value lhs, Value rhs, "
     "Value acc, ArrayAttr indexingMaps, ArrayAttr iteratorTypes">,
     OpBuilder<
       "Builder *builder, OperationState &result, Value lhs, Value rhs, "
       "Value acc, ArrayRef<ArrayRef<AffineExpr>> indexingExprs, "
let description = [{
     Some reductions (add/mul for fp) also allow an optional fused
     accumulator.
 
     Note that these operations are restricted to 1-D vectors to remain
     close to the corresponding LLVM intrinsics:
 
     http://llvm.org/docs/LangRef.html#experimental-vector-reduction-intrinsics
 
-    Examples:
+    Example:
+
     ```mlir
-      %1 = vector.reduction "add", %0 : vector<16xf32> into f32
 
-      %3 = vector.reduction "xor", %2 : vector<4xi32> into i32
 
-      %4 = vector.reduction "mul", %0, %1 : vector<16xf32> into f32
     ```
   }];
   let extraClassDeclaration = [{
     VectorType getVectorType() {
       return vector().getType().cast<VectorType>();
     }
   }];
 }
* a k-D vector <s_1 x .. x s_k x type> can be broadcast to
                                s_1     x .. x s_j x .. x s_k
                <duplication>         <potential stretch>
        ```
     The source operand is duplicated over all the missing leading dimensions
     and stretched over the trailing dimensions where the source has a non-equal
     dimension of 1. These rules imply that any scalar broadcast (k=0) to any
     shaped vector with the same element type is always legal.
 
-    Examples:
+    Example:
+
     ```mlir
-      %0 = constant 0.0 : f32
-      %1 = vector.broadcast %0 : f32 to vector<16xf32>
-      %2 = vector.broadcast %1 : vector<16xf32> to vector<4x16xf32>
     ```
   }];
   let extraClassDeclaration = [{
     Type getSourceType() { return source().getType(); }
     VectorType getVectorType() {
       return vector().getType().cast<VectorType>();
     }
   }];
* the two operands and the result must have the same rank and trailing
               v1 : <s_1 x s_2 x .. x s_k x type> and
               v2 : <t_1 x t_2 x .. x t_k x type>
       we have s_i = t_i for all 1 < i <= k
     * the mask length equals the leading dimension size of the result
     * numbering the input vector indices left to right across the operands, all
       mask values must be within range, viz. given two k-D operands v1 and v2
       above, all mask values are in the range [0,s_1+t_1)
 
-    Examples:
+    Example:
+
     ```mlir
     %0 = vector.shuffle %a, %b[0, 3]
                : vector<2xf32>, vector<2xf32>       ; yields vector<2xf32>
     %1 = vector.shuffle %c, %b[0, 1, 2]
                : vector<2x16xf32>, vector<1x16xf32> ; yields vector<3x16xf32>
     %2 = vector.shuffle %a, %b[3, 2, 1, 0]
                : vector<2xf32>, vector<2xf32>       ; yields vector<4xf32>
-
     ```
   }];
   let builders = [OpBuilder<"Builder *builder, OperationState &result,"
                             "Value v1, Value v2, ArrayRef<int64_t>">];
   let extraClassDeclaration = [{
     static StringRef getMaskAttrName() { return "mask"; }
     VectorType getV1VectorType() {
       return v1().getType().cast<VectorType>();
def Vector_ExtractElementOp :
   let description = [{
     Takes an 1-D vector and a dynamic index position and extracts the
     scalar at that position. Note that this instruction resembles
     vector.extract, but is restricted to 1-D vectors and relaxed
     to dynamic indices. It is meant to be closer to LLVM's version:
     https://llvm.org/docs/LangRef.html#extractelement-instruction
 
     Example:
+
     ```mlir
-      %c = constant 15 : i32
-      %1 = vector.extractelement %0[%c : i32]: vector<16xf32>
     ```
   }];
   let extraClassDeclaration = [{
     VectorType getVectorType() {
       return vector().getType().cast<VectorType>();
     }
   }];
 
   let assemblyFormat = [{
     $vector `[` $position `:` type($position) `]` attr-dict `:` type($vector)
   }];
 }
 
 def Vector_ExtractOp :
   Vector_Op<"extract", [NoSideEffect,
      PredOpTrait<"operand and result have same element type",
                  TCresVTEtIsSameAsOpBase<0, 0>>]>,
     Arguments<(ins AnyVector:$vector, I64ArrayAttr:$position)>,
     Results<(outs AnyType)> {
   let summary = "extract operation";
   let description = [{
     Takes an n-D vector and a k-D position and extracts the (n-k)-D vector at
     the proper position. Degenerates to an element type in the 0-D case.
 
-    Examples:
+    Example:
+
     ```mlir
-      %1 = vector.extract %0[3]: vector<4x8x16xf32>
-      %2 = vector.extract %0[3, 3, 3]: vector<4x8x16xf32>
     ```
   }];
   let builders = [OpBuilder<
     "Builder *builder, OperationState &result, Value source,"
     "ArrayRef<int64_t>">];
   let extraClassDeclaration = [{
     static StringRef getPositionAttrName() { return "position"; }
     VectorType getVectorType() {
let description = [{
     generating the unrolling of 'vector' shape, which has all slices of shape
     'sizes' except for slices at dimension boundaries when 'vector' dimension
     sizes are not a multiple of 'sizes'.
 
     Each slice is returned at the tuple element index corresponding to the
     linear index of the slice w.r.t the unrolling scheme represented by 'sizes'.
     Currently, only unit strides are supported.
 
-    Examples:
+    Example:
+
     ```mlir
-      %0 = vector.transfer_read ...: vector<4x2xf32>
 
-      %1 = vector.extract_slices %0, [2, 2], [1, 1]
-        : vector<4x2xf32> into tuple<vector<2x2xf32>, vector<2x2xf32>>
 
-      // Example with partial slices at dimension boundaries.
-      %2 = vector.transfer_read ...: vector<4x3xf32>
 
-      %3 = vector.extract_slices %2, [2, 2], [1, 1]
-        : vector<4x3xf32> into tuple<vector<2x2xf32>, vector<2x1xf32>,
-                                     vector<2x2xf32>, vector<2x1xf32>>
     ```
   }];
   let builders = [OpBuilder<
     "Builder *builder, OperationState &result, TupleType tupleType, " #
     "Value vector, ArrayRef<int64_t> sizes, " #
     "ArrayRef<int64_t> strides">];
   let extraClassDeclaration = [{
     VectorType getSourceVectorType() {
let description = [{
     pointwise multiply-and-accumulate: `$result = `$lhs * $rhs + $acc`.
     All operands and result have the same vector type. The semantics
     of the operation correspond to those of the `llvm.fma`
     [intrinsic](https://llvm.org/docs/LangRef.html#int-fma). In the
     particular case of lowering to LLVM, this is guaranteed to lower
     to the `llvm.fma.*` intrinsic.
 
     Example:
+
     ```mlir
-      %3 = vector.fma %0, %1, %2: vector<8x16xf32>
     ```
   }];
   // Fully specified by traits.
   let verifier = ?;
   let assemblyFormat = "$lhs `,` $rhs `,` $acc attr-dict `:` type($lhs)";
   let builders = [OpBuilder<
     "Builder *b, OperationState &result, Value lhs, Value rhs, Value acc",
     "build(b, result, lhs.getType(), lhs, rhs, acc);">];
let description = [{
     Takes a scalar source, an 1-D destination vector and a dynamic index
     position and inserts the source into the destination at the proper
     position.  Note that this instruction resembles vector.insert, but
     is restricted to 1-D vectors and relaxed to dynamic indices. It is
     meant to be closer to LLVM's version:
     https://llvm.org/docs/LangRef.html#insertelement-instruction
 
     Example:
+
     ```mlir
-      %c = constant 15 : i32
-      %f = constant 0.0f : f32
-      %1 = vector.insertelement %f, %0[%c : i32]: vector<16xf32>
     ```
   }];
   let extraClassDeclaration = [{
     Type getSourceType() { return source().getType(); }
     VectorType getDestVectorType() {
       return dest().getType().cast<VectorType>();
     }
   }];
Vector_Op<"insert", [NoSideEffect,
      Arguments<(ins AnyType:$source, AnyVector:$dest, I64ArrayAttr:$position)>,
      Results<(outs AnyVector:$res)> {
   let summary = "insert operation";
   let description = [{
     Takes an n-D source vector, an (n+k)-D destination vector and a k-D position
     and inserts the n-D source into the (n+k)-D destination at the proper
     position. Degenerates to a scalar source type when n = 0.
 
-    Examples:
+    Example:
+
     ```mlir
-      %2 = vector.insert %0, %1[3]:
-        vector<8x16xf32> into vector<4x8x16xf32>
-      %5 = vector.insert %3, %4[3, 3, 3]:
-        f32 into vector<4x8x16xf32>
+    %2 = vector.insert %0, %1[3] : vector<8x16xf32> into vector<4x8x16xf32>
+    %5 = vector.insert %3, %4[3, 3, 3] : f32 into vector<4x8x16xf32>
     ```
   }];
   let assemblyFormat = [{
     $source `,` $dest $position attr-dict `:` type($source) `into` type($dest)
   }];
 
   let builders = [OpBuilder<
     "Builder *builder, OperationState &result, Value source, " #
let description = [{
     generating the unrolling of 'vector' shape, which has all slices of shape
     'sizes' except for slices at dimension boundaries when 'vector' dimension
     sizes are not a multiple of 'sizes'.
 
     Each slice in 'vectors' is at the tuple element index corresponding to the
     linear index of the slice w.r.t the unrolling scheme represented by 'sizes'.
     Currently, only unit strides are supported.
 
-    Examples:
+    Example:
+
     ```mlir
-      %0 = vector.extract_slices %0, [2, 2], [1, 1]
-        : vector<4x2xf32> into tuple<vector<2x2xf32>, vector<2x2xf32>>
 
-      %1 = vector.insert_slices %0, [2, 2], [1, 1]
-        : tuple<vector<2x2xf32>, vector<2x2xf32>> into vector<4x2xf32>
 
-      // Example with partial slices at dimension boundaries.
-      %3 = vector.extract_slices %2, [2, 2], [1, 1]
-        : vector<4x3xf32> into tuple<vector<2x2xf32>, vector<2x1xf32>,
-                                     vector<2x2xf32>, vector<2x1xf32>>
 
-      %4 = vector.insert_slices %3, [2, 2], [1, 1]
-        : tuple<vector<2x2xf32>, vector<2x1xf32>,
-                vector<2x2xf32>, vector<2x1xf32>> into vector<4x3xf32>
     ```
   }];
 
   let extraClassDeclaration = [{
     TupleType getSourceTupleType() {
       return vectors().getType().cast<TupleType>();
     }
     VectorType getResultVectorType() {
let description = [{
     into the n-D destination vector.
 
     At the moment strides must contain only 1s.
 
     Returns an n-D vector that is a copy of the n-D destination vector in which
     the last k-D dimensions contain the k-D source vector elements strided at
     the proper location as specified by the offsets.
 
-    Examples:
+    Example:
+
     ```mlir
-      %2 = vector.insert_strided_slice %0, %1
-          {offsets = [0, 0, 2], strides = [1, 1]}:
-        vector<2x4xf32> into vector<16x4x8xf32>
     ```
   }];
 
   let assemblyFormat = [{
     $source `,` $dest attr-dict `:` type($source) `into` type($dest)
   }];
 
   let builders = [OpBuilder<
let description = [{
 
     An optional extra 2-D vector argument may be specified in which case the
     operation returns the sum of the outer-product and the extra vector. In this
     multiply-accumulate scenario, the rounding mode is that obtained by
     guaranteeing that a fused-multiply add operation is emitted. When lowered to
     the LLVMIR dialect, this form emits `llvm.intr.fma`, which is guaranteed to
     lower to actual `fma` instructions on x86.
 
-    Examples:
-    ```mlir
-      %2 = vector.outerproduct %0, %1: vector<4xf32>, vector<8xf32>
-      return %2: vector<4x8xf32>
+    Example:
+
+    ```
 
-      %3 = vector.outerproduct %0, %1, %2:
-        vector<4xf32>, vector<8xf32>, vector<4x8xf32>
-      return %3: vector<4x8xf32>
     ```
   }];
   let extraClassDeclaration = [{
     VectorType getOperandVectorTypeLHS() {
       return lhs().getType().cast<VectorType>();
     }
     VectorType getOperandVectorTypeRHS() {
       return rhs().getType().cast<VectorType>();
let description = [{
     across fixed vector shapes. For example, if a vector has a valid data
     shape [6] with fixed vector size [8], then the valid data elements are
     assumed to be stored at the beginning of the vector with the remaining
     vector elements undefined.
 
     In the examples below, valid data elements are represented by an alphabetic
     character, and undefined data elements are represented by '-'.
 
-    Example:
-    ```mlir
+    Example
+
       vector<1x8xf32> with valid data shape [6], fixed vector sizes [8]
 
                 input: [a, b, c, d, e, f]
 
            layout map: (d0) -> (d0 floordiv 8, d0 mod 8)
 
         vector layout: [a, b, c, d, e, f, -, -]
 
-    ```
-    Example:
-    ```mlir
+    Example
+
       vector<2x8xf32> with valid data shape [10], fixed vector sizes [8]
 
                 input: [a, b, c, d, e, f, g, h, i, j]
 
            layout map: (d0) -> (d0 floordiv 8, d0 mod 8)
 
         vector layout: [[a, b, c, d, e, f, g, h],
                         [i, j, -, -, -, -, -, -]]
-    ```
-    Example:
-    ```mlir
+
+    Example
+
       vector<2x2x2x3xf32> with valid data shape [3, 5], fixed vector sizes
       [2, 3]
 
                 input: [[a, b, c, d, e],
                         [f, g, h, i, j],
                         [k, l, m, n, o]]
 
            layout map: (d0, d1) -> (d0 floordiv 3, d1 floordiv 5,
                                     d0 mod 3, d1 mod 5)
 
         vector layout: [[[[a, b, c],
                           [f, g, h]]
                          [[d, e, -],
                           [i, j, -]]],
                         [[[k, l, m],
                           [-, -, -]]
                          [[n, o, -],
                           [-, -, -]]]]
-    ```
-    Example:
-    ```mlir
+
+    Example
+
       %1 = vector.reshape %0, [%c3, %c6], [%c2, %c9], [4]
         : vector<3x2x4xf32> to vector<2x3x4xf32>
 
              input: [[a, b, c, d, e, f],
                      [g, h, i, j, k, l],
                      [m, n, o, p, q, r]]
 
         layout map: (d0, d1) -> (d0, d1 floordiv 4, d1 mod 4)
 
 
       Input vector:  [[[a, b, c, d],
                        [e, f, -, -]],
                       [[g, h, i, j],
                        [k, l, -, -]],
                       [[m, n, o, p],
                        [q, r, -, -]]]
 
       Output vector:  [[[a, b, c, d],
                         [e, f, g, h],
                         [i, -, -, -]],
                        [[j, k, l, m],
                         [n, o, p, q],
                         [r, -, -, -]]]
-    ```
   }];
 
   let extraClassDeclaration = [{
     VectorType getInputVectorType() {
       return vector().getType().cast<VectorType>();
     }
     VectorType getOutputVectorType() {
       return getResult().getType().cast<VectorType>();
let description = [{
 
     At the moment strides must contain only 1s.
     // TODO(ntv) support non-1 strides.
 
     Returns an n-D vector where the first k-D dimensions match the `sizes`
     attribute. The returned subvector contains the elements starting at offset
     `offsets` and ending at `offsets + sizes`.
 
-    Examples:
+    Example:
+
     ```mlir
-      %1 = vector.strided_slice %0
-          {offsets = [0, 2], sizes = [2, 4], strides = [1, 1]}:
-        vector<4x8x16xf32> to vector<2x4x16xf32>
-    ```
 
     // TODO(ntv) Evolve to a range form syntax similar to:
     %1 = vector.strided_slice %0[0:2:1][2:4:1]
       vector<4x8x16xf32> to vector<2x4x16xf32>
+    ```
   }];
   let builders = [OpBuilder<
     "Builder *builder, OperationState &result, Value source, " #
     "ArrayRef<int64_t> offsets, ArrayRef<int64_t> sizes, " #
     "ArrayRef<int64_t> strides">];
   let extraClassDeclaration = [{
     static StringRef getOffsetsAttrName() { return "offsets"; }
     static StringRef getSizesAttrName() { return "sizes"; }
let description = [{
 
     where `broadcast` broadcasts from element 0 to all others along the
     specified dimension. This time, the temporary storage footprint is `3 * 5`
     values which is the same amount of data as the `3 * 5` values transferred.
     An additional `1` broadcast is required. On a GPU this broadcast could be
     implemented using a warp-shuffle if loop `j` were mapped to `threadIdx.x`.
 
     Syntax
-    ```mlir
+    ```
     operation ::= ssa-id `=` `vector.transfer_read` ssa-use-list
       `{` attribute-entry `} :` memref-type `,` vector-type
     ```
 
-    Examples:
+    Example:
+
     ```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
     for %i0 = 0 to %0 {
       affine.for %i1 = 0 to %1 step 256 {
         affine.for %i2 = 0 to %2 step 32 {
           %v = vector.transfer_read %A[%i0, %i1, %i2], (%f0)
let description = [{
     'store' because the super-vector granularity is generally not representable
     with a single hardware register. A `vector.transfer_write` is thus a
     mid-level abstraction that supports super-vectorization with non-effecting
     padding for full-tile-only code. It is the responsibility of
     `vector.transfer_write`'s implementation to ensure the memory writes are
     valid. Different lowerings may be pertinent depending on the hardware
     support.
 
-    Syntax:
-
-    ```mlir
-    operation ::= `vector.transfer_write` ssa-use-list `{` attribute-entry `} :
-      ` vector-type ', ' memref-type '
-    ```
-
-    Examples:
+    Example:
 
     ```mlir
     // write vector<16x32x64xf32> into the slice
     //   `%A[%i0, %i1:%i1+32, %i2:%i2+64, %i3:%i3+16]`:
     for %i0 = 0 to %0 {
       affine.for %i1 = 0 to %1 step 32 {
         affine.for %i2 = 0 to %2 step 64 {
           affine.for %i3 = 0 to %3 step 16 {
let description = [{
     It is currently assumed that this operation does not require moving data,
     and that it will be folded away before lowering vector operations.
 
     There is an exception to the folding expectation when targeting
     llvm.intr.matrix operations. We need a type conversion back and forth from a
     2-D MLIR vector to a 1-D flattened LLVM vector.shape_cast lowering to LLVM
     is supported in that particular case, for now.
 
-    Examples:
+    Example:
 
     ```mlir
     // Example casting to a lower vector rank.
     %1 = vector.shape_cast %0 : vector<5x1x4x3xf32> to vector<20x3xf32>
 
     // Example casting to a higher vector rank.
     %3 = vector.shape_cast %2 : vector<10x12x8xf32> to vector<5x2x3x4x8xf32>
 
let description = [{
     Performs a conversion from a memref with scalar element to a memref with a
     *single* vector element, copying the shape of the memref to the vector. This
     is the minimal viable operation that is required to makeke
     super-vectorization operational. It can be seen as a special case of the
     `view` operation but scoped in the super-vectorization context.
 
     Syntax:
 
-    ```mlir
+    ```
     operation ::= `vector.type_cast` ssa-use : memref-type to memref-type
     ```
 
     Example:
 
     ```mlir
     %A  = alloc() : memref<5x4x3xf32>
     %VA = vector.type_cast %A : memref<5x4x3xf32> to memref<vector<5x4x3xf32>>
let description = [{
     within a hyper-rectangular region specified by the 'mask_dim_sizes'
     array attribute argument. Each element of the 'mask_dim_sizes' array,
     specifies an exclusive upper bound [0, mask-dim-size-element-value)
     for a unique dimension in the vector result. The conjunction of the ranges
     define a hyper-rectangular region within which elements values are set to 1
     (otherwise element values are set to 0).
 
     Example:
-    ```
-      create a constant vector mask of size 4x3xi1 with elements in range
-      0 <= row <= 2 and 0 <= col <= 1 are set to 1 (others to 0).
 
-      %1 = vector.constant_mask [3, 2] : vector<4x3xi1>
+    ```mlir
+    // create a constant vector mask of size 4x3xi1 with elements in range
+    // 0 <= row <= 2 and 0 <= col <= 1 are set to 1 (others to 0).
 
-      print %1
-                    columns
-                  0    1    2
-                |------------
-              0 | 1    1    0
-        rows  1 | 1    1    0
-              2 | 1    1    0
-              3 | 0    0    0
     ```
   }];
 
   let extraClassDeclaration = [{
     static StringRef getMaskDimSizesAttrName() { return "mask_dim_sizes"; }
   }];
   let assemblyFormat = "$mask_dim_sizes attr-dict `:` type(results)";
 }
 
 def Vector_CreateMaskOp :
   Vector_Op<"create_mask", [NoSideEffect]>,
     Arguments<(ins Variadic<Index>:$operands)>, Results<(outs VectorOf<[I1]>)> {
   let summary = "creates a vector mask";
   let description = [{
     Creates and returns a vector mask where elements of the result vector
     are set to '0' or '1', based on whether the element indices are contained
     within a hyper-rectangular region specified by the operands. Specifically,
     each operand specifies a range [0, operand-value) for a unique dimension in
     the vector result. The conjunction of the operand ranges define a
     hyper-rectangular region within which elements values are set to 1
     (otherwise element values are set to 0).
 
     Example:
-    ```
-      create a vector mask of size 4x3xi1 where elements in range
-      0 <= row <= 2 and 0 <= col <= 1 are set to 1 (others to 0).
 
-      %1 = vector.create_mask %c3, %c2 : vector<4x3xi1>
+    ```mlir
+    // create a vector mask of size 4x3xi1 where elements in range
+    // 0 <= row <= 2 and 0 <= col <= 1 are set to 1 (others to 0).
 
-      print %1
-                    columns
-                  0    1    2
-                |------------
-              0 | 1    1    0
-        rows  1 | 1    1    0
-              2 | 1    1    0
-              3 | 0    0    0
     ```
   }];
 
   let hasCanonicalizer = 1;
   let assemblyFormat = "$operands attr-dict `:` type(results)";
 }
 
 def Vector_TupleOp :
   Vector_Op<"tuple", [NoSideEffect]>,
     Arguments<(ins Variadic<AnyVector>:$vectors)>,
     Results<(outs TupleOf<[AnyVector]>)> {
   let summary = "make tuple of vectors operation";
   let description = [{
     Returns a tuple of its operands 'vectors'.
 
     Note that this operation is used during the vector op unrolling
     transformation and should be removed before lowering to lower-level
     dialects.
 
-    Examples:
+
+    Example:
+
     ```mlir
-      %0 = vector.transfer_read ... : vector<2x2xf32>
-      %1 = vector.transfer_read ... : vector<2x1xf32>
-      %2 = vector.transfer_read ... : vector<2x2xf32>
-      %3 = vector.transfer_read ... : vector<2x1xf32>
 
-      %4 = vector.tuple %0, %1, %2, %3
-        : vector<2x2xf32>, vector<2x1xf32>, vector<2x2xf32>, vector<2x1xf32>
-
     ```
   }];
 
   let extraClassDeclaration = [{
     TupleType getResultTupleType() {
       return getResult().getType().cast<TupleType>();
     }
   }];
 }
 
 def Vector_TransposeOp :
   Vector_Op<"transpose", [NoSideEffect,
     PredOpTrait<"operand and result have same element type",
                  TCresVTEtIsSameAsOpBase<0, 0>>]>,
     Arguments<(ins AnyVector:$vector, I64ArrayAttr:$transp)>,
     Results<(outs AnyVector:$result)> {
   let summary = "vector transpose operation";
   let description = [{
     Takes a n-D vector and returns the transposed n-D vector defined by
     the permutation of ranks in the n-sized integer array attribute.
     In the operation
+
     ```mlir
-      %1 = vector.transpose %0, [i_1, .., i_n]
-        : vector<d_1 x .. x d_n x f32>
-        to vector<d_trans[0] x .. x d_trans[n-1] x f32>
     ```
+
     the transp array [i_1, .., i_n] must be a permutation of [0, .., n-1].
 
     Example:
+
     ```mlir
     %1 = vector.transpose %0, [1, 0] : vector<2x3xf32> to vector<3x2xf32>
 
      [ [a, b, c],       [ [a, d],
        [d, e, f] ]  ->    [b, e],
                           [c, f] ]
     ```
   }];
def Vector_TupleGetOp :
   let summary = "vector tuple get operation";
   let description = [{
     Returns the tuple element of 'vectors' at 'index'.
 
     Note that this operation is used during the vector op unrolling
     transformation and should be removed before lowering to lower-level
     dialects.
 
-    Examples:
+    Example:
+
     ```mlir
-      %4 = vector.tuple %0, %1, %2, %3
-        : vector<2x2xf32>, vector<2x1xf32>, vector<2x2xf32>, vector<2x1xf32>>
 
-      %5 = vector.tuple_get %4, 1
-        : tuple<vector<2x2xf32>, vector<2x1xf32>,
-                vector<2x2xf32>, vector<2x1xf32>>
     ```
   }];
 
   let extraClassDeclaration = [{
     VectorType getResultVectorType() {
       return getResult().getType().cast<VectorType>();
     }
     int64_t getIndex() {
       return getAttrOfType<IntegerAttr>("index").getValue().getSExtValue();
     }
     static StringRef getIndexAttrName() { return "index"; }
   }];
   let hasFolder = 1;
 }
 
 def Vector_PrintOp :
   Vector_Op<"print", []>, Arguments<(ins AnyType:$source)> {
   let summary = "print operation (for testing and debugging)";
   let description = [{
     Prints the source vector (or scalar) to stdout in human readable
     format (for testing and debugging). No return value.
 
-    Examples:
+    Example:
+
     ```mlir
-      %0 = constant 0.0 : f32
-      %1 = vector.broadcast %0 : f32 to vector<4xf32>
-      vector.print %1 : vector<4xf32>
 
-      when lowered to LLVM, the vector print is unrolled into
-      elementary printing method calls that at runtime will yield
 
-      ( 0.0, 0.0, 0.0, 0.0 )
 
-      on stdout when linked with a small runtime support library,
-      which only needs to provide a few printing methods (single
-      value for all data types, opening/closing bracket, comma,
-      newline).
     ```
   }];
   let verifier = ?;
   let extraClassDeclaration = [{
     Type getPrintType() {
       return source().getType();
     }
   }];
let description = [{
     The ‘vector.matrix_multiply’ op treats `lhs` as matrix with <lhs_rows> rows
     and <lhs_columns> columns, `rhs` as matrix with <lhs_columns> rows and
     <rhs_columns> and multiplies them. The result matrix is returned embedded in
     the result vector.
 
     Example:
 
     ```mlir
-      %C = vector.matrix_multiply %A, %B
-        { lhs_rows = 4: i32, lhs_columns = 16: i32 , rhs_columns = 3: i32 } :
-        (vector<64xf64>, vector<48xf64>) -> vector<12xf64>
     ```
   }];
   let builders = [
    OpBuilder<"Builder *builder, OperationState &result, Value lhs, Value rhs, "
              "unsigned lhsRows, unsigned lhsColumns, unsigned rhsColumns",
    [{
      result.addOperands({lhs, rhs});
      result.addAttribute("lhs_rows", builder->getI32IntegerAttr(lhsRows));