[Pass] SunmmioPipelinePlanning & InjectSunmmioPipeline#98
Open
wanghz18 wants to merge 32 commits intoSUNMMIO:tilelang_mesh_mainfrom
Open
[Pass] SunmmioPipelinePlanning & InjectSunmmioPipeline#98wanghz18 wants to merge 32 commits intoSUNMMIO:tilelang_mesh_mainfrom
wanghz18 wants to merge 32 commits intoSUNMMIO:tilelang_mesh_mainfrom
Conversation
… into pipeline_rebase
… into pipeline_rebase
…nto pipeline_rebase
… into pipeline_rebase
…nto pipeline_rebase
1 task
… into pipeline_rebase
…nto pipeline_rebase
… into pipeline_rebase
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
SunmmioPipelinePlanning
Step 1: 为语句打上生产者、消费者或者both的label。
使用Tilelang的原生识别模式。对于非复合语句,如果语句中shared buffer的写入是由global完成的,那么该语句是生产者,否则为消费者。生产者的模式只会出现在
dma_copy和BufferStore中。对于IfStmt或者SeqStmt这样的复合语句,如果子树上的label不同,那么该语句的label为both。Step 2: 根据Step 1中生成的label寻找需要multi-version的buffer。
Step 2.1: 使用Tilelang的原生识别模式。1、如果一个buffer被生产者写且被消费者读,那么该buffer可以multi-version。2、如果一个buffer的first write在last read之前,且该buffer的first write语句只读了global buffer,那么该buffer可以multi-version。
Step 2.2: 我们发现Tilelang原生的识别模式过于保守,因此加入了改进策略。如果一个buffer的所有写都是由可multi-version的buffer或者global进行的,那么该buffer也可以multi-version。根据这一策略可以拓展multi-version buffer的集合,直至其不再变化。
Step 3: 寻找Prologue语句。第0个iteration中的生产者语句即为Prologue。
Step 4: 寻找Body语句。设当前为原循环中的第k次循环,展开num_stages次,那么当前unroll包含第k次的非producer语句,第k+1到第k+num_stages-1次循环中的全部语句,第k+num_stages次循环中的producer语句。
Step 5: 寻找Epilogue语句。根据原循环次数与unroll次数的模关系,将尚未分配的语句加入Epilogue。
Step 6: 根据每个阶段对应的语句序列建立Dependency Graph。由于这里显式做了unroll,因此这是一个有向无环图。语句为节点,依赖关系为有向边。

Step 7: 根据cost model(暂时缺失)为DAG中的节点预测delay,也就是weight。根据DAG计算每个节点的bottom level,这个值可以看作是优先级,值越高在DAG中越关键,越需要优先完成。上图中附带了计算出的优先级。
Step 8: 根据上述计算出的优先级,计算出语句在流水线中的位置。如果一个device(ODMA、TensorCore、VectorCore)闲置,则运行当前可以运行的优先级最高的语句。


我们在实验中注意到,这样可能会导致额外的空泡出现,如下图所示。因为我们在这里贪心地要求设备空闲后即运行语句,但是这样并不一定是最优的。我们需要运行优先级较高的语句,但是当优先级高的语句都无法运行时就会运行一些不着急的语句,比如下一个iteration的数据预取,其实这些什么时候做都是可以的,不需要现在做抢占设备。
为此,我们做了另一个启发式的改进。下一个iteration的数据预取并不参加流水线排布,而是等待当前流水线排好后,再根据依赖关系插入流水线的空洞中。这样的结果如下:
Step 9: 将必要的信息记录在annotation中,传入下一步。本Pass结束。
InjectSunmmioPipeline
Step 1: 将需要multi-version的buffer在全局中多版本化,多版本后新增维度的index全部设置为0。
Step 2: 展开Prologue。多版本维度的index全部为0。
Step 3: 展开Body。原有的循环变量被替换为 unroll次数 * for_node->loop_var + 当前语句的iteration + for_node->min。多版本维度的index为 当前语句的iteration % unroll次数。这里“当前语句的iteration”可以参考上面流水线的排布图,例如“1-2”指的是第一个iteration的第二条语句,因此为1。
Step 4: 展开Epilogue。原有的循环变量被替换为 新的extent * unroll次数+ 当前语句的iteration + for_node->min。多版本维度的index为当前语句的iteration。