Repository navigation
[Ascend] Port pto-dev PTO/VMI increments onto the PR #3310 baseline - #1
Open
jimmychou0 wants to merge 19 commits into
Open
jimmychou0 wants to merge 19 commits into
jimmychou0 wants to merge 19 commits into
Conversation
Add the VMI (virtual machine instruction) intrinsics and PTO codegen:
- tilelang/ascend/language/vmi.py: the T.vmi.* surface, exposed as T.vmi
through tilelang.ascend.language - vector loads/stores, masks,
arithmetic/bitwise/shift ops, vcvt (incl. packed FP4 sources),
vinterpret_cast, vzip/vunzip, precision-aware vdiv, gather/scatter,
histogram ops, and register allocation helpers
- 54 tl.vmi.* TIR builtins registered in src/op/builtin.{h,cc}
- PTO codegen: pto.vmi.* emission with typed-literal wrapping, FP4
pair expansion, eager vreg allocation for vector local.var, the
signed/unsigned integer type split (pto.si*/pto.ui*), and a clear
AscendC rejection diagnostic; generated kernels import
tilelang.contrib.ptodsl before ptodsl so this build's native
libraries load ahead of the PTODSL MLIR modules
- KernelParam.storage_shape: packed-storage shapes for sub-byte dtypes
Co-authored-by: KurrinQu <qukelin1991@163.com>
Co-authored-by: Jay Zhuang <80731350+learning-chip@users.noreply.github.com>
Co-authored-by: Zhang Zhendong <zhendong404@gmail.com>
Co-authored-by: cj <erhsh_165@126.com>
Add the T.Parallel -> PTO VMI vectorization pipeline: - Verify -> Legalize -> Vectorize pass pipeline ahead of AscendSimdVFLowerParallel, with pass registration, the Python interface, and routing tests (SIMT and hand-written VMI regions are bypassed inside the passes; control-flow rewriting stays out of scope) - Verify: shared region scan and inverse lane-mapping analysis, the LayoutInference-for-PTO branch, 1D/2D lane mapping, and T.SimdVF(lanes=...) frontend support, with contract tests - Vectorize: chunk loops with dynamic prefix tail masks, remaining state reset per row/column and per sequential unit, and dtype/cast codegen for element-wise add/sub/mul/div, unary abs/exp/log/sqrt, and f32 <-> f16/bf16 conversions; runtime lane-uniform float computations are rejected (they would stay scalar inside SIMD_VF) - Device tests: element-wise, tail-mask, and FP8 suites Co-authored-by: qukelin <qukelin1991@163.com> Co-authored-by: jimmychou <47636600+jimmychou0@users.noreply.github.com>
The si32 reducer strategy serves both signed si32 and signless i32 carriers, but created its initial-value constant as signless pto.i32. With TIR signedness now carried into PTODSL, scalar.select mixing an si32 carrier with an i32 constant is rejected by current PTOAS (pto.select integer arms require matching types). Both creation points (_emit_warp_hw_reduce and _emit_cross_warp_reduce) now share _identity_const: the strategy key still selects the constant's numeric value, while its integer type follows the carrier value's own signedness (signed carriers get pto.si32, signless keep pto.i32).
Signed-off-by: wangchen <wangchen965@huawei.com>
Signed-off-by: wangchen <wangchen965@huawei.com>
feat(pto): default disable LegalizePTOVMIMemoryAccess
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
Ports the pto-dev PTO/VMI increments from WenboCodes/tilelang-deepseek
pto-devonto the merged PR tile-ai#3310 baseline: 47 commits fromfa11da8f([PTO] Add PTO codegen backend support tile-ai#140) through34c6d3ea, plus 7 port-integration fixes (latest3d70ede9).What's included
T.vmi,vzip/vunzip,vgather, packed FP4 vcvt, width-changingvinterpret_castpto.vecscope, unroll-hint forwarding, PTO perf-regression CI jobPort-integration fixes (validated on hardware)
50bb9af9relaxascend_load_cbuf_to_ca/cbto variable arg count (9/16 MX companion)e6ae351cpoint PTO codegen at the relocatedcodegen_pysourcesd4d04fb7passcompile_formattoBinaryCache.load/save90054e32restore 7 pre-fa11da8f PTO example modules0774ecffrestoreT.blockscaled_gemm(delegates sfa/sfb to the 16-slottl.tileop.gemm_blockscaledprotocol)3d70ede9restore thescale=MX companion kwarg onT.copyValidation (238, x86_64 + Ascend950PR; pto-as vmi-v0.1.9
b9da12032; CANN 9.1 runtime + 9.2 OPP; py3.11/torch 2.10 — deviation from the aarch64 CI env noted)*Known issue A — pto-as
mad_rawlowering regression (31 cases)All
test_blockscaled_gemm_l0_*,test_gemm_padding_irregular_shapes(fp8 blockscaled) andtest_gemm_splitk/test_gemm_ub_mergefailures arepto.mad→pto.mad_raw"explicitly marked illegal" in pto-as v0.1.8..master (tested e61df455d, stack-v3 98a338585, stack7, b9da12032). The identical generated.ptolowers fine on the pre-0915 PTOAS build, and the.ptotext is byte-identical to what original pto-dev emits — no tilelang-side regression. Needs a pto-as fix or a matched release; tracked separately.*Known issue B — cross-file state contamination (~7 cases, non-deterministic)
test_vecadd[simd/simt],test_while_pipelined,test_int64_stride_vectorize_load,test_sync_inter_autoschedule,test_per_token_cast_to_fp8fail only when mixed with other files in one pytest process; each passes standalone and per-file (test_vecadd.pywhole file: 2 passed). Test-order-sensitive runtime state; recommend per-file isolation in CI while investigating.Notes
3rdparty/tvmsubmodule stays on the main pointer (4211874e), a descendant of pto-dev's907a88c8on the same fork.