Skip to content

fix: LL proto ar dispatch - #4753

Open
TennyWang1223 wants to merge 1 commit into
mainfrom
LL_ar_new_dispatch
Open

fix: LL proto ar dispatch#4753
TennyWang1223 wants to merge 1 commit into
mainfrom
LL_ar_new_dispatch

Conversation

@TennyWang1223

Copy link
Copy Markdown
Contributor

Motivation

DeepSeek decode uses allreduce shapes (16, 7168) and (32, 7168) bf16 (224/448 KB). These fell through to the barrier-based ar_gfx1250_naive_unroll4
because the LL threshold was only 128 KiB. Raise it so these cases route to the faster flag-in-data ar_ll_gfx1250 kernel.

Technical Details

Changes in csrc/include/custom_all_reduce_gfx1250.cuh:

  • kLLArMaxBytes: 128 KiB → 4 MiB
  • kLLScratchCapBytes: 256 KiB → 4 MiB (scratch capacity to match)
  • allreduce_ll block size: bytes < 65536 → 512 threads, otherwise 256
  • ar_ll_gfx1250 __launch_bounds__: (256, 2)(512, 1)

Test Plan

python op_tests/multigpu_tests/gfx1250_poc/test_gfx1250_allreduce.py
-s 16,7168 -s 32,7168 -d bf16 -g true -t 4 --check-arch

Test Result

tp_size shape dtype size_kb min_us max_us err
4 (16, 7168) torch.bfloat16 224 8.884 9.172 0.000183
4 (32, 7168) torch.bfloat16 448 10.066 10.415 0.000166

Submission Checklist

Signed-off-by: TennyWang1223 <tennwang@amd.com>
@TennyWang1223
TennyWang1223 requested review from a team and valarLip August 14, 2026 06:55
@github-actions

Copy link
Copy Markdown
Contributor

🏷️ CI Guide

Runs automatically on every PR:

  • ✅ Pre-checks (submodule verification, code formatting)
  • ✅ Aiter op tests (gfx942 + gfx950)
  • ✅ Triton tests on MI35X (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label Tests
ci:gfx1250-ffm-triton Run the five-shard gfx1250 FFM Triton test suite
ci:triton-300x Run an additional Triton test job on MI300X in PRs; main branch always runs both MI35X and MI300X
ci:sglang SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm vLLM benchmark: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
ci:all All standard extended tests (excludes ci:atom_full)

Only add ci:atom_full for FlyDSL or Triton upgrades.
Add labels via the sidebar or gh pr edit 4753 --add-label <label>

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