docs: align VPTO MTE memory naming by WenboCodes · Pull Request #341 · mouliangyu/PTOAS

WenboCodes · 2026-05-08T09:58:40Z

Summary

reorganize MTE DMA copy docs around public pto.mte_<src>_<dst> op naming
move cube load/store movement docs into the DMA copy chapter and split GM-to-L1 burst/fractal forms
align VPTO spec references and address-space conventions with L1/L0A/L0B/L0C/BT names and aliases

旧 wrapper 接口	新 MTE 接口	数据方向	说明
`pto.dma_load`	`pto.mte_gm_ub`	GM -> UB	Structured GM-to-UB DMA load wrapper；`nburst(...)` / `loop(...)` / `pad(...)` 保留
`pto.dma_store`	`pto.mte_ub_gm`	UB -> GM	Structured UB-to-GM DMA store wrapper；`nburst(...)` / `loop(...)` 保留
`pto.dma_copy`，dst 是 UB	`pto.mte_ub_ub`	UB -> UB	UB 内部 copy wrapper
`pto.dma_copy`，dst 是 L1 / `mat`	`pto.mte_ub_l1`	UB -> L1	UB-to-L1 copy wrapper
`pto.cube_load`	`pto.mte_gm_l1_burst`	GM -> L1	Structured GM-to-L1 burst copy wrapper
`pto.cube_load_frac`	`pto.mte_gm_l1_fractal`	GM -> L1	Structured GM-to-L1 fractal load wrapper；用 `nd2nz` / `dn2nz` mode 区分
`pto.left_load`	`pto.mte_l1_l0a`	L1 -> L0A	Structured L1-to-L0A wrapper
`pto.right_load`	`pto.mte_l1_l0b`	L1 -> L0B	Structured L1-to-L0B wrapper
`pto.left_load_mx`	`pto.mte_l1_l0a_mx`	L1 -> L0A	MX-mode L1-to-L0A wrapper
`pto.right_load_mx`	`pto.mte_l1_l0b_mx`	L1 -> L0B	MX-mode L1-to-L0B wrapper
`pto.bias_load`	`pto.mte_l1_bt`	L1 -> BT	Structured L1-to-BT bias load wrapper
`pto.acc_store`，dst 是 L1 / `mat`	`pto.mte_l0c_l1`	L0C -> L1	Structured L0C-to-L1 wrapper
`pto.acc_store_gm`	`pto.mte_l0c_gm`	L0C -> GM	Structured L0C-to-GM writeback wrapper
`pto.acc_store_ub`	`pto.mte_l0c_ub`	L0C -> UB	Structured L0C-to-UB wrapper

Explain block/subblock runtime queries in workload-partitioning terms and remove redundant supported-forms wording from conversion ops docs. Co-Authored-By: Claude Opus 4.6 (1M context) <[email protected]>

- Add detailed mode parameter documentation (mode=0 vs mode=1) - Add 'Why get_buf/rls_buf is More Programmer-Friendly' section: - No manual priming/draining for ping/pong loops - No loop peeling for complex/nested loop dependencies - Simpler mental model (buffer ID + program order) - Add quick example comparison showing set_flag overhead vs get_buf simplicity - Update Example 2 and 3b with explicit mode=0 in code - Update comparison table with 'Loop peeling' row

- set_flag/wait_flag: 2 IDs per buffer (1 forward + 1 reverse pipe-pair) - get_buf/rls_buf: 1 ID per buffer (handles both directions automatically) - 8 per pipe-pair is HW limit, not a formula

- set_flag/wait_flag: 8 IDs per pipe-pair direction (HW limit) - get_buf/rls_buf: 1 buffer ID per shared resource (HW limit: 32 global), same ID used across all pipelines

- Event ID mgmt: each buffer occupies 1 ID per direction (removed misleading 4 IDs calc) - Drain example: use concrete EVT_*_0/EVT_*_1 instead of {(N-1)%2} expressions

- 4 set_flag + 4 wait_flag (not 8) - 4 IDs = 2 pipe-pair directions × 2 ping/pong buffers

- set_flag/wait_flag: 1 MTE2 load, 8 Vector slices — must peel set/wait outside loop - get_buf/rls_buf: same pattern but acquire/release can stay inside or outside

- Acquire/release per slice inside loop - Iteration 0 blocks until MTE2 done, iterations 1-7 proceed immediately

Add the merged v0.3 PTO micro-instruction release spec document for A5, including ISA group references and updated synchronization notes. Co-Authored-By: Claude Sonnet 4.6 <[email protected]>

