Skip to content

feat: sync AITER operators, build fixes, and public dependencies - #13

Merged
lyt-ghub merged 43 commits into
mainfrom
sync_dcudas
Sep 29, 2026
Merged

lyt-ghub merged 43 commits into
mainfrom
sync_dcudas

Conversation

@longbisheng

Copy link
Copy Markdown
Contributor

Updates AITER sampling, top-k transforms, paged MQA logits, GDN, MoE, and collective communication paths. Includes gfx946 build fixes, W4A8 expert-parallel handling and configuration fixes, CI result reporting fixes, tuned configurations, and operator documentation.

Preserves the public GitHub Actions workflows and existing public CI customizations, including the current quality-gate workflow from main.

Dependencies

Validation

  • All 42 translated source snapshots passed metadata and tree verification.
  • Complete Linux checkout passed, including both submodules at their exact SHA, tree, and public origin.
  • Strict Git object checks and review-branch ancestry checks passed.
  • GPU compilation, operator correctness, and performance tests were not rerun for this synchronization. GitHub CI results are tracked on this PR.

xiaowei.zhang and others added 30 commits September 14, 2026 17:46
…he max_size limit for custom Allreduce input tensor
Migrate the fused metadata op from the Lightop Triton implementation to a
native HIP kernel and optimize it for paged-attention inference:

- fuse Phase-1 prefix scan (cache_seqlens_int32 / cu_seqlens_k) and
  Phase-2 page-table gather (page_table / swa_page_table) into a single
  kernel launch
- Phase 1: double-buffered LDS Hillis-Steele scan (B <= 64), row-group
  parallel scan with scalar group chain (B <= 1024), serial fallback for
  XL batches
- Phase 2: three tile variants (2048 / 1024 / 256 pages per block) with
  host-side workgroup-count routing calibrated at the measured U-curve
  bottom (WG ~ 256); no runtime branches inside the kernels
- zero-copy dtype dispatch over 8 int32/int64 combinations; USE_SWA and
  SHIFT (page_size = 1 vs > 1) resolved as template parameters
- add module_kvcache jit module, pybind entry, python wrapper and
  correctness/benchmark tests (op_tests/test_fused_metadata.py)

perf (12 configs, B=1..64, P=512..4096, page_size=1/32, swa on/off):
- vs Triton: 0.060-0.073ms -> 0.007-0.009ms, avg 8.73x, max 9.56x
- vs LightOP: (pre-migration): 0.008-0.012ms -> 0.007-0.009ms, avg 7.86x, max 9.07x
Document the fused_metadata operator (module_kvcache) and its three-way
comparison benchmark. The operator prepares paged-attention KV cache
metadata in a single kernel launch: cache_seqlens_int32, cu_seqlens_k
exclusive prefix sum, page_table gather, and an optional SWA page table.

- docs/fused_metadata.md: operator description, parameter contract,
  kernel variant routing and the AITER / LightOp / Triton perf report
- aiter算子库说明.md: new chapter 9 "KVCache" covering the operator
- op_tests/test_fused_metadata.py: group cases into v1-v6, add
  RuntimeError contract cases (invalid page_size, missing outputs,
  bad dtype/shape), and share one case definition between the pytest
  parametrization and the CLI

Benchmark result: AITER passes all 56 cases; LightOp fails 12
configurations, so its kernel time is excluded from its speedup stats.

Co-Authored-By: Claude Code <noreply@anthropic.com>
The fused_metadata report and its benchmark referred to the
pre-migration reference implementation by that project's internal
name. The name must not enter the open-sourced repository, so replace
every mention with a neutral one: 原实现 in Chinese prose, origin in
English comments and docstrings.

- docs/fused_metadata.md: 27 mentions renamed, CJK spacing cleaned up,
  and the resulting "源实现/源测试" duplications collapsed
- op_tests/test_fused_metadata.py: 9 mentions renamed; comments,
  docstrings and --case help text only, no identifier or import path
  affected, so behaviour and test semantics are unchanged
- aiter算子库说明.md: unify the KVCache chapter wording with the report

Text-only change. py_compile passes on the test module.

Co-Authored-By: Claude Code <noreply@anthropic.com>
@hygon-ai-reviewer

hygon-ai-reviewer Bot commented Sep 29, 2026 •

Copy link
Copy Markdown

AI Review

发现 1 个问题。

[P3] 索引按 int32 计算,超大张量下 offset 溢出导致越界访问

  • 位置:aiter/ops/triton/add_swiglu.py:21
  • 触发条件:当输出元素总数 N 超过 2^30(fp16 下 base/delta 各约 4 GiB)时触发
  • 证据与影响:_forward/_backward 中 i 与 offset = row*(2*D)+col 均为 int32 运算(tl.program_id/arange 默认 int32),A/B 的最大寻址偏移约为 2N-1,2N 超过 2^31 后回绕为负值或错位值,掩码 i < N 无法保护被回绕的偏移 对 A/B/DX 发生越界读写,产生静默数据损坏或非法内存访问
  • 修改建议:在计算 row/col/offset 前将 i 转为 int64(i.to(tl.int64)),前向与反向同根因一并修复

部分结果不可用:1/129 批未完整完成,已保留可用问题。

审查未完整完成:以 GitHub 返回的 PR diff 为主体,关联上下文仅用于核验,未运行代码。

覆盖情况:共 112 个文件,已审查 34 个,部分审查 2 个,待审查 73 个,跳过 3 个。

跳过原因:无 patch 0 个;文件类型过滤 3 个。

进度已保存。正在继续处理。

@lyt-ghub
lyt-ghub merged commit 2f1c2fc into main Sep 29, 2026
2 of 5 checks passed
@longbisheng
longbisheng deleted the sync_dcudas branch October 13, 2026 04:34
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.

4 participants