[mlir][linalg] Tiling implementation for `linalg.generic` will result in assert/crash if indexing expressions have negative multiplicative coefficients

[christopherbate]

Any indexing map that contains negative multiplicative coefficients will crash due to an assertion here. Such maps are sometimes produced to represent reversed iteration over some input dimension.

The condition should be checked much higher up in the call stack (e.g. in Linalg/Transforms/TilingInterfaceImpl.cpp); I think there's not much awareness of this limitation.

Here's a test case:

#map = affine_map<(d0) -> (d0)>
#map1 = affine_map<(d0) -> (-d0 + 7)>

func.func @test(%arg0: tensor<8xi8>, %arg1: tensor<8xi8>) -> tensor<8xi8> {
  %0 = linalg.generic {
    indexing_maps = [#map1, #map],
    iterator_types = ["parallel"]}
  ins(%arg0 : tensor<8xi8>)
  outs(%arg1 : tensor<8xi8>) {
  ^bb0(%in: i8, %out: i8):
    linalg.yield %in : i8
  } -> tensor<8xi8>
  return %0 : tensor<8xi8>
}


module attributes {transform.with_named_sequence} {
  transform.named_sequence @__transform_main(%arg1: !transform.any_op {transform.readonly}) {
    %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!transform.any_op) -> !transform.any_op
    %1:2 = transform.structured.tile_using_for
      %0 tile_sizes [2] : (!transform.any_op) -> (!transform.any_op, !transform.any_op)
     transform.yield
  }
}

[llvmbot]

@llvm/issue-subscribers-mlir-linalg

Author: Christopher Bate (christopherbate)

Any indexing map that contains negative multiplicative coefficients will crash due to an assertion [here](https://github.com/llvm/llvm-project/blob/main/mlir/lib/Dialect/Linalg/Utils/Utils.cpp#L68). Such maps are sometimes produced to represent reversed iteration over some input dimension.

The condition should be checked much higher up in the call stack (e.g. in Linalg/Transforms/TilingInterfaceImpl.cpp); I think there's not much awareness of this limitation.

Here's a test case:

#map = affine_map<(d0) -> (d0)>
#map1 = affine_map<(d0) -> (-d0 + 7)>

func.func @<!-- -->test(%arg0: tensor&lt;8xi8&gt;, %arg1: tensor&lt;8xi8&gt;) -&gt; tensor&lt;8xi8&gt; {
  %0 = linalg.generic {
    indexing_maps = [#map1, #map],
    iterator_types = ["parallel"]}
  ins(%arg0 : tensor&lt;8xi8&gt;)
  outs(%arg1 : tensor&lt;8xi8&gt;) {
  ^bb0(%in: i8, %out: i8):
    linalg.yield %in : i8
  } -&gt; tensor&lt;8xi8&gt;
  return %0 : tensor&lt;8xi8&gt;
}


module attributes {transform.with_named_sequence} {
  transform.named_sequence @<!-- -->__transform_main(%arg1: !transform.any_op {transform.readonly}) {
    %0 = transform.structured.match ops{["linalg.generic"]} in %arg1 : (!transform.any_op) -&gt; !transform.any_op
    %1:2 = transform.structured.tile_using_for
      %0 tile_sizes [2] : (!transform.any_op) -&gt; (!transform.any_op, !transform.any_op)
     transform.yield
  }
}

[siddhesh195]

A possible place to add a check can be during computation of tiling factor in mlir/lib/Dialect/Linalg/Transforms/Tiling.cpp

[siddhesh195]

I would like to attempt a patch if no one else is working on it

[MaheshRavishankar]

This was just brought to my attention but afaics such a linalg op is illegal. I'll am afk right now, but will explain more in the morning

[christopherbate]

The case mentioned in the description is produced by the StableHLO-to-linalg pass maintained in the StableHLO repo (conversion of "stablehlo.reverse").

Edit: (Not saying it's valid linalg op for that reason, I wasn't aware that was a limitation.) If we can't have negative multiplication coefficients then that seems easy to check in the verifier.... or is the invalid-ness due to something else?