Introduce a three-pass pipeline that lowers PTO tile ops to vector-level implementations via TileLang DSL templates: - ExpandTileOp: invokes TileLang Python DSL to instantiate template functions and replaces tile ops with func.call. SpecKey covers all operands; tile_buf operands are passed through without bridging. - PTOInlineLibCall: extended to recognize tilelang instance functions via the attribute set by the DSL frontend. - FoldTileBufIntrinsics: resolves pto.tile_buf_addr / tile_valid_rows / tile_valid_cols, including dynamic valid-shape via pto.bind_tile chain tracing. - MemrefToTileBuf: recovers tile_buf types from memref + bind_tile metadata after PlanMemory/InsertSync. - PTOViewToMemref: insert pto.bind_tile anchors for tile_buf function args so MemrefToTileBuf can recover them. Adds new PTO ops (tile_buf_addr/tile_valid_rows/tile_valid_cols), ptoas pipeline wiring, design docs, and unit tests.

Signed-off-by: FangRui <[email protected]>

…amming Model - Define fractal NZ layout (K1M1M0K0 / K1N1K0N0 / N1M1M0N0) for L1/L0A/L0B/L0C - Document full GM->L1->L0A/B->L0C->GM data flow pipeline with ASCII diagrams - Clarify copy_gm_to_cbuf_multi_nd2nz vs dn2nz (nd2nz preferred for GEMM; dn2nz for NCHW/conv; A2/A3 only has nd2nz so nd2nz is backward compatible) - Clarify L0A layout: FRACTAL_NZ K1M1M0K0 on A5 (FRACTAL_ZZ M1K1M0K0 on A3) - Clarify load_cbuf_to_ca/cb: each burst = one 512B fractal z-block (16x16 bf16); inner-box transpose for B done on-the-fly during MTE L1->L0B transfer - Add copy_matrix_cc_to_ub writeback path (A5 only, fixed-point datapath)

Add new subsection under Intra-Cluster Data Paths in Cluster Programming Model: - Define fractal NZ layout (K1M1M0K0 / K1N1K0N0 / N1M1M0N0) for L1/L0A/L0B/L0C - Per-buffer NZ layout table with copy ops - L0A: FRACTAL_NZ K1M1M0K0 on A5 / FRACTAL_ZZ M1K1M0K0 on A3 - Full GM->L1->L0A/B->L0C->GM ASCII pipeline diagram - load_cbuf_to_ca/cb: each burst = one 512B fractal z-block; B transpose on-the-fly - copy_matrix_cc_to_ub writeback (A5 only, fixed-point datapath) - nd2nz preferred for GEMM; dn2nz for NCHW/conv; A2/A3 has no dn2nz (backward compat)

* [Feat] enhance vcvt behavior for rowsum * [Feat] enhance vcvt behavior for rowsum * [Feat] enhance vcvt behavior for rowsum * update attributes

* [Add] tcmp trem tfmod * [Fix] license-header check --------- Co-authored-by: zwd060924 <[email protected]>

* fix tfillpad_inplace_template * fix tfillpad_inplace * fix tfillpad_inplace * fix tfillpad_inplace * fix tfillpad_inplace * fix tfillpad_inplace * fix tfillpad_inplace --------- Co-authored-by: kangjiaming <[email protected]>

…angyu#332) * Delete useless code and abbreviate tile_buf type in test cases * Delete useless code and abbreviate tile_buf type in test cases * Delete useless code and abbreviate tile_buf type in test cases * Delete useless code and abbreviate tile_buf type in test cases --------- Co-authored-by: kangjiaming <[email protected]>

* add tcolargmax/min tileops lib implementation * fix tcolargmax/min implementations and tile_buf declaration

* Add OP For TSort32 * review-fix: tile_buf use compact-mode * review-fix: tmp param should be optional --------- Co-authored-by: caojian5 <[email protected]>

mouliangyu and others added 30 commits April 28, 2026 04:36

feat: add vpto backend

97cbb4b

clarify block query docs and trim conversion section

cd451b4

Explain block/subblock runtime queries in workload-partitioning terms and remove redundant supported-forms wording from conversion ops docs. Co-Authored-By: Claude Opus 4.6 (1M context) <[email protected]>

fix: correct event ID explanation in comparison table

46f1a8a

- set_flag/wait_flag: 2 IDs per buffer (1 forward + 1 reverse pipe-pair) - get_buf/rls_buf: 1 ID per buffer (handles both directions automatically) - 8 per pipe-pair is HW limit, not a formula

fix: clarify event ID management in comparison table

3df56a7

- set_flag/wait_flag: 8 IDs per pipe-pair direction (HW limit) - get_buf/rls_buf: 1 buffer ID per shared resource (HW limit: 32 global), same ID used across all pipelines

fix: simplify event ID explanation and drain example

31540f2

- Event ID mgmt: each buffer occupies 1 ID per direction (removed misleading 4 IDs calc) - Drain example: use concrete EVT_*_0/EVT_*_1 instead of {(N-1)%2} expressions

fix: correct set_flag/wait_flag count in quick example

440925c

