docs: add v0.5 release spec bundles#325
Open
KurrinQu wants to merge 209 commits into
Open
Conversation
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) <noreply@anthropic.com>
- 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 <noreply@anthropic.com>
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.
Co-authored-by: mouliangyu <mouliangyu@huawei.com>
(cherry picked from commit 34b7864)
(cherry picked from commit ecc78bf)
Signed-off-by: FangRui <fangrui_95@163.com>
…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)
mouliangyu
force-pushed
the
feature-vpto-backend
branch
from
5e223fb to
42b74f9
Compare
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
17 participants
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.
No description provided.