Add MLIR support (#1044)

Author: lutzroeder

source/mlir-metadata.json

@@ -1250,10 +1250,24 @@
     "assemblyFormat": "$linear_index `into`\n    custom<DynamicIndexList>($dynamic_basis, $static_basis, \"{}\", \"::mlir::AsmParser::Delimiter::Paren\")\n    attr-dict `:` type($multi_index)"
   },
   {
-    "name": "affine.dma_start"
+    "name": "affine.dma_start",
+    "summary": "affine dma start operation",
+    "description": "The `affine.dma_start` op starts a non-blocking DMA operation that\n    transfers data from a source memref to a destination memref. The source and\n    destination memref need not be of the same dimensionality, but need to have\n    the same elemental type. The operands include the source and destination\n    memref's each followed by its indices, size of the data transfer in terms of\n    the number of elements (of the elemental type of the memref), a tag memref\n    with its indices, and optionally at the end, a stride and a\n    number_of_elements_per_stride arguments. The tag location is used by an\n    `affine.dma_wait` to check for completion. The indices of the source memref,\n    destination memref, and the tag memref have the same restrictions as any\n    affine.load/store. In particular, index for each memref dimension must be an\n    affine expression of loop induction variables and symbols.\n\n    The optional stride arguments should be of 'index' type, and specify a\n    stride for the slower memory space (memory space with a lower memory space\n    id), transferring chunks of number_of_elements_per_stride every stride until\n    %num_elements are transferred. Either both or no stride arguments should be\n    specified. The value of 'num_elements' must be a multiple of\n    'number_of_elements_per_stride'. If the source and destination locations\n    overlap the behavior of this operation is not defined.\n\n    Example:\n\n    ```mlir\n    %num_elements = arith.constant 256\n    %idx = arith.constant 0 : index\n    %tag = memref.alloc() : memref<1xi32, 4>\n    affine.dma_start %src[%i + 3, %j], %dst[%k + 7, %l], %tag[%idx],\n      %num_elements :\n        memref<40x128xf32, 0>, memref<2x1024xf32, 1>, memref<1xi32, 2>\n\n    // If %stride and %num_elt_per_stride are specified, the DMA is expected to\n    // transfer %num_elt_per_stride elements every %stride elements apart from\n    // memory space 0 until %num_elements are transferred.\n    affine.dma_start %src[%i, %j], %dst[%k, %l], %tag[%idx], %num_elements,\n      %stride, %num_elt_per_stride : ...\n    ```",
+    "attributes": [
+      { "name": "src_map", "type": "AffineMapAttr" },
+      { "name": "dst_map", "type": "AffineMapAttr" },
+      { "name": "tag_map", "type": "AffineMapAttr" }
+    ],
+    "hasCustomAssemblyFormat": true
   },
   {
-    "name": "affine.dma_wait"
+    "name": "affine.dma_wait",
+    "summary": "affine dma wait operation",
+    "description": "The `affine.dma_wait` op blocks until the completion of a DMA operation\n    associated with the tag element `%tag[%index]`. `%tag` is a memref, and\n    `%index` has to be an index with the same restrictions as any load/store\n    index. In particular, index for each memref dimension must be an affine\n    expression of loop induction variables and symbols. `%num_elements` is the\n    number of elements associated with the DMA operation.\n\n    Example:\n\n    ```mlir\n    affine.dma_start %src[%i, %j], %dst[%k, %l], %tag[%index], %num_elements :\n      memref<2048xf32, 0>, memref<256xf32, 1>, memref<1xi32, 2>\n    ...\n    affine.dma_wait %tag[%index], %num_elements : memref<1xi32, 2>\n    ```",
+    "attributes": [
+      { "name": "tag_map", "type": "AffineMapAttr" }
+    ],
+    "hasCustomAssemblyFormat": true
   },
   {
     "name": "affine.for",
@@ -27306,7 +27320,7 @@
   {
     "name": "memref.expand_shape",
     "summary": "operation to produce a memref with a higher rank.",
-    "description": "The `memref.expand_shape` op produces a new view with a higher rank whose\n    sizes are a reassociation of the original `view`. The operation is limited\n    to such reassociations, where a dimension is expanded into one or multiple\n    contiguous dimensions. Such reassociations never require additional allocs\n    or copies.\n\n    A reassociation is defined as a grouping of dimensions and is represented\n    with an array of DenseI64ArrayAttr attributes.\n\n    Example:\n\n    ```mlir\n    %r = memref.expand_shape %0 [[0, 1], [2]] output_shape [%sz0, %sz1, 32]\n        : memref<?x32xf32> into memref<?x?x32xf32>\n    ```\n\n    If an op can be statically proven to be invalid (e.g, an expansion from\n    `memref<10xf32>` to `memref<2x6xf32>`), it is rejected by the verifier. If\n    it cannot statically be proven invalid (e.g., the full example above; it is\n    unclear whether the first source dimension is divisible by 5), the op is\n    accepted by the verifier. However, if the op is in fact invalid at runtime,\n    the behavior is undefined.\n\n    The source memref can be zero-ranked. In that case, the reassociation\n    indices must be empty and the result shape may only consist of unit\n    dimensions.\n\n    For simplicity, this op may not be used to cast dynamicity of dimension\n    sizes and/or strides. I.e., if and only if a source dimension is dynamic,\n    there must be a dynamic result dimension in the corresponding reassociation\n    group. Same for strides.\n\n    The representation for the output shape supports a partially-static\n    specification via attributes specified through the `static_output_shape`\n    argument.  A special sentinel value `ShapedType::kDynamic` encodes that the\n    corresponding entry has a dynamic value.  There must be exactly as many SSA\n    inputs in `output_shape` as there are `ShapedType::kDynamic` entries in\n    `static_output_shape`.\n\n    Note: This op currently assumes that the inner strides are of the\n    source/result layout map are the faster-varying ones.",
+    "description": "The `memref.expand_shape` op produces a new view with a higher rank whose\n    sizes are a reassociation of the original `view`. The operation is limited\n    to such reassociations, where a dimension is expanded into one or multiple\n    contiguous dimensions. Such reassociations never require additional allocs\n    or copies.\n\n    A reassociation is defined as a grouping of dimensions and is represented\n    with an array of DenseI64ArrayAttr attributes.\n\n    Example:\n\n    ```mlir\n    %r = memref.expand_shape %0 [[0, 1], [2]] output_shape [%sz0, %sz1, 32]\n        : memref<?x32xf32> into memref<?x?x32xf32>\n    ```\n\n    If an op can be statically proven to be invalid (e.g, an expansion from\n    `memref<10xf32>` to `memref<2x6xf32>`), it is rejected by the verifier. If\n    it cannot statically be proven invalid (e.g., the full example above; it is\n    unclear whether the first source dimension is divisible by 5), the op is\n    accepted by the verifier. However, if the op is in fact invalid at runtime,\n    the behavior is undefined.\n\n    The source memref can be zero-ranked. In that case, the reassociation\n    indices must be empty and the result shape may only consist of unit\n    dimensions.\n\n    For simplicity, this op may not be used to cast dynamicity of dimension\n    sizes and/or strides. I.e., if and only if a source dimension is dynamic,\n    there must be a dynamic result dimension in the corresponding reassociation\n    group. Same for strides.\n\n    The representation for the output shape supports a partially-static\n    specification via attributes specified through the `static_output_shape`\n    argument. A special sentinel value `ShapedType::kDynamic` encodes that the\n    corresponding entry has a dynamic value. Both the number of SSA inputs in\n    `output_shape` and the number of `ShapedType::kDynamic` entries in\n    `static_output_shape` match the number of dynamic dimensions in the result\n    type.\n\n    Note: This op currently assumes that the inner strides are of the\n    source/result layout map are the faster-varying ones.",
     "operands": [
       { "name": "src", "type": "AnyStridedMemRef" },
       { "name": "output_shape", "type": "Variadic<Index>" }
@@ -59923,7 +59937,7 @@
   {
     "name": "tensor.expand_shape",
     "summary": "operation to produce a tensor with a higher rank",
-    "description": "The `tensor.expand_shape` op produces a tensor of higher (or equal)\n    rank than the operand `src` whose dimension sizes are a reassociation of\n    `src`.\n\n    A reassociation is defined as a continuous grouping of dimensions and is\n    represented with an array of DenseI64ArrayAttr attribute.  The reassociation\n    maps applied to the result tensor with the higher rank must result in the\n    operand tensor with the smaller rank.\n\n    The representation for the output shape supports a partially-static\n    specification via attributes specified through the `static_output_shape`\n    argument.  A special sentinel value `ShapedType::kDynamic` encodes that the\n    corresponding entry has a dynamic value.  There must be exactly as many SSA\n    inputs in `output_shape` as there are `ShapedType::kDynamic` entries in\n    `static_output_shape`.\n\n    Example:\n\n    ```mlir\n    // Dimension expansion i -> (i', j') and (k) -> (k')\n    %b = tensor.expand_shape %a [[0, 1], [2]] output_shape [%sz0, %sz1, 32]\n        : tensor<?x32xf32> into tensor<?x?x32xf32>\n    ```",
+    "description": "The `tensor.expand_shape` op produces a tensor of higher (or equal)\n    rank than the operand `src` whose dimension sizes are a reassociation of\n    `src`.\n\n    A reassociation is defined as a continuous grouping of dimensions and is\n    represented with an array of DenseI64ArrayAttr attribute. The reassociation\n    maps applied to the result tensor with the higher rank must result in the\n    operand tensor with the smaller rank.\n\n    The representation for the output shape supports a partially-static\n    specification via attributes specified through the `static_output_shape`\n    argument. A special sentinel value `ShapedType::kDynamic` encodes that the\n    corresponding entry has a dynamic value. Both the number of SSA inputs in\n    `output_shape` and the number of `ShapedType::kDynamic` entries in\n    `static_output_shape` match the number of dynamic dimensions in the result\n    type.\n\n    Example:\n\n    ```mlir\n    // Dimension expansion i -> (i', j') and (k) -> (k')\n    %b = tensor.expand_shape %a [[0, 1], [2]] output_shape [%sz0, %sz1, 32]\n        : tensor<?x32xf32> into tensor<?x?x32xf32>\n    ```",
     "operands": [
       { "name": "src", "type": "AnyRankedTensor" },
       { "name": "output_shape", "type": "Variadic<Index>" }
@@ -116379,6 +116393,15 @@
     ],
     "assemblyFormat": "`(` $sem_addr `,` $val `)` attr-dict `:` functional-type(operands, results)"
   },
+  {
+    "name": "ttkernel.pack_reconfig_l1_acc",
+    "summary": "Reconfigure packer to L1 accumulation mode.",
+    "description": "Reconfigures the packer to accumulate when packing from DST to L1.",
+    "operands": [
+      { "name": "l1_acc_en", "type": "I32" }
+    ],
+    "assemblyFormat": "`(` $l1_acc_en `)` attr-dict `:` functional-type(operands, results)"
+  },
   {
     "name": "ttkernel.pack_tile",
     "summary": "PackTile op.",