- 4 set_flag + 4 wait_flag (not 8) - 4 IDs = 2 pipe-pair directions × 2 ping/pong buffers

fix: add concrete 1:N example for loop peeling comparison

a7e5451

- set_flag/wait_flag: 1 MTE2 load, 8 Vector slices — must peel set/wait outside loop - get_buf/rls_buf: same pattern but acquire/release can stay inside or outside

fix: show get_buf/rls_buf inside scf loop for 1:N example

4d5dac8

- Acquire/release per slice inside loop - Iteration 0 blocks until MTE2 done, iterations 1-7 proceed immediately

refactor vpto llvm emiter

dcd8586

add online softmax (q * k is ready) case

227da17

docs: add VPTO spec v0.3 release draft

dbfe96b

Add the merged v0.3 PTO micro-instruction release spec document for A5, including ISA group references and updated synchronization notes. Co-Authored-By: Claude Sonnet 4.6 <[email protected]>

add PTO-Gym submodule

a77c8b2

feat: add PTO-Gym guide skill

949a928

Add tilelang dsl implementation

495a02c

Update openspec

f5ee676

Support more syntax in tilelang dsl

47a102f

Update openspec

5ac3dba

tilelang dsl first working version

0fdf4d3

Update openspec

0d8e71e

Support more dsl syntax

5d96301

Amend

0945be0

Update docs

bd86d21

docs: add TileOp expand design and demo

11bc085

brough back the lost pass and pto op after rebase

f640129

feat: tile lang dsl

40b2c93

Update DSL user guide

203ea91

Support more tile attributes

51112bb

Support more event IDs

30dd61a

Zhendong404 and others added 25 commits April 30, 2026 14:56

fix(vpto): force V300 ctrl mode for sat-sensitive vcvt

6670954

feat: add cube dma laod/store ops

0626c66

bugfix: fixup mix fatobj

33905f4

fix(test/vpto): use pto.acc_store_gm for ACC to GM in cube vpto ST

e8014e9

feat: add left/right_load_mx

f9b523b

Signed-off-by: FangRui <[email protected]>

[Feat] enhance vcvt behavior for rowsum (mouliangyu#295)

b1ebb9f

* [Feat] enhance vcvt behavior for rowsum * [Feat] enhance vcvt behavior for rowsum * [Feat] enhance vcvt behavior for rowsum * update attributes

[Add] tcmp tfmod trem (mouliangyu#321)

68a91a7

* [Add] tcmp trem tfmod * [Fix] license-header check --------- Co-authored-by: zwd060924 <[email protected]>

fix tfillpad_inplace_template (mouliangyu#299)

821189e

* fix tfillpad_inplace_template * fix tfillpad_inplace * fix tfillpad_inplace * fix tfillpad_inplace * fix tfillpad_inplace * fix tfillpad_inplace * fix tfillpad_inplace --------- Co-authored-by: kangjiaming <[email protected]>

Add trems/tcmps/tfmods (mouliangyu#283)

85dea6e

add tcolargmax/min tileops lib implementation (mouliangyu#294)

5d5fcf2

* add tcolargmax/min tileops lib implementation * fix tcolargmax/min implementations and tile_buf declaration

Add OP For TSort32 (mouliangyu#258)

8e7a0d5

* Add OP For TSort32 * review-fix: tile_buf use compact-mode * review-fix: tmp param should be optional --------- Co-authored-by: caojian5 <[email protected]>

Fix default render dtype selection for AnyInt

7e6ddc6

Support scalar/index auto conversion

5a51620

Add precision_mode attr plumbing for tile math ops

b3acd8e

Fix PrecisionMode Python bindings

14ccac5

Remove PrecisionMode binding smoke test

96a27e0

docs: reorganize MTE DMA op naming

f51507a

docs: rename MTE ops to hardware memory tokens

202ce37

docs: use L1 wording for MTE buffer docs

ef3ac3b

docs: split GM to L1 MTE forms

0a6c653

docs: align VPTO spec with MTE op docs

d828699

docs: document PTO address space aliases

9ef8b4e

WenboCodes requested review from mouliangyu and peanutchan May 9, 2026 01:14

mouliangyu reviewed May 9, 2026

View reviewed changes

Comment thread docs/isa/micro-isa/02-dma-copy.md Outdated

docs: shorten GM to L1 MTE op names

1251f85

mouliangyu force-pushed the feature-vpto-backend branch from 5e223fb to 42b74f9 Compare May 14, 2026 00:19

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

docs: align VPTO MTE memory naming#341

docs: align VPTO MTE memory naming#341
WenboCodes wants to merge 227 commits into
mouliangyu:feature-vpto-backendfrom
WenboCodes:codex/vpto-mte-docs

WenboCodes commented May 8, 2026 •

edited

Loading

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

17 participants

Conversation

WenboCodes commented May 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

17 participants

WenboCodes commented May 8, 2026 •

edited

Loading