support vpto backend#699
Merged
zhangstevenunity merged 266 commits intoMay 29, 2026
Merged
Conversation
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
There was a problem hiding this comment.
Code Review
This pull request significantly expands the project's documentation by adding detailed ISA specifications, design documents for core operations like mad and acc_store, and a framework for TileLang ST validation. It also integrates the PTO-Gym submodule, updates build and Docker configurations for TileLang DSL support, and introduces a script for automated release document generation. The review feedback identifies several areas for improvement in the generation script, including more robust version string parsing, avoiding incorrect heading demotion within code blocks, replacing fragile string matching with structured markers, and ensuring anchor uniqueness to prevent broken links in merged documentation.
| notes[version] = "Release refresh" | ||
|
|
||
| def key_fn(item: str) -> tuple[int, ...]: | ||
| return tuple(int(part) for part in item.split(".")) |
There was a problem hiding this comment.
This list comprehension will raise a ValueError if any part of the version string is not a digit (e.g., "0.4-beta"). Since this script is used for release documentation generation, consider adding a check to ensure only numeric parts are converted to integers to avoid runtime crashes.
| new_level = min(6, len(hashes) + levels) | ||
| return f"{'#' * new_level} {heading}" | ||
|
|
||
| return re.sub(r"^(#{1,6})\s+(.*)$", replace, text, flags=re.MULTILINE) |
There was a problem hiding this comment.
The regex r"^(#{1,6})\s+(.*)$" used with re.MULTILINE matches any line starting with a hash followed by a space. This can incorrectly demote comments inside fenced code blocks (e.g., in Python or Bash snippets) if they are at the start of a line. It is recommended to use a more robust way to identify headings that are not part of code blocks, such as splitting the text by code block markers before applying the substitution.
Comment on lines
+175
to
+178
| if "For detailed semantics, C-style pseudocode, and CCE mappings" in line: | ||
| continue | ||
| if "CCE correspondence" in line or "builtin mapping" in line.lower(): | ||
| continue |
There was a problem hiding this comment.
This logic for stripping unwanted lines relies on exact string matches for prose, which is fragile. If the source document (vpto-spec.md) is updated with slightly different wording, these lines will no longer be stripped. Consider using more stable markers or a structured approach to identify content to be removed.
Comment on lines
+212
to
+213
| if anchor_suffix: | ||
| return f"]({anchor_suffix})" |
There was a problem hiding this comment.
When an intra-bundle link contains an anchor (e.g., [Foo](chapter.md#anchor)), it is rewritten to ](#anchor). This assumes that the anchor name is unique across all inlined chapters in the merged document. If multiple chapters use the same anchor names, these links may break or point to the wrong location. Consider prefixing internal anchors with the chapter name during the merge process to ensure uniqueness.
Codex Review该评论由 review 机器人自动更新。
Summary检查到 PR #699 存在 4 个高优先级问题： Findings
这里先在 1647-1652 行保留了解码后模块自带的
Linux workflow 这里把
|
mouliangyu
force-pushed
the
feature-vpto-backend-merge
branch
from
2a4ccae to
db2bf1f
Compare
Contributor
|
/run a3 |
|
已接收
页面会自动刷新，可以直接看当前阶段、排队情况和最近结果。 |
A3 板测失败
失败用例
|
A3 板测失败详情：PR #699syncall_binding
tprefetch_async_binding
|
Contributor
|
/run a5 |
|
已接收
页面会自动刷新，可以直接看当前阶段、排队情况和最近结果。 |
A5 板测失败
失败用例
|
A5 板测失败详情：PR #699tprefetch_async_binding
syncall_binding
test_intercore_sync_a5_dyn
matMul
cmps
cmp
|
Collaborator
|
/run a5 |
|
已接收
页面会自动刷新，可以直接看当前阶段、排队情况和最近结果。 |
A5 板测失败
失败用例
|
A5 板测失败详情：PR #699tprefetch_async_binding
syncall_binding
rope_kv_cache
plan_memory_loop_in_if
cmps
cmp
|
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.
Move lit-discovered .pto tests out of test/basic into test/lit/pto or test/lit/vpto, and update related documentation paths. Relocate misplaced VPTO sim cases under the micro-op hierarchy, remove a duplicated nested sim case, and add local guidance for choosing PTOAS test frameworks.
Update VPTO lit tests to inspect explicit VPTO IR with --emit-vpto where they check VPTO-level rewrites. Route *_vpto_llvm tests through the VPTO LLVM lowering dump and check llvm.hivm calls from LLVM dialect MLIR instead of stale VPTO output.
Relocate FileCheck-based VPTO .pto tests from test/vpto to test/lit/vpto so they are discovered by the lit framework. Wrap auto-vecscope tests in a vector kernel submodule to match the current VPTO container form.
* add tcolargmax/min tileops lib implementation * fix tcolargmax/min implementations and tile_buf declaration * add init texp high precision implementation * add texp high precision implementation * remove exp high precision code in math.py(already in custom.py) * add empty lines in math.py * fix texp high precision test case's eps * fix texp test case * move exp hp from custom to exp_hp.py
* Rename VPTO MTE semantic ops Input: existing VPTO memory-transfer semantic ops and tests using dma/cube/acc_store-style names. Output: public semantic ops use mte_<src>_<dst> names, with ptr addrspace aliases for the new memory-space spelling. Key steps: update ODS/parser/lowering/expand/ptr-normalize paths; migrate lit, VPTO, and TileLang test inputs; refresh VPTO/ISA docs and legacy raw-op wrapper mappings. --------- Co-authored-by: mouliangyu <mouliangyu@huawei.com>
* Add OP for TMrgSort * Add OP for TMrgSort | fix review * Add constraints to tmov template for UB2UB ND2ND only Add constraint function to restrict template_tmov_basic to only support: - UB2UB: Both src and dst must be in Unified Buffer (memory_space="ub") - ND2ND: Both tiles must have N-dimensional layout (s_layout=NONE_BOX) Other scenarios (GM2UB, UB2GM, specialized layouts) require different implementation paths and are explicitly rejected by this constraint. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> --------- Co-authored-by: caojian5 <caojian5@huawei.com> Co-authored-by: Claude Opus 4.7 <noreply@anthropic.com>
* Add HP (HIGH_PRECISION) support for TDiv * Add HP for divs recip rowexpanddiv colexpanddiv * [fix] Remove redundant else branch in HIGH_PRECISION mode since MLIR validation guarantees f16/f32 only * [fix] Add missing HIGH_PRECISION test kernels for tdivs ST test HP test cases were defined in cases.py but missing from implementation files, causing FileNotFoundError when compare.py tried to read output.bin. Added 12 HP kernels across three synced layers: - tdivs.pto: HIGH_PRECISION kernels with precision_mode attribute - main.cpp: kCases[] entries and launch wrapper declarations - launch.cpp: Launch functions with correct IEEE 754 scalar values Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> --------- Co-authored-by: caojian5 <caojian5@huawei.com> Co-authored-by: Claude Opus 4.7 <noreply@anthropic.com>
mouliangyu
force-pushed
the
feature-vpto-backend-merge
branch
from
db2bf1f to
21ad259
Compare
A5 板测成功
|
A3 板测失败
失败用例
|
A3 板测失败详情：PR #699syncall_binding
tprefetch_async_binding
|
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
20 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.
Summary
This PR adds the VPTO backend to PTOAS. The new backend is selected with
--pto-backend=vpto; the default backend remainsemitc.Main additions:
tools/ptoas, including VPTO module normalization, TileLang template expansion, VPTO pointer normalization, wrapper-op expansion, LLVM lowering, host stub emission, and fatobj emission.ptoas.Impact on non-VPTO code paths
The default
ptoaspath still uses theemitcbackend. VPTO split/normalize, TileLang expand, VPTO pointer cleanup, VPTO LLVM emission, and fatobj emission are only entered through the VPTO backend branch or explicitly registered passes.This PR does change shared PTO IR and shared passes used by non-VPTO flows:
!pto.ptrnow carries a PTO memory space. The parser/printer accepts forms such as!pto.ptr<T, gm>and!pto.ptr<T, ub>, with GM as the default.getPTOAddressSpaceAttr(Type)now reads the memory space fromPtrType, andPTOViewToMemrefpreserves that memory space when lowering ptr types to memrefs.set_flag,wait_flag,get_buf, andrls_bufnow use custom parsers/printers.get_buf/rls_bufalso acceptpipeattrs in addition to the previous sync op type attrs.castptr,get_tensor_view_stride, andmem_barare added as shared IR surface.precision_modeattr:tcolexpanddiv,tdiv,tdivs,texp,tlog,trecip,trowexpanddiv,trsqrt, andtsqrt. The verifier requiresHIGH_PRECISIONto be used only with f16/f32 element types.castptrvalidates supported cast forms and memory-space consistency; A5tload/tstorevalidate vec layouts; shift-like tile ops require src0/src1/dst element types to match; scalar tile ops such astadds/tmuls/tsubsrequire valid rows to match on A5;tpart*ops validate source valid shapes against destination valid shapes;tselsvalidates row-major layout for mask/src/dst on A5.PTOViewToMemrefpreserving ptr memory spaces, insertingpto.bind_tileanchors for original tile_buf function arguments, loweringget_tensor_view_stride, and preserving attrs when rebuilding selected tile ops.PTOPlanMemoryhandles null alias inputs defensively.PTOToEmitCmaps the extendedReluPreModevaluesScalarRelu,VectorRelu, andPwl.Validation
Local validation run on this branch:
ninja -C build-ci-local ptoas ptobc installninja -C build-ci-local check-pto(416/416)OK=263 FAIL=0 SKIP=16)docker/test_ptoas_cli.sh312tests)281/281)85/85)docker/collect_ptoas_dist.shwas checked locally and failed because the script only collects LLVM libraries from the CI LLVM workspace layout; the local LLVM installation is under the actions-runner tool cache.