TileAndFuse configuration gives numeric error

[nirvedhmeshram]

For this pre-configured aligned matmul

compile with

iree-compile small_matmul_aligned_withconfig.mlir   -o output_tileandfuse_aligned.vmfb 

and run

iree-run-module --device=hip --module=output_tileandfuse_aligned.vmfb --input=64x64xf32=1 --input=64x16xf32=1

the expected answer is all 64s but we get

EXEC @matmul
result[0]: hal.buffer_view
64x16xf32=[16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16][16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16]
[16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16][16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16][16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16]
[16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16][16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16][16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16]
[16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16][16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16][16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16]
[16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16][16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]
[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0][0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0]

The heuristics will not automatically choose this configuration for aligned shape but they do for unaligned shape where this bug was originally discovered.

Note that this is using c promotion which is new for tile and fuse pipeline. However even for aligned case this configuration is yet to be supported without c promotion. The missing patterns are being added for that.

[nirvedhmeshram]

cc @jerryyin @qedawkins

[nirvedhmeshram]

Here is the IR dump

[nirvedhmeshram]

Update: taking the fix being added in llvm/llvm-project#117340, changes the numeric error where 16 in the above dump becomnes 64 (whihc is the right value, but we still have the 0's) . I thought of doing that becuase I noticed insert_slices's in the IR after pack/unpack decomposition.

[nirvedhmeshram]

I think I do see the source of the bug and it is being added in the thread level tiling, it generates this scf.forall

      %13 = scf.forall (%arg4, %arg5) = (0, 0) to (16, 4) step (1, 4) shared_outs(%arg6 = %12) -> (tensor<16x4xf32>) {
        %18 = affine.apply affine_map<()[s0, s1] -> (s0 + s1)>()[%arg4, %arg0]
        %19 = affine.apply affine_map<(d0)[s0] -> (d0 * 4 + s0)>(%arg2)[%arg5]
        %extracted_slice = tensor.extract_slice %3[%18, %19] [1, 4] [1, 1] : tensor<64x64xf32> to tensor<1x4xf32>
        %extracted_slice_3 = tensor.extract_slice %arg6[%arg4, %arg5] [1, 4] [1, 1] : tensor<16x4xf32> to tensor<1x4xf32>
        %20 = linalg.copy {lowering_config = #iree_gpu.derived_thread_config} ins(%extracted_slice : tensor<1x4xf32>) outs(%extracted_slice_3 : tensor<1x4xf32>) -> tensor<1x4xf32>
        scf.forall.in_parallel {
          tensor.parallel_insert_slice %20 into %arg6[%arg4, %arg5] [1, 4] [1, 1] : tensor<1x4xf32> into tensor<16x4xf32>
        }
      } {mapping = [#gpu.thread<linear_dim_1>, #gpu.thread<linear_dim_0>]}

Notice that it has a step and an indexing offset, so its stepping by 16 when it wants to step by 4, causing the bug.

EDIT: Actually on a second pass, this seems fine, it is however different than a simialr case which works where the thread distributed loop had tile size of 1 vs tile size of 4 here which might be cuasing issue, as we didnt expect that.

[nirvedhmeshram]

So the reaming error was becuase I had |workgroup_size| in translation_info > |workgroup| in lowering_config. In practice I believe such a config doesnt make sense so closing this issue.