Repository navigation
[Ascend] Port T.Parallel ST suite (SV/CF/SP) onto pto-dev - #2
Open
peanutchan wants to merge 63 commits into
Open
peanutchan wants to merge 63 commits into
peanutchan wants to merge 63 commits into
Conversation
* Add PTO PTODSL codegen backend support Co-authored-by: LLMZhangYC <zhangyuchen39@huawei.com> Co-authored-by: benyang0506 <yb15295528526@163.com> * Move PTO codegen under Ascend backend --------- Co-authored-by: benyang0506 <yb15295528526@163.com> (cherry picked from commit 73bf5d5)
tile-ai#181) * Support PTO GEMM codegen * Align PTO GEMM pipeline scope * Refine PTO GEMM codegen checks * Add PTO GEMM pytest coverage * Align PTO GEMM template pipeline control * Align PTO GEMM disabled unit flag pipeline * Clean up PTO GEMM adaptation and add validation checks * Guard PTO local var buffer lowering * Clean up PTO GEMM examples and formatting * Fix PTO GEMM test coverage and vector loop lowering (cherry picked from commit c4748e2)
* Add PTO RMSNorm SIMT_VF codegen support Lower Ascend RMSNorm SIMT_VF blocks to PTODSL helpers, add PTO scalar/vector buffer access handling, pass dynamic UB size through the PTO host wrapper, and cover AscendC/PTO RMSNorm correctness in one test. * ci: isolate PTO example tests --------- Co-authored-by: LLMZhangYC <zhangyuchen39@huawei.com> (cherry picked from commit 2918b4a)
… PTO backend (tile-ai#211) * feat(pto): codegen for loading / storing scalars from / to the global scope * test(pto): add example tests for PTO the target * chore: revise code format * test(pto): Adapt test forms & clean redundant files * ci: update and rebuild PTOAS from latest main * Revert "ci: update and rebuild PTOAS from latest main" This reverts commit 3cb1fccf1a4f2c42d585446169094ade9b34467e. (cherry picked from commit 790110c)
…py tests (tile-ai#223) * Support FP8 inputs for PTO GEMM * Fix PTO GM-to-L1 NZ destination stride * Add PTO coverage for GEMM tests * Format PTO codegen (cherry picked from commit f2233a9)
* vmi: add T.vmi language surface and tl.vmi.* builtins * vmi: emit pto.vmi.* calls in PTO codegen * vmi: reject tl.vmi.* ops in Ascend CCE codegen with clear error * vmi: add tests and examples * fix(vmi): align create_mask/vcadd/... surface with ptodsl * fix(pto): add l2 cache ctrl argument for mte_ub_gm codegen * feat(vmi): add T.vmi.alloc_var() interface * vmi: add mhc/quant vmi examples * vmi: add T.vmi.alloc_local interface * test(vmi): update pto_vmi.py/pre_apply_mix_vmi.py * fix lint fails * fix(vmi): treat Tilelang int8 as pto.si8 for T.vmi.vcvt * fix(vmi): extract dtype from pointer annotations * feat(vmi): remove the offset argument for vmi.vload/vmi.vstore when using BufferLoad/BufferStore * test(vmi): register pto marker * test(vmi): use cli options instead of env variables * fix(pto): remove max_regs argument in the pto.simt decorator * ci: bump ptoas version to vmi 0.1.3 * test(vmi): remove vmi examples * fix: remove env options and torch.npu.set_device(0) * fix(vmi): mark test_vmi_tail_convert as pto test and remove others --------- Co-authored-by: Zhendong404 <zhendong404@gmail.com> (cherry picked from commit 9cd295b)
…O backend (tile-ai#272) * feat(pto): enhance ability of handling dynamic / symbolic parameters * test(pto): add test case test_int64_stride_vectorize_load[pto] (cherry picked from commit ff1a03a)
* feat(pto): add codegen for ascend_get_buf & ascend_rls_buf * test(pto): add test case test_simtvf_vecadd_mutex[pto] * test(pto): add test case test_simtvf_vecadd_auto[pto] test_simtvf_vector_add[pto] (cherry picked from commit a5748df)
(cherry picked from commit ae6d6db)
* feat(pto): support BF16 GEMM output * feat(pto): support FP32 GEMM inputs * feat(pto): bind HF32 mode per GEMM * fix(pto): preserve HF32 mode after positive loops * test(pto): streamline HF32 mode coverage --------- Co-authored-by: caojian5 <caojian5@huawei.com> (cherry picked from commit cb7b036)
…ch (tile-ai#305) * ascend(vmi): allow width-changing vinterpret_cast when bit totals match * test(ascend): cover width-changing vmi.vinterpret_cast bit totals (cherry picked from commit b7873e5)
(cherry picked from commit 4e1a2ed)
* feat(ascend): support FP8 SIMTVF PTO codegen Add FP8 SIMTVF dtype, cast, and local contiguous memory codegen for the PTO backend. Emit PTO floating-point min/max operations and add per-token FP8 cast test coverage. * feat(ascend): support compress and MHC kernels on PTO * feat(ascend): support postproc callbacks for PTO (cherry picked from commit c52f453)
* feat(pto): support multi-kernel JIT modules * fix(pto): preserve host kernel launch order (cherry picked from commit b8796c6775fe5b034084cee3c4f9a3c606c9d762)
* Add PTO SIMD VF support * Share Ascend and PTO SIMD VF tests * Fix shared VINTLV kernel scope * Fix PTO SIMD VF review comments * fix(pto): unroll topk_gate selection loop for VMI carry After asc tile-ai#336, serial loops inside PTO SimdVF lower to device-side pto.for_, and VMI register list updates do not carry across iterations. Explicitly unroll the k-selection loop so the PTO twin keeps working. * style: apply ruff format to PTO test files * refactor: restore VMI language module as vmi.py * test: use tilelang.testing.main for parameterized SimdVF/PTO entries * fix: address SimdVF/PTO twin review nits Unify backend= factory kwargs, preserve si* on unpack vload annotations, align vintlv N to covered lanes, clamp vselr indices, and compare on CPU. * test: import tilelang.testing for standalone __main__ entries * refactor: restore VmiPair and VMI scope helper names * fix(vmi): reject vinterpret_cast results with illegal PTO lane counts * style: apply ruff format to vmi lane-count helper * refactor: share SimdVF kernels via is_pto and speed up PTO CI Merge bitwise/vintlv ASC+PTO twins into one kernel when only the SimdVF body differs, and raise the PTO pytest xdist workers to 32 after rebase onto native Ascend targets. * refactor: share ASC non-VF shells across PTO SimdVF twins Lift PTO vecadd/vabsdif onto the ASC multi-core tile schedule and merge remaining twins into one factory with VF-only is_pto branches, keeping documented exceptions for scalar carry and topk pad/fill. * style: branch SimdVF twins on backend==pto directly Drop the is_pto locals and restore Ascend wording in shared path validation messages for both AscendC and PTO backends. * test: use shared DEFAULT_N=2**30 for ASC and PTO vecadd Keep AscendC and PTO on the same problem size so latency and bandwidth numbers are comparable. * ci: split PTO pytest into examples -n 4 and unit -n 32 * test: keep VMI AscendC reject check past SimdVF boolx256 gate (cherry picked from commit 449562df887709ec0416864bae2e0bc7f3042b0b)
(cherry picked from commit cabe6a09c28b17646fc72b1153739224333ac74d)
* feat(ascend): add packed FP4 VMI vcvt support * feat(ascend): add packed FP4 VMI support * fix(ascend): pack multidimensional FP4 VMI offsets * fix(ascend): validate packed FP4 VMI access modes * test(ascend): cover physical FP4 mask groups * fix(pto): scope packed host ABI to PTO (cherry picked from commit cb77d5a4397dd525cacab2ae0bcbd180b2ed65c8)
* fix(pto): carry vector local.var across T.serial Emit Python range() for T.serial loops that store an outer vector local.var so PTODSL ast_rewrite builds scf.for iter_args. Pin wheel.py-api to cp39 so Cython 3.3 can compile this revision. * style(pto): clang-format loop-carried local var helper (cherry picked from commit 102beab08d00018ccd5ce663e8b752d71118ba91)
* feat(ascend): migrate mixed GEMM PTO codegen * fix(ascend): rebind local vars in PTO sections * fix(ascend): correct mixed GEMM PTO codegen * feat(pto): remove ssa value workaround * fix(pto): add back kernel_kind for pto.jit. default is vector, section set None * feat(pto): GEMM support dual_copy(mixed=true) * feat(ascend): add PTO blockscaled GEMM codegen * feat(ascend): lower L0 blockscaled GEMM to PTO * test(ascend): harden PTO blockscaled GEMM coverage * refactor(ascend): centralize blockscaled MAD flag handling * feat(ascend): add FP4 blockscaled GEMM PTO support * feat(ascend): preserve explicit PTO L1-to-L0 controls * fix(pto): lower explicit T.set_atomic in PTO codegen Wire VisitExpr_ for tl.ascend_set_atomic / set_atomic_none on target=pto so kernels can arm/disarm store-atomic state without annotation-based T.copy(..., atomic_op=...). AscendC path already existed; no store-atomic copy lowering changes. * fix(example_gemm): include initial c0 in atomic-accumulate reference For acc=True the kernel computes C = c0 + A@B^T by accumulating into the pre-initialized output via store-atomic. The checking code compared against ref_program(x, w) which omitted c0, so max_diff equaled the magnitude of the random c0 (~5-6) and the case falsely FAILed. Pass c=c0 to the reference. It is a no-op for acc=False (c0 is None). * feat(ascend): GEMM L0 PTO codegen support * fix(ascend): allow float16 in PTO GEMM * feat(example): add target option for annotate unlimit memory * feat(ascend): support explicit FP4 PTO L1-to-L0 loads * style(ascend): format PTO mixed-kernel changes * test(ascend): unify Ascend and PTO GEMM coverage Add explicit target selection to GEMM examples and performance helpers. Parameterize auto, manual mixed, and blockscaled GEMM tests across the Ascend and PTO backends, including mixed, HF32, atomic, FP8, and FP4 paths. * fix(pto): address mixed GEMM codegen review findings * fix(pto): validate mixed section topology * fix(pto): reject unpaired physical sections * fix(pto): validate mixed kernel AIV count --------- Co-authored-by: jimmychou <47636600+jimmychou0@users.noreply.github.com> Co-authored-by: caojian5 <caojian5@huawei.com> Co-authored-by: andodo <andodo@Mac.local> Co-authored-by: andodo <andodo@andododeMacBook-Air.local> (cherry picked from commit a0ec76e11d68b8c18a345b6161642e9a78fdd291)
* fix(pto): inline Let expressions in Python/PTO codegen After tile-ai#382, shape >= 0 leaves nested Lets in Persistent For extents. CodeGenTileLangPY had no Let visitor, so PTO failed with tirx.Let. * fix(pto): emit Let as once-eval assignment above For headers Replace Substitute inlining with a single binding statement (CuTeDSL-style) and print For bounds before the loop header so nested Lets land on their own lines without duplicating side effects. * fix(pto): keep Let once-eval emit in PTO codegen only Move the Let visitor out of CodeGenTileLangPY so the shared Python base and CUDA/CuTeDSL paths stay unchanged. PTO still binds Let once above For headers after tile-ai#382. (cherry picked from commit 38cd68acda7380432f8bad80eb3dc8d64eb97651)
* feat(ascend): add PTO SIMTVF RNG codegen (cherry picked from commit 49b17b723fff0f2098a4652d35ce79d79a8fb6ac) * feat(ascend): support irregular GEMM PTO padding codegen (cherry picked from commit ef72cb54b94a4e83f86ed3f6f85db0fbb1aa1ae7) * feat(pto): improve mixed-kernel codegen for GEMM Split-K Add Nd2Nz lowering, safer cube/section kernel_kind selection, scalar min/max and carried local.var loop emission, plus NPU device-aware Cython launch so Split-K PTO kernels can target non-default devices. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 64ccd9a063e1246f3c34c8c5d6d9e9a36ac3bd2a) * Adapt example merge ub. Mod test args. Cleancode. (cherry picked from commit 3fe520e360a717bd64bf5694612500e058b6e435) * Fix PTO if_then_else literal select codegen fix(pto): directly lower while and loop break Add target option to GEMM various shapes/while pipeline example & test. Narrow shift codegen. (cherry picked from commit 87cf8922dae532d12f911d0d6ac10f1ceb7d2f27) * chore(ci): upgrade PTOAS to vmi-v0.1.5 * style(ascend): format PTO integration codegen * fix(pto): finalize PTOAS 0.1.5 integration * fix(pto): generate python-native for instead of pto.for_ * feat(pto): support dcache bypass and padded copies * fix(pto): harden PTOAS integration codegen --------- Co-authored-by: jimmychou <47636600+jimmychou0@users.noreply.github.com> Co-authored-by: csjlchen <chenjinlin1@huawei.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: andodo <andodo@Mac.local> Co-authored-by: andodo <andodo@andododeMacBook-Air.local> Co-authored-by: Zhendong404 <zhendong404@gmail.com> (cherry picked from commit 854806b75164c942bfff89a71f3979ea239565d2)
…e-ai#386) * feat(pto): generate pto.vecscope context manager for each SimdVF * ci: update PTOAS VMI wheel to v0.1.6 * test: mark empty SimdVF PTO regression (cherry picked from commit 229e194466fae98073cf6f65ecde2aafa79095e5)
* Enable PTO scalar GM DCache-bypass test. * Require actual PTO GM DCache-bypass call expressions in tests. * Format dcache-bypass analysis test for ruff. (cherry picked from commit 86bb3b46ad5319fafc9be2871507d06609282121)
* feat(pto): complete SIMD and VMI codegen support * fix(pto): preserve vdup scalar element types * fix(pto): handle cross-domain vdup conversions * test(pto): add vdup cross-type numeric check * style: format vdup PTO changes * test(pto): add int to float vdup numeric check (cherry picked from commit 069f85224cf1b41b56e16e4ad629e08a0385a121)
Adapt the PTO-specific portions of tile-ai#348 and tile-ai#356 to the refreshed target architecture.
Extract only the PTO-specific portions of tile-ai#420 instead of replaying the large upstream synchronization commit.
Adapt tile-ai#252 to the refreshed CI and pin the PTOAS vmi-v0.1.8 wheels used by the PTO backend.
After the refreshed target, tilelang.language only exposes the CUDA facade; Ascend and PTO tests must import tilelang.ascend.language.
Use the refreshed environment API for temporary files and hold TargetContext at direct lower() test call sites.
Move the cross-workitem all-reduce into tilelang/contrib/ptodsl/simt.py and stop monkey-patching the ptodsl-internal _allreduce module. The reducer/shuffle carrier wrappers (_redux_add/_redux_max/_redux_min/_shuffle_bfly) move with it: the low-level ops accept signless integer carriers only, so the authored-signedness adaptation at those boundaries belongs to the caller and now lives next to the dispatch tables that use it. simt.py keeps exporting simt_allreduce_sum/max/min (plus the existing scalar_div/scalar_rsqrt/vectorize_* helpers) so the generated kernels' 'from tilelang.contrib.ptodsl.simt import ...' contract is unchanged; no C++ rebuild is required. All scalar.* emissions are unified to pto.* and dtype descriptors to pto.float32/pto.si32 per the unified scalar/SIMT surface; redux_add is called without the retired signedness keyword. Co-authored-by: andodo <andodo@Mac.local>
* feat(ascend): accept precision='exact' on T.vmi.vdiv. SwiGLU residual search needs the CANN zero-ulp path. Hardware divide without the kwarg stays unchanged. * fix(pto): emit signless integer atomic_add. Current PTOAS rejects signedness= and expects i32/i64 pointers. Allow 64-bit integer atomics for SwiGLU clamp counters.
* fix(pto): pass signedness on integer atomic_max/min. PTOAS still requires signedness= for integer extremum ops. atomic_add stays signless. (cherry picked from commit 39272b12a74682484aacc2d2f9055dabbdc6a36c) * style: clang-format the atomic codegen region The pre-push clang-format hook (v23.1.0) reformats two lines in EmitAtomicOp that predate this branch. Keep it as a separate commit so the cherry-pick above stays faithful. --------- Co-authored-by: learning-chip <jiawei.zhuang@outlook.com>
* feat(ascend): expose tl.vmi.vunzip / tl.vmi.vzip
Register the lane-preserving element split/merge VMI ops and wrap them as
T.vmi.vunzip / T.vmi.vzip: the lane count is preserved while every element is
split into its low and high half-width halves (NxT -> 2 x Nx(T/2)); vzip is
the inverse.
Both wrappers take an optional to_dtype naming the produced element type
(defaulting to the unsigned half/double-width integer) and forward it to
PTODSL as a to_dtype annotation, so the requested type reaches the emitted
pto.vmi.vunzip / pto.vmi.vzip op instead of being re-derived there. The
parameter matches vcvt / vinterpret_cast, which already use to_dtype.
Also register the interleaved-channel counterparts channel_split /
channel_merge, which the same VMI surface exposes.
* fix(ascend): keep tl.vmi to vunzip/vzip and mirror PTODSL width checks
- drop channel_split/channel_merge: builtin registrations, declarations and
the two language wrappers; this change only needs vunzip/vzip
- vunzip requires a 16- or 32-bit source element (8 -> 4 bit pairs have no VMI
element type); vzip requires 8- or 16-bit halves plus matching lane count and
element type, mirroring PTODSL's checks and messages
- document that to_dtype is forwarded to the emitted PTO call, where PTODSL
derives the result element type from it and re-checks it against the source
width
- add a unit test covering the default unsigned halves, an explicit half
element type and the rejected width/operand combinations
The new test was not executed locally: it needs a rebuilt Ascend tilelang.
* test(ascend): stub _wrap_pair in the vunzip/vzip width test
test_pto_vunzip_vzip_mirror_ptodsl_width_rules mocked require_vmi_scope and
_call_vmi but still ran the real _wrap_pair(), which rejects calls outside an
active IRBuilder ("VMI pair results require an active IRBuilder"). The pair
wrapper is guarded independently of the mocked require_vmi_scope, so the test
could not get past its first vunzip call; mocking the scope check alone is not
enough.
Stub _wrap_pair the way test_pto_vload_dintlv_returns_pair does, and assert the
returned pair type/dtype (plus the vzip result dtype) so the happy path is
still covered rather than only the recorded _call_vmi arguments.
Still not executed here: this checkout has no runnable Ascend tilelang
("import tvm" fails and 3rdparty/tvm carries no tirx module), so the check is
py_compile plus reading _wrap_pair / IRBuilder.is_in_scope.
* feat(vmi): support ui8->ui16 vmi.vgather * fix(vmi): enforce uint16 offsets for 8-bit vgather
* Allow reverse packed-FP4 vcvt so cast-back can convert in-register. Stock pto-dev rejected float4_e2m1fn sources. Emit a scratch-free vinterpret_cast of the bf16x2 pair so packed e2m1 can widen to BF16/FP32. * fix(pto): load contrib helpers without importing tilelang. The PTO compile child ran tilelang.contrib imports from kernel.ptodsl.py, which loaded torch_npu and spun at 100% CPU while the parent held the NPU. * fix(pto): validate reverse FP4 vcvt on logical lanes and reject pmode.
pto-dev's PTO codegen emits 9- or 16-arg ascend_load_cbuf_to_ca/cb calls (16 args carry the MX scale-factor companion, see EmitAscendLoadCbufToL0). Relax set_num_inputs to -1 so IR construction accepts both forms, while keeping the standalone ascend_load_ca_sf/cb_sf builtins from PR tile-ai#3310 intact.
main moved codegen_py.cc/.h to src/backend/common/codegen/ and compiles them unconditionally, while pto-dev's src/ascend/CMakeLists.txt still appended the old src/cuda/codegen copy that fa11da8-era trees had. Drop the stale append and retarget the codegen_pto.h include at the relocated header (contents are identical).
main's BinaryCache.load/save take an explicit compile_format argument, while pto-dev's AscendBinaryCache overrides relied on a defaulted parameter that no longer exists. Pass binary_format explicitly at the backend.py call sites.
The ported tests (test_simtvf_*, test_blockscaled_gemm, test_int64_stride_vectorize_load) import example modules that already existed on pto-dev before fa11da8 and never landed on main, leaving pytest collection with 7 import errors. Restore the seven modules verbatim from pto-dev.
pto-dev's Ascend dialect exposes blockscaled_gemm (L0-input MXFP8 GEMM with optional trailing scale regions), while main's PR tile-ai#3310 baseline ships the same capability under gemm_blockscaled. The ported examples call T.blockscaled_gemm, which was missing. Add the pto-dev entry verbatim alongside gemm_blockscaled so both surfaces stay available.
pto-dev's block-scaled L0 pipeline loads MX scale factors through T.copy(x_l1, x_l0, scale=sf_region), which emits a three-region tl.tileop.ascend_copy that the PTO codegen lowers to asc_copy_l12l0a_mx/asc_copy_l12l0b_mx alongside the data load. The kwarg was missing from the ported surface, breaking the blockscaled_gemm L0 examples. Restore it verbatim from pto-dev.
Tree-copy examples/ascend/st_simtvf_parallel from Wenbo tip 3828bbcf onto PTO-ISA pto-dev (3d70ede). Adapt imports to tilelang.ascend.language and harness SimtVF FFI; document CANN MicroAPI opsim blocker on pto-b10.
Co-authored-by: peanutchan <peanutchan@users.noreply.github.com>
peanutchan
changed the title
SP4d walks every expert row and zeros pad holes with a full-lane vsel. SP5d sideband gathers V and Sf; interleave gathers one Pack slot then splits the scale tail. Oneshot, README, and SUMMARY now list 21+21. VF numbers for the new twins stay TODO. Co-authored-by: peanutchan <peanutchan@users.noreply.github.com>
SP4d pad25 VF=265 (1.98us); SP5d sideband VF=583 (1.87us), interleave VF=699 (1.94us). All PASS maxabs=0 on tip fa8ffee.
KurrinQu
reviewed
Oct 9, 2026
| out_ub = T.alloc_shared((E,), "float32") | ||
| T.copy(A, s) | ||
| T.copy(Thresh, thr_ub) | ||
| with T.SimtVF(threads=threads): |
There was a problem hiding this comment.
目前kernels是SimtVF形态，我们最终验证的路径是SimdVF→PTO，入库前是否需要调整成SimdVF？
Author
There was a problem hiding this comment.
同意最终验证路径是 SimdVF→PTO，但入库前不打算把整套 T.Parallel ST 直接改成只留 SimdVF。
计划是：同一套 T.Parallel case 通过环境变量切换 T.SimtVF / T.SimdVF（例如 ST_VF_MODE=simt|simd），再跑 regress 汇总各 case PASS/FAIL（哪些 SimdVF 已通、哪些还挂）。现有 kernels_ptodsl/*d（手写 PTO-DSL）仍保留，作为 sensitivity 对照，不是替代这条开关路径。
正在往本 PR 分支落地。
|
|
||
| ```bash | ||
| # Env (typical pto-b10 paths) | ||
| export ASCEND_HOME_PATH=/mnt/fluxdata/Ascend/cann_91b3/cann-9.1.0-beta.3 |
Author
There was a problem hiding this comment.
已参数化：oneshot / harness 通过 PY、ASCEND_HOME_PATH、TILELANG_DEPS 等环境变量取值，未设置则直接报错退出（见 tip 1006caa）。
README 里仍有一段 pto-b10 示例绝对路径；按你的意见会再清掉，只保留占位符说明，不再写死 /mnt/fluxdata/Ascend/...。
Oneshots and harnesses take PY, ASCEND_HOME_PATH, and the existing dependency prefixes from the environment so a reviewer can run the suite without editing pto-b10 lab paths. Co-authored-by: Cursor Agent <cursoragent@cursor.com> Co-authored-by: peanutchan <peanutchan@users.noreply.github.com>
Keep one T.Parallel source for the Ascend ST kernels and choose the VF frame from ST_VF_MODE. Add a simt/simd PASS/FAIL regress and drop checked-in lab absolute paths from the suite README. Co-authored-by: Cursor Agent <cursoragent@cursor.com> Co-authored-by: peanutchan <peanutchan@users.noreply.github.com>
ST_VF_MODE now chooses SimtVF or SimdVF and compile target ascend or pto. simt and simd stay as aliases for the ascend modes. The VF-mode regress runs all four by default and records PASS/FAIL, including a PTO compile failure, without remapping that compile to ascend. Co-authored-by: Cursor Agent <cursoragent@cursor.com> Co-authored-by: peanutchan <peanutchan@users.noreply.github.com>
Document tested versions and live VF-mode rollup (simt-asc 56/61, simt-pto 33/61, simd-pto 0/61; keep alloc_fragment). README + SUMMARY only; no kernel changes.
jimmychou0
referenced
this pull request
in jimmychou0/tilelang
Oct 12, 2026
Per review: the example came in only as an import dependency of test_simdvf_vecadd_pto_run.py from the ported vectorize commit. Element-wise device coverage stays with the self-contained test_vectorize_parallel_pto_run.py (vecadd 1D/2D, tail masks, casts, auto-schedule kept enabled); the element-wise example from PR #2 (t_parallel_st_suite sv1) becomes the canonical showcase once that suite lands.
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.
Summary
Ascend SimtVF sensitivity suite at
examples/ascend/t_parallel_st_suiteon PTO-ISA/tilelangpto-dev.3828bbcfpto-dev@3d70ede9Reviewer links
Primary metric in SUMMARY/deck: total VF cycles (Simt
T.Parallelvs best SIMD (PTO-DSL) schedule).Live VF-mode regress (cann_92b2 + PTOAS nightly 0.67, 2026-10-10)
Four-way
ST_VF_MODEswitch is on this tip (simt-asc/simt-pto/simd-asc/simd-pto). Live arms:/mnt/fluxdata/Ascend/cann_92b2/cann-9.2.0-beta.2(Ascend-cann-toolkit 9.2.0-beta.2)314834389d0e0472629702827ddf4d0b1aa5f740onfeat/st-simtvf-parallel-sv1-sv9(basepto-dev@3d70ede9)ptoas0.67 in.camodel_deps_nightly067— not CANN-bundled PTOAS; not the older camodelptoas-vmi 0.1.8(that stack was 0/61 on simt-pto via missingpto.max)ST_VF_MODE/mnt/fluxdata/happybot/st_out_cann92b2_20261010_104824. MicroAPI wall cleared vs cann_91b3. Residuals:sv4 remat_idx(eval_vec_), sv6/sv7 harness alias rows. VF vs Wenbo:/mnt/fluxdata/happybot/WENBO_VS_CANN92B2_SIMT_ASC_VF_20261010.md/mnt/fluxdata/happybot/st_out_nightly067_simt_pto_20261010_170837· tableVF_MODE_PASS_TABLE.txt. Was 0/61 on ptoas-vmi 0.1.8 (pto.max). Remaining FAIL compile: f16→f32EmitInlineSimtVF(~11), sv4 keep_idxllvm.store(~8), SV9 store arity (~3),eval_vec_/ SP5 / SP3 e8m0. VF vs Wenbo + simt-asc:/mnt/fluxdata/happybot/NIGHTLY067_SIMT_PTO_VS_BASELINES_20261010.md/mnt/fluxdata/happybot/st_out_simd_pto_20261010_163238(and SV1 recheck under nightly067). All FAIL compile. Dominant:Unsupported storage scope in PTO pointer emission: localfrom STT.alloc_fragment→ RF/localunder SimdVF→PTO (codegen_pto.cc). Decision: keep fragment RF as-is (do not rewrite kernels to shared-only); simd-pto needs a TileLang codegen fix forlocal/ SimdVF→PTO. Nightly 0.67 does not unblock SV1 here.kernels_ptodsl/(hand-written SIMD PTO-DSL) unchanged — SUMMARY SIMD VF numbers still come from that path.Test case PASS / FAIL (opsim Ascend950PR_9599)
Status key: PASS = compile + camodel gold OK on recorded tip.
Simt / most SIMD VF: Wenbo tip
3828bbcf(2026-10-07–08). Historical note: Simt oneshots oncann_91b3/92b1hit missing CANN MicroAPI; cleared on cann_92b2 forsimt-asc(see Live VF-mode section above).Exception (live on this tip): SP4d / SP5d SIMD PASS @
fa8ffeeb(2026-10-08).SV1–SV9
CF1–CF6
SP1–SP6
Rollup
Adaptations (examples only)
import tilelang.language as T→import tilelang.ascend.language as Tcommon_asc_harness.patch_simtvf_ffiuses stock Ascend 3-argSimtVF(no missingtilelang.language.vf)TILELANG_DEPSoverride for in-treelibtilelang.soNo broad
src// compiler patches. ST kernels keepT.alloc_fragment(RF/local) — shared-only rewrite is out of scope; simd-pto unblock is codegen.Test plan
simt-ptoon PTOAS nightly 0.67 — 33/61 PASS (was 0/61 on 0.1.8)oneshot_ptodsl_sp1d_sp6d.shrefresh on this tip.39gold=got when reachablesimd-ptoVF-mode regress — 0/61 PASS (localstorage scope / SimdVF→PTO; keepalloc_fragment)simd-ascVF-mode regress on cann_92b2