Skip to content

fix(ascend): insert castptr for mismatched MTE pointers - #14

Open
Zhendong404 wants to merge 10 commits into
pto-devfrom
codex/unify-asc-mte-pointer-lowering
Open

Zhendong404 wants to merge 10 commits into
pto-devfrom
codex/unify-asc-mte-pointer-lowering

Conversation

@Zhendong404

Copy link
Copy Markdown

Summary

  • Normalize GM/UB MTE copy pointer element types in TileLang PTO codegen.
  • Emit byte pointers when source and destination element widths differ, while keeping padded GM->UB copies element-typed.
  • Avoid redundant same-type pto.castptr calls.

Validation

  • cmake --build build -j 8
  • pytest -q testing/ascend/language/test_tilelang_ascend_pto_mte_copy.py (3 passed)
  • Full tilelang.compile(..., target="ascend") for the aliased u16 UB -> u8 GM case
  • Existing copy-related tests: 135 passed, 2 skipped; 8 execution tests could not initialize NPU in this environment

jimmychou0 and others added 9 commits October 20, 2026 16:42
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).
@Zhendong404
Zhendong404 force-pushed the codex/unify-asc-mte-pointer-lowering branch 2 times, most recently from 6361f7e to 79cedf3 Compare October 22, 2026 08:17
Keep PTO MTE verifier contracts strict. TileLang now compares actual GM/UB handle types and inserts non-trivial castptr only when needed, avoiding redundant same-type casts and preserving padded GM->UB element semantics.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants