feat(vpto): prepare PTODSL pipe expansion - #986
Conversation
f77ab73 to
3b3b6c4
Compare
jimmychou0
changed the title
3b3b6c4 to
fc00205
Compare
fc00205 to
0b8266d
Compare
| case PIPE::PIPE_MTE1: | ||
| case PIPE::PIPE_MTE2: | ||
| case PIPE::PIPE_MTE3: | ||
| case PIPE::PIPE_V: |
There was a problem hiding this comment.
DSL侧允许set_intra_flag/wait_intra_flag设置PIPE_S，和这里不统一，看一下应该改哪里
There was a problem hiding this comment.
已修复 DSL 与 A5 verifier 的 endpoint 不一致：SyncSetOp 和 SyncWaitOp 现均接受 PIPE_S，与 PTODSL set_intra_flag / wait_intra_flag 的表层契约一致。
| 编译器侧路径仅由以下开关启用： | ||
|
|
||
| ```text | ||
| --enable-pipe-tilelib-expand |
There was a problem hiding this comment.
这个特性不需要加开关吧，应该是一个基础能力
There was a problem hiding this comment.
已删除flag, a5 的vpto 后端自动开启相关能力。
| `pto.tdrain` 是只由 PipeState 路径插入的内部 operation： | ||
|
|
||
| ```text | ||
| pto.tdrain(%pipe, %state : !pto.pipe, !pto.struct<i32, i32>) { split = <0|1> } |
There was a problem hiding this comment.
这个新增OP的语义、命名可能需要再讨论下。建议组织个会讲一次吧
| global pipe entry 可以直接使用 `pto.declare_global` 的结果，也可以使用从该声明 | ||
| 经过一个或多个 `pto.tassign` 得到的结果。传给 pipe op 的仍是重绑定后的结果值。 | ||
|
|
||
| ## 9. TileLib 交接 |
There was a problem hiding this comment.
似乎缺少了DSL侧为Pipe提供的操作抽象API定义，即lib开发者在他写的template中可以对pipe对象进行哪些操作？
There was a problem hiding this comment.
已补充第 4.3 节：定义了 PTODSL Pipe 的构造器、初始化、alloc/push/pop/free、属性及方向/split 约束。同时明确 TileLib template 不接收 Python Pipe 对象；它只通过 PipeSpec、资源、PipeState、entry 和 subblock ABI 实现各 transaction 的运行时语义。
| static std::unique_ptr<Pass> createSerialFrontendPipeLoweringPass() { | ||
| return std::make_unique<SerialFrontendPipeLoweringPass>(); | ||
| static std::unique_ptr<Pass> createRecursivePipeStateMaterializationPass() { | ||
| return std::make_unique<RecursivePipeStateMaterializationPass>(); |
There was a problem hiding this comment.
Materialize pipe state的位置要考虑下，放在哪里对现有pass的影响最小
There was a problem hiding this comment.
已将 MaterializePipeState 后移到 PTOResolveReservedBuffers 后、pipe-only InsertTemplateAttributes 前。这样 Layout、Fusion、PlanMemory 和 reserved-buffer 相关既有 pass 仍消费原有 pipe IR；PipeState 只在 metadata/expansion 需要前生成，且仅插入到 A5 VPTO 的 pipe transaction 路径。
2 participants
Summary
tdrain, resolved pipe metadata,ExpandTileOpspecialization, and VPTO LLVM lowering.splitvalue; derive terminal-drain split from resolved pipe-levelnosplit.ptodsl/tilelib/**, TileLib templates, and TileLib-ST/runtime validation are intentionally excluded and owned by the TileLib follow-up.Fixes #966
Validation
ninja -C build-local-vpto ptoas PTOPythonModulesptodsl/tests/test_vector_cube_ops.py -v(38 passed)expand_tile_op_tilelang_tadds.ptogit diff --check