Skip to content

third_party/pto-isa: Update reference to recent state - #663

Open
ChristosMatzoros wants to merge 1 commit into
sgl-project:mainfrom
ChristosMatzoros:update_pto_isa
Open

third_party/pto-isa: Update reference to recent state#663
ChristosMatzoros wants to merge 1 commit into
sgl-project:mainfrom
ChristosMatzoros:update_pto_isa

Conversation

@ChristosMatzoros

Copy link
Copy Markdown
Contributor

Update third_party/pto-isa: bump af66b6ad6081a2d3

One-line change: the third_party/pto-isa submodule reference.

The current pin is from 2026-05-06 — three months old — and predates the whole fused
multiply-add instruction family, including TMULADDDST, which
PR #601 needs for the A5 causal_conv1d
path.

from af66b6ada43c40ad9583731883fb9bef65ac840a (2026-05-06)
to 6081a2d3master tip (2026-08-06)
upstream https://gitcode.com/cann/pto-isa.git

Notable additions

New instructions, for both npu/a5 and npu/a2a3:

  • fused multiply-addTMulAddDst, TFusedMulAdd, TFusedMulAddRelu
  • fused activation/quant epiloguesTSubRelu, TSubReluConv, TAddDeqRelu, TAddReluConv
  • async prefetchTPrefetchAsync
  • reductions / layoutTPairReduceSum, TTernOp, and (a5) TInterleave, TDeInterleave

Plus dtype coverage extensions (TRowMax/TRowMin uint8/int8, TABS int8/16/32, TCmp bf16),
A5 TMOV DN→ZZ, fixpipe support in TPUSH, and precision fixes to the new fused ops.

TMULADDDST is the one PR #601 depends on: an element-wise fused multiply-accumulate,
dst = src0 * src1 + dst, lowering to a single instruction on both families (vmula on dav-c310,
vmla on dav-c220).

Validation

mega_chunk_gdn — 910B2 builds clean, tests pass, 0 errors
codegen stability the .so built at 6081a2d3 is byte-identical to one built at b963dcb6, 143 commits earlier — the intervening range does not touch the a2a3 codegen these kernels use

Signed-off-by: Christos Konstantinos Matzoros <christos.konstantinos.matzoros@h-partners.com>
EdwardXuy pushed a commit to EdwardXuy/sgl-kernel-npu that referenced this pull request Aug 10, 2026
1. Always install triton-ascend 3.2.1 (per dev advice) — fixes triton.language.extra.cann missing on CANN 8.5.0

2. sglang stub: uninstall partial sglang first, then create clean stub providing only is_npu (real sglang __init__ needs orjson etc.)

3. Merge PR sgl-project#663: bump pto-isa to 6081a2d3 (tests pass per developer)

4. F injection kept (fixes NameError in test_swiglu_quant; test now fails on precision assertion, not missing import)
EdwardXuy pushed a commit to EdwardXuy/sgl-kernel-npu that referenced this pull request Aug 10, 2026
1. Revert pto-isa to af66b6a (pre-sgl-project#663) — sgl-project#663's 6081a2d3 causes SYNC_AIV_ONLY_ALL ambiguity in mega_kernel.cpp, breaking Build/Daily/DeepEP

2. Fix sglang stub: uninstall partial sglang (--no-deps installs real pkg whose __init__ needs orjson) then create clean stub providing only is_npu

3. F injection kept (works, fixes NameError in test_swiglu_quant)

No triton version change, no pto-isa bump — safe incremental fix on 9cb432b base
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.

1 participant