Skip to content

feat: add x_attention_v2 operator use pypto as compile backend. - #2286

Open
Yangxu-China wants to merge 1 commit into
xLLM-AI:mainfrom
Yangxu-China:x_attention_v2
Open

feat: add x_attention_v2 operator use pypto as compile backend.#2286
Yangxu-China wants to merge 1 commit into
xLLM-AI:mainfrom
Yangxu-China:x_attention_v2

Conversation

@Yangxu-China

@Yangxu-China Yangxu-China commented Aug 29, 2026

Copy link
Copy Markdown

Description

Add the x_attention_v2 PyPTO Pro kernel to the xllm NPU kernel package, along with its JIT test suite.

This PR introduces a new DSL sub-package xllm/python/kernels_npu/pypto/ (parallel to the existing tilelang/ and triton/ sub-packages) containing the x_attention_v2 operator — a beam-decode flash attention kernel for Ascend A5 (Ascend950), implementing shared KV + unshared KV dual-path merged attention with online softmax combine.

The kernel is written in the PyPTO Pro DSL (pypto_pro.language) and launched via the @pl.jit bracket-launch path. A Python host wrapper (x_attention_v2) computes tiling (core split) and launches the @pl.jit kernel (x_attention_v2_kernel).

Key additions:

  • xllm/python/kernels_npu/pypto/__init__.py — pypto DSL sub-package entry (license + docstring only, not registered in kernels_npu/__init__.py, consistent with tilelang/ and triton/)
  • xllm/python/kernels_npu/pypto/x_attention_v2.py@pl.jit kernel + Python host wrapper (1892 lines, self-contained)
  • tests/python/test_x_attention_v2_jit.py — 144 parametrized test cases (3 GQA configs × 2 batch sizes × 3 beam sizes × 4 KV lengths × 2 unshared lengths) with self-contained golden cache mechanism
  • .gitignore — entries for tests/python/golden_cache/ and tests/python/msprof*/

Verification:

  • 144/144 tests pass on Ascend950 — 13.6s with golden cache present, ~186s on first run without cache (CPU reference computation)
  • JIT compilation is cached per-process per-shape (first call ~2.4s compile, subsequent ~16ms cached)
  • msprof performance: total 6981us across 144 cases, avg 48.5us per case

Related Issues

N/A — new operator integration, no tracking issue yet.

Change Type

  • Bug fix
  • New feature
  • Performance improvement
  • Refactor
  • Documentation
  • Test
  • Build or CI

Pull Request Checklist

Thank you for contributing to xLLM. Before requesting review, please make sure the following items are complete.

PR Title and Commit Messages

  • The PR title and each commit message follow the xLLM commit format: <type>: <subject>.

Allowed types: feat, bugfix, docs, test, refactor, chore, style, revert, perf, model, build, release.
The subject should use clear English, start with a verb, include at least 4 words, and end with ..

Pre-commit Checks

  • I have installed pre-commit by running pip install pre-commit or an equivalent command.
  • I have installed the hooks with pre-commit install.
  • I have run pre-commit run --all-files and fixed any reported issues.

If you are unsure how to set up pre-commit, see the pre-commit documentation.

Self Review

  • I have self-reviewed the code according to .agents/skills/code-review/references/custom-code-style.md, especially code written or assisted by AI.
  • I have rebased this PR onto the latest main branch.

Build and Test Coverage

  • Tests have been added or updated as needed.
  • CUDA: python setup.py build test has passed on a CUDA machine.
  • NPU: python setup.py build test has passed on an NPU machine.
  • MLU: python setup.py build test has passed on an MLU machine.

Reviewer Notes

  • pypto_pro is a CANN-bundled package (installed at /home/developer/Ascend/cann-9.2.0/python/site-packages/pypto_pro/), not declared as a pip dependency — same model as tilelang (conditionally installed via CANN setup). The kernel file imports it lazily inside the host wrapper, so the module is importable without pypto_pro installed.
  • The golden_cache/ directory (144 .pt files) is gitignored — the test generates cache on first run if absent (~186s for all 144 CPU reference computations), subsequent runs load from disk (~14s).
  • msprof*/ directories are also gitignored — profiling results are not committed.
  • The kernel source (x_attention_v2.py) is self-contained with zero external path dependencies — all imports are standard library, pypto_pro, or torch.
  • The .gitignore change adds two entries for test artifacts.
  • Its operator prototype is consistent with the existing x_attention operator. It only supports the Ascend950 hardware and uses PyPTO as the backend compiler.
  • Good performance gains are observed across all 144 test cases within the test scope. Lab test results show that its average performance reaches 1.4× of the original x_attention Ascend950 implementation.

@Yangxu-China
Yangxu-China force-pushed the x_attention_v2 branch 2 times, most recently from 54927f2 to 5eb86b0 Compare August 29, 2026 03:26
@Yangxu-China Yangxu-China changed the title feat: Add x_attention_v2 operator feat: Add x_attention_v2 operator use pypto as compile backend. Aug 29, 2026
@liutongxuan liutongxuan changed the title feat: Add x_attention_v2 operator use pypto as compile backend. feat: add x_attention_v2 operator use pypto as compile backend. Aug 29, 2026
Comment thread .gitignore

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

这个不需要提交吧?

@zhang-minchao
zhang-minchao self-requested a review September 1, 2026 05:22
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.

2 participants