Questions about lowering config

[chiachia36]

Hi,

I try to find a proper way of tiling for a model.
I found there are some samples using tile_sizes in lowering_config to set the tile size, but I didn't find the definition of it.
I check the example here to understand the usage of tile_sizes in lowering_config.
I still confuse about the meaning of tile_sizes.
For example, in the sample I mention above
lowering_config = <tile_sizes = [[32, 32], [8, 8, 0], [0, 0, 8]]>,
I want to know what each part represents.
Could you give me some hint or references?

Thanks!

[MaheshRavishankar]

Hey @chiachia36 . Sorry for the late response. This is not a forum I personally monitor, so I missed it. This is a bit of a "ninja" feature, in the sense this was added to unblock some development/performance tuning, but with the hope that there will be a more long term stable solution (which is in the works). Having said that the meaning of the tile sizes list depends on the pipeline that is used for lowering. If you look at the entire attribute

#compilation0 = #iree_codegen.compilation_info<
    lowering_config = <tile_sizes = [[32, 32], [8, 8, 0], [0, 0, 8]]>,
    translation_info = <CPUDoubleTilingPadExpert>,
    workgroup_size = []>

this attribute is specifying the use of the CPUDoubleTilingPadExpert. That name only makes sense to people deeply involved in code-generation work, but broadly it describes a certain sequence of transformations that take the operation and lower them into scalar+vector code. The tile_sizes are parameters used by this sequence at different steps.
In this example this is telling the CPU backend to use

iree/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.cpp

Line 248 in c6616e0

void addDoubleTilingPadExpertPassPipeline(OpPassManager &passManager) {

sequence of transformations. The tile_sizes list are

1. tile sizes used to distribute work to threads (one level and coarsest level of tiling)
2. tile sizes used to tile the parallel dimensions of the result of the first level tiling (i.e. second level of tiling)
3. tile sizes used to tile the reduction dimension of the previous tiling. This is also second level of tiling, but is just tiling the K dimension instead of M and N dimensions that were tiled earlier. The split here helps with fusion.

All this was put in before we had better mechanisms in MLIR to control transformations. The Transform dialect work upstream will eventually supersede this and make things much nicer! But there is a long path to get there. In the meantime this has been useful.