From c2581d95d71905dc04522fff58f9b46bce48b97b Mon Sep 17 00:00:00 2001 From: x Date: Sat, 10 Oct 2026 14:56:24 +0000 Subject: [PATCH 01/14] docs(agent): Qwen3-30B-A3B FP8 tensorwise Turbo campaign notes and scripts Handoff for the 8x MI355X Qwen3-30B-A3B FP8 tensorwise campaign (29,474 -> 40,283 tokens/s/GPU): cumulative e2e table, the Primus switches and Turbo/Primus branches the final stack needs, PR order and ready-to-paste PR bodies, how to rebuild the stack, rejected options and open items, plus the round log, the microbenchmark / trace / e2e scripts and the task prompt. --- .../qwen3_fp8_tw_gfx950_20261010/README.md | 143 ++++++++++ .../bench_attention_qwen3.py | 118 ++++++++ .../bench_deepep_qwen3.py | 98 +++++++ .../bench_dense_layouts.py | 66 +++++ .../bench_fp8_permute.py | 268 ++++++++++++++++++ .../bench_grouped_gemm_qwen3.py | 111 ++++++++ .../bench_permute_qwen3.py | 92 ++++++ .../bench_qk_rmsnorm_rope.py | 97 +++++++ .../bench_quant_chain.py | 111 ++++++++ .../dbg_aiter_fwd.py | 42 +++ .../dbg_permute_probs.py | 55 ++++ .../dbg_rope128.py | 39 +++ .../dbg_rope128_isa.py | 22 ++ .../dbg_rope128b.py | 30 ++ .../debug_fp8_worst.py | 98 +++++++ .../e2e/amend_msg.py | 49 ++++ .../e2e/e2e_ab.sh | 35 +++ .../e2e/e2e_chain.sh | 20 ++ .../e2e/e2e_queue2.sh | 9 + .../e2e/e2e_queue3.sh | 8 + .../e2e/e2e_queue4.sh | 8 + .../e2e/e2e_queue5.sh | 8 + .../e2e/e2e_run.sh | 13 + .../e2e/env_turbo_base.sh | 2 + .../e2e/env_turbo_opt.sh | 2 + .../e2e/env_turbo_opt_qknorm.sh | 3 + .../e2e/parse_e2e.py | 40 +++ .../pr_bodies/fix-moe-permute-padding-rows.md | 54 ++++ .../perf-moe-fp8-permute-tensorwise.md | 82 ++++++ .../pr_bodies/perf-moe-permute-default-hip.md | 60 ++++ .../qwen3_fp8_tw_gfx950_20261010/prompt.md | 104 +++++++ .../quant_bitexact.py | 48 ++++ .../quant_variants/quant_variants.hip | 160 +++++++++++ .../quant_variants/quant_variants_hip.hip | 160 +++++++++++ .../quant_variants/run_variants.py | 53 ++++ .../qwen3_fp8_tw_gfx950_20261010/qv_real.py | 15 + .../trace_attrib.py | 59 ++++ .../trace_kernels.py | 37 +++ .../trace_shapes.py | 60 ++++ 39 files changed, 2479 insertions(+) create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_attention_qwen3.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_deepep_qwen3.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_dense_layouts.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_fp8_permute.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_grouped_gemm_qwen3.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_permute_qwen3.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_qk_rmsnorm_rope.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_quant_chain.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_aiter_fwd.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_permute_probs.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128_isa.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128b.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/debug_fp8_worst.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/amend_msg.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_ab.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_chain.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue2.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue3.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue4.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue5.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_base.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt_qknorm.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/parse_e2e.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-moe-permute-padding-rows.md create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-fp8-permute-tensorwise.md create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-permute-default-hip.md create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/prompt.md create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_bitexact.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/quant_variants.hip create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/quant_variants_hip.hip create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/run_variants.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/qv_real.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_attrib.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_kernels.py create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_shapes.py diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md new file mode 100644 index 000000000..4512253d7 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -0,0 +1,143 @@ +# Qwen3-30B-A3B FP8 tensorwise on 8x MI355X: Turbo kernel campaign (handoff) + +Task: [prompt.md](prompt.md). Log of every round: [logs/optimize.md](logs/optimize.md). + +Status (2026-10-10): the 40,000 tokens/s/GPU target is reached. The baseline was 29,474 tokens/s/GPU and the best stack runs at 40,283 (mean of 4 runs). + +## End-to-end + +Qwen3-30B-A3B FP8 tensorwise pretrain on 8x MI355X: EP8, MBS 8, GBS 512, seq 4096, even routing, 20 iterations, mean of iterations 11-20. tokens/s/GPU = 262144 / (s/iter). Each row adds one change on top of the row above. + +| run | ms/iter | TFLOP/s/GPU | tokens/s/GPU | mem | loss@20 | +|---|---|---|---|---|---| +| baseline: submitted config, Turbo e2d9f1d7 | 8894.0 | 670.6 | 29,474 | 277 GB | 11.33393 | +| + `use_turbo_fused_act_with_probs` | 7915.6 | 753.5 | 33,117 | 223 GB | 11.33435 | +| Turbo 1103b2df (base of the Turbo branches) | 7925.4 | 752.5 | 33,077 | 222 GB | 11.33427 | +| + flat tensorwise FP8 quant kernel | 7870.9 | 757.7 | 33,305 | 222 GB | 11.33415 | +| + per-token `get_dispatch_layout` | 7774.2 | 767.2 | 33,720 | 222 GB | 11.33431 | +| + HIP permute in `DeepEPTokenDispatcher` | 7643.5 | 780.3 | 34,296 | 222 GB | 11.33423 | +| + `turbo_deepep_num_cu 160` | 7353.4 | 811.1 | 35,649 | 223 GB | 11.33456 | +| + fused qk RMSNorm + RoPE, head_dim 128 | 6987.1 | 853.6 | 37,518 | 224 GB | 11.33373 | +| + `turbo_fused_grouped_gemm` (2 runs: 6820.2, 6815.1) | 6817.7 | 874.9 | 38,451 | 224.5 GB | 11.33412 | +| + `turbo_fp8_permute` (4 runs: 6442.4, 6515.0, 6523.8, 6549.1) | 6507.6 | 916.7 | 40,283 | 223.3 GB | 11.33404 | + +Run-to-run loss noise is about 4e-4 at iteration 20. + +## Settings for the last row + +The final configuration is `examples/megatron/configs/MI355X/qwen3_30B_A3B-FP8-pretrain.yaml` plus: +- `--use_turbo_fused_act_with_probs True` +- `--turbo_deepep_num_cu 160` +- `--turbo_fused_grouped_gemm True` +- `--turbo_fp8_permute True` +- env `PRIMUS_FUSED_QK_RMSNORM_ROPE=1` + +Code it needs: +- Turbo: all five Turbo branches below. Building on `perf/moe/fp8-permute-tensorwise` also brings in the permute-default change. +- Primus: `perf/qwen3-qk-rmsnorm-rope-hd128`, which sits on `perf/qwen3-30b-a3b-tuning`, plus `perf/megatron/turbo-fp8-permute`. + +See [e2e/e2e_run.sh](e2e/e2e_run.sh) and [e2e/env_turbo_opt_qknorm.sh](e2e/env_turbo_opt_qknorm.sh). Their paths are for the old machine. + +## Branches and PRs + +Turbo (`AMD-AGI/Primus-Turbo`): + +| branch | head | PR | notes | +|---|---|---|---| +| `fix/moe/permute-padding-rows` | f1a59fb9 | #553 merged | TURBO permute padding rows with worst-case buffers + pad | +| `perf/moe/permute-default-hip` | 30b6cc7c | open next | body: [pr_bodies/perf-moe-permute-default-hip.md](pr_bodies/perf-moe-permute-default-hip.md) | +| `perf/moe/fp8-permute-tensorwise` | 83e712a9 | after permute-default | stacked on it; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | +| `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | not opened | independent; flat no-pad quant kernel | +| `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | not opened | independent; `get_dispatch_layout` 340 -> 45 us | +| `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | not opened | independent; needed by the Primus qk-norm patch | +| `fix/fp8-attn-strided-layout` | c825b13e | not opened | FP8 attention with sbhd/bhsd storage | + +Before opening the FP8 permute PR, rebase `perf/moe/fp8-permute-tensorwise` onto main once permute-default has merged: +- Drop a04d011b, its duplicate of #553. +- Put the permute-default PR number into the body. +- The commit message quotes the single-run gain (-5.54%); the PR body uses the 4-run mean (-4.55%). + +Primus (`AMD-AGI/Primus`): + +| branch | head | PR | notes | +|---|---|---|---| +| `perf/qwen3-30b-a3b-tuning` | 82ebf666 | not opened | Qwen3 FP8 config tuned to Turbo tensorwise | +| `perf/qwen3-qk-rmsnorm-rope-hd128` | bfa227f2 | not opened | on top of the tuning branch; fused qk-norm patch for head_dim 128 | +| `perf/megatron/turbo-fp8-permute` | 69a32f64 | not opened | on main; `turbo_fp8_permute` flag | + +## Rebuilding the stack on a new machine + +All of these merges are clean as of 2026-10-10. + +Turbo: + +```bash +git fetch origin && git checkout -b integ/qwen3-fp8 origin/main +for b in perf/quantization/tensorwise-fp8-qwen3 perf/deep_ep/dispatch-layout-per-token \ + perf/flydsl/qk-rmsnorm-rope-hd128 perf/moe/fp8-permute-tensorwise; do + git merge --no-edit origin/$b || break +done +GPU_ARCHS=gfx950 MAX_JOBS=64 python setup.py build_ext --inplace +``` + +Primus: + +```bash +git fetch origin && git checkout -b integ/qwen3-fp8 origin/perf/qwen3-qk-rmsnorm-rope-hd128 +git merge --no-edit origin/perf/megatron/turbo-fp8-permute +git submodule update --init third_party/Megatron-LM +``` + +Then edit the paths in `e2e/env_turbo_opt_qknorm.sh` and `e2e/e2e_run.sh` for the new machine, and run, for example: + +```bash +e2e/e2e_run.sh g1 --turbo_fused_grouped_gemm True --turbo_fp8_permute True +``` + +Summarize the runs with `e2e/parse_e2e.py `, which reads `/tmp/q3_e2e_.log`. + +Run Python from the campaign directory with `PYTHONPATH=`. Otherwise the pip-installed Turbo is imported. + +## Findings and rejected options + +- **MegaMoE** (bf16 and mxfp8) runs out of memory at MBS 8. When the worst-case dispatch pool allocation fails, 261 GB (bf16, failing on 3 GiB) or 264 GB (mxfp8, failing on 1.5 GiB) is already allocated. Using it would need a MegaMoE change, a smaller MBS, or recompute. +- **`turbo_deepep_num_cu`**: + - 192 gives a NaN forward loss on rank 2 at iteration 5, a DeepEP bug at high channel counts. + - 256 is slower (7028 ms). + - 160 is best. +- **Sync-free MoE stage 2** costs +300 ms over stage 1. `num_worst_tokens` makes every post-dispatch step scan 262144 receive rows instead of 32768 tokens. +- **`turbo_grouped_gemm_without_padding`** needs `moe_router_padding_for_quantization False`, which sync-free stage 1 forces to True. +- **Attention**: TE CK is the fastest. The Turbo AITER forward is 2x slower (causal kernel without tile skipping), so keep `use_turbo_attention` false. +- **#455** made TRITON the generic `moe_permute` default, and it is slower than HIP at every measured shape. `perf/moe/permute-default-hip` switches the dispatcher back. +- **MXFP4 tests**: 20 MXFP4 quantization tests also fail on the original build (`x.is_contiguous()` assertion in MXFP4 dequant). Unrelated to this work. +- **TURBO permute requires routed tokens to come first.** DeepEP guarantees this. The worst-case + pad padding-row bug is fixed in #553. TURBO unpermute still leaves rows past the routed count unwritten; this is harmless because DeepEP combine reads only received rows. + +## Open items + +- **Unexplained e2e jitter with FP8 permute.** The first run took 6442 ms; three later runs took 6515-6549 ms with more per-iteration jitter. These causes are ruled out: + - the Primus change: a same-code rerun gave 6515 ms; + - the padding fix: permute kernel times are unchanged; + - the node: the `turbo_fused_grouped_gemm` baseline reran at 6815 ms. + The cause is not identified. +- **Recompute**: `turbo_fp8_permute` rejects selective recompute of `moe_act`, because the checkpointed fused-MLP path is not validated with `QuantizedTensor` inputs and gradients. +- **Untested on the final build**: the 12 MXFP8 tests in `tests/pytorch/ops/test_grouped_mlp_fp8.py`. The MXFP8 path is untouched, and FlyDSL compilation takes more than 30 minutes. +- **Remaining hot spots** (R3 profile, rank 0, one step, taken before the fused MLP and the FP8 permute): + + | kernel group | ms | + |---|---| + | attention | 1583 | + | grouped GEMM | 1268 | + | hipBLASLt | 1003 | + | DeepEP | 876 | + | quant chain | 739 | + | permute / unpermute | 346 | + | SwiGLU | 262 | + +## Files + +- `bench_*.py`: per-item microbenchmarks at Qwen3 shapes. `bench_fp8_permute.py --time` gives the kernel breakdown of the FP8 permute. +- `trace_*.py`: kernel-time attribution from torch profiler traces. +- `debug_*.py` and `dbg_*.py`: one-off reproducers. +- `quant_variants/`: tensorwise quant kernel variants, source only. +- `e2e/`: run, queue and A/B scripts, the env files, `parse_e2e.py`, and `amend_msg.py`, which writes the cumulative table into a commit message. +- `pr_bodies/`: PR descriptions, ready to paste. diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_attention_qwen3.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_attention_qwen3.py new file mode 100644 index 000000000..165e8736c --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_attention_qwen3.py @@ -0,0 +1,118 @@ +"""Qwen3-30B-A3B attention (sbhd, GQA 32/4, hd128, causal, bf16): TE CK vs Turbo backends. + +Env: SHAPE=S,B,HQ,HKV,D (default 4096,8,32,4,128), BACKENDS=te,flydsl,aiter +""" + +import os + +os.environ.setdefault("NVTE_FLASH_ATTN", "0") +os.environ.setdefault("NVTE_FUSED_ATTN", "1") +os.environ.setdefault("NVTE_CK_USES_BWD_V3", "1") + +import torch # noqa: E402 (torch before TE: LLVM option clash otherwise) + +S, B, HQ, HKV, D = (int(v) for v in os.environ.get("SHAPE", "4096,8,32,4,128").split(",")) +BACKENDS = os.environ.get("BACKENDS", "te,flydsl,aiter").split(",") +ITERS = int(os.environ.get("ITERS", "20")) +dev = "cuda" + + +def timed(fn, iters=ITERS, warmup=5): + for _ in range(warmup): + fn() + torch.cuda.synchronize() + ts = [] + for _ in range(iters): + s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + s.record() + fn() + e.record() + torch.cuda.synchronize() + ts.append(s.elapsed_time(e) * 1e3) + return sorted(ts)[len(ts) // 2] + + +def make_te(): + import transformer_engine.pytorch as te + + attn = te.DotProductAttention( + HQ, D, num_gqa_groups=HKV, attention_dropout=0.0, attn_mask_type="causal", qkv_format="sbhd" + ) + + def run(q, k, v): + return attn(q, k, v).view(S, B, HQ, D) + + return run + + +def make_turbo(backend_name): + from primus_turbo.pytorch.core.backend import BackendType, GlobalBackendManager + from primus_turbo.pytorch.ops import flash_attn_func + + backend = getattr(BackendType, backend_name.upper()) + + def run(q, k, v): + GlobalBackendManager.set_attn_backend(backend) + o = flash_attn_func( + q.permute(1, 0, 2, 3), k.permute(1, 0, 2, 3), v.permute(1, 0, 2, 3), causal=True + ) + GlobalBackendManager.set_attn_backend(None) + return o.permute(1, 0, 2, 3) + + return run + + +torch.manual_seed(0) +q0 = torch.randn(S, B, HQ, D, device=dev, dtype=torch.bfloat16) +k0 = torch.randn(S, B, HKV, D, device=dev, dtype=torch.bfloat16) +v0 = torch.randn(S, B, HKV, D, device=dev, dtype=torch.bfloat16) +do = torch.randn(S, B, HQ, D, device=dev, dtype=torch.bfloat16) + +fwd_flops = 4 * B * HQ * S * S * D / 2 +results, outs = {}, {} +for name in BACKENDS: + run = make_te() if name == "te" else make_turbo(name) + q, k, v = (t.clone().requires_grad_() for t in (q0, k0, v0)) + try: + o = run(q, k, v) + o.backward(do) + except Exception as exc: # noqa: BLE001 + print(f"{name}: unsupported ({type(exc).__name__}: {str(exc)[:200]})") + continue + outs[name] = (o.detach().float(), q.grad.float(), k.grad.float(), v.grad.float()) + + def fwd(): + with torch.no_grad(): + run(q0, k0, v0) + + def fwdbwd(): + q.grad = k.grad = v.grad = None + run(q, k, v).backward(do) + + t_f, t_fb = timed(fwd), timed(fwdbwd) + results[name] = (t_f, t_fb) + if os.environ.get("PROFILE") == "1": + with torch.profiler.profile(activities=[torch.profiler.ProfilerActivity.CUDA]) as prof: + for _ in range(5): + fwdbwd() + torch.cuda.synchronize() + print(f"--- {name} kernels (5 x fwd+bwd)") + print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=12, max_name_column_width=70)) + +print(f"shape S={S} B={B} HQ={HQ} HKV={HKV} D={D} causal bf16") +for name, (t_f, t_fb) in results.items(): + t_b = t_fb - t_f + print( + f"{name:8s} fwd {t_f:8.1f} us ({fwd_flops / t_f / 1e6:6.0f} TF/s) bwd {t_b:8.1f} us " + f"({2.5 * fwd_flops / t_b / 1e6:6.0f} TF/s) fwd+bwd {t_fb:8.1f} us" + ) +if "te" in outs: + ref = outs["te"] + for name, got in outs.items(): + if name == "te": + continue + diffs = " ".join( + f"{tag} max {(a - b).abs().max().item():.3g}" + for tag, a, b in zip(("o", "dq", "dk", "dv"), ref, got) + ) + print(f"{name} vs te: {diffs}") diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_deepep_qwen3.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_deepep_qwen3.py new file mode 100644 index 000000000..741d44c18 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_deepep_qwen3.py @@ -0,0 +1,98 @@ +"""Item 3 microbench: Turbo DeepEP intranode at Qwen3-30B-A3B EP8 (32768 tokens/rank, hidden 2048, +128 experts, top-8). Default dispatch/combine configs (what training uses); sweeps num_sms and +compares BF16 vs FP8 (per-128 scales) dispatch, under random top-k and Primus "even" routing. +Each phase is timed separately with CUDA events (median over ITERS, max over ranks); the +non-cached dispatch includes its CPU wait on the notify kernel, the cached one does not. +""" + +import os +import statistics +import sys + +import torch +import torch.distributed as dist + +sys.path.insert(0, os.path.join(os.path.dirname(__file__), "../../../benchmark/ops/training")) +from deep_ep.utils import init_dist # noqa: E402 + +from primus_turbo.pytorch import deep_ep # noqa: E402 + +T, H, E, K = 32768, 2048, 128, 8 +SMS = [int(s) for s in os.environ.get("SMS", "64,80,96").split(",")] +ITERS = 20 + + +def topk_for(routing, rank): + if routing == "even": # PrimusTopKRouter._force_even_routing + return (torch.arange(T * K, device="cuda").view(T, K) % E).to(torch.int64) + g = torch.Generator(device="cuda").manual_seed(rank) + scores = torch.randn((T, E), device="cuda", generator=g).abs() + 1 + return torch.topk(scores, K, dim=-1, sorted=False)[1] + + +def timed(fn): + s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + torch.cuda.synchronize() + s.record() + out = fn() + e.record() + torch.cuda.synchronize() + return out, s.elapsed_time(e) * 1e3 + + +def worker(local_rank, num_local_ranks): + rank, num_ranks, group = init_dist(local_rank, num_local_ranks) + buf = deep_ep.Buffer(group, int(2e9), 0, low_latency_mode=False, explicitly_destroy=True) + x = torch.randn(T, H, device="cuda", dtype=torch.bfloat16) + x_fp8 = x.to(torch.float8_e4m3fn) + x_scales = torch.ones(T, H // 128, device="cuda", dtype=torch.float32) + rows = [] + for routing in ("random", "even"): + topk_idx = topk_for(routing, rank) + topk_w = torch.rand(T, K, device="cuda", dtype=torch.float32) + for sms in SMS: + deep_ep.Buffer.set_num_sms(sms) + for prec in ("bf16", "fp8"): + inp = x if prec == "bf16" else (x_fp8, x_scales) + ts = {"layout": [], "dispatch": [], "dispatch_cached": [], "combine": []} + for it in range(ITERS + 3): + dist.barrier() + (ntr, _, nte, itr, _), t_l = timed(lambda: buf.get_dispatch_layout(topk_idx, E)) + dist.barrier() + (recv, _, rw, _, handle, _), t_d = timed( + lambda: buf.dispatch( + inp, + num_tokens_per_rank=ntr, + is_token_in_rank=itr, + num_tokens_per_expert=nte, + topk_idx=topk_idx, + topk_weights=topk_w, + ) + ) + dist.barrier() + _, t_dc = timed(lambda: buf.dispatch(inp, handle=handle)) + r = recv[0] if isinstance(recv, tuple) else recv + t_c = 0.0 + if prec == "bf16": + dist.barrier() + _, t_c = timed(lambda: buf.combine(r, handle, topk_weights=rw)) + if it >= 3: + for k, v in zip(ts, (t_l, t_d, t_dc, t_c)): + ts[k].append(v) + med = torch.tensor([statistics.median(v) for v in ts.values()], device="cuda") + dist.all_reduce(med, op=dist.ReduceOp.MAX) + rows.append((routing, sms, prec, r.shape[0], med.tolist())) + if rank == 0: + print("routing sms prec recv_tokens | layout dispatch dispatch_cached combine (us, max over ranks)") + for routing, sms, prec, nrecv, m in rows: + print( + f"{routing:6s} sms={sms:3d} {prec:4s} recv={nrecv:6d} | " + f"layout {m[0]:6.0f} dispatch {m[1]:6.0f} cached {m[2]:6.0f} combine {m[3]:6.0f}" + ) + buf.destroy() + dist.barrier() + dist.destroy_process_group() + + +if __name__ == "__main__": + torch.multiprocessing.spawn(worker, args=(8,), nprocs=8) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_dense_layouts.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_dense_layouts.py new file mode 100644 index 000000000..00c868257 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_dense_layouts.py @@ -0,0 +1,66 @@ +"""hipBLASLt FP8 tensorwise: NT-with-explicit-transposes (current gfx950 bwd) vs NN / TN. + +Shapes: Qwen3-30B-A3B dense projections per microbatch (M = 32768 tokens). +""" + +import torch + +from primus_turbo.pytorch.core.backend import BackendType +from primus_turbo.pytorch.core.low_precision import ScalingGranularity +from primus_turbo.pytorch.kernels.gemm.gemm_fp8_impl import gemm_fp8_impl + +TW = ScalingGranularity.TENSORWISE.value +HB = BackendType.HIPBLASLT.value +E4, E5 = torch.float8_e4m3fn, torch.float8_e5m2 + + +def bench(fn, iters=30): + for _ in range(5): + fn() + s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + torch.cuda.synchronize() + s.record() + for _ in range(iters): + fn() + e.record() + torch.cuda.synchronize() + return s.elapsed_time(e) * 1e3 / iters + + +def tr(x): + return torch.ops.primus_turbo_cpp_extension.transpose_2d(x, 0, 1) + + +def g(a, ta, b, tb, tc=False): + one = torch.ones((), device="cuda") + return gemm_fp8_impl(a, one, ta, b, one, tb, torch.bfloat16, tc, granularity=TW, default_backend=HB) + + +M = 32768 +for name, N, K in [("qkv", 5120, 2048), ("proj", 2048, 4096)]: + x = (torch.randn(M, K, device="cuda") * 0.5).to(E4) + w = (torch.randn(N, K, device="cuda") * 0.5).to(E4) + go = (torch.randn(M, N, device="cuda") * 0.5).to(E5) + fl = 2 * M * N * K + + fwd = bench(lambda: g(x, False, w, True)) + # dgrad: current = transpose(w) + NT ; alt = NN + w_t = tr(w) + d_nt = bench(lambda: g(go, False, w_t, True)) + d_tr = bench(lambda: tr(w)) + d_nn = bench(lambda: g(go, False, w, False)) + # wgrad: current = transpose(x) [fwd] + transpose(go) + NT(a_t, go_t, trans_c) ; alt = TN(go, x) + x_t, go_t = tr(x), tr(go) + w_nt = bench(lambda: g(x_t, False, go_t, True, True)) + w_tr = bench(lambda: tr(x)) + bench(lambda: tr(go)) + w_tn = bench(lambda: g(go, True, x, False)) + ref_d = g(go, False, w_t, True).float() + alt_d = g(go, False, w, False).float() + ref_w = g(x_t, False, go_t, True, True).float() + alt_w = g(go, True, x, False).float() + ed = ((ref_d - alt_d).norm() / ref_d.norm()).item() + ew = ((ref_w - alt_w).norm() / ref_w.norm()).item() + print(f"[{name}] M={M} N={N} K={K}") + print(f" fwd NT : {fwd:7.1f} us {fl / fwd / 1e6:6.0f} TFLOP/s") + print(f" dgrad NT+transpose: {d_nt:7.1f} + {d_tr:5.1f} us NN: {d_nn:7.1f} us ({fl / d_nn / 1e6:4.0f} TF/s) rel diff {ed:.1e}") + print(f" wgrad NT+transpose: {w_nt:7.1f} + {w_tr:5.1f} us TN: {w_tn:7.1f} us ({fl / w_tn / 1e6:4.0f} TF/s) rel diff {ew:.1e}") diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_fp8_permute.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_fp8_permute.py new file mode 100644 index 000000000..b7f340dd1 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_fp8_permute.py @@ -0,0 +1,268 @@ +"""Validate + time tensorwise-FP8 permute (fwd) and FP8 unpermute grad (bwd) at Qwen3 shapes. + +Per-rank shape (EP8, MBS8, seq 4096): 32768 dispatched tokens, H=2048, 16 local experts, +topk 8 with all 8 experts on one rank (even routing), pad 16. +""" + +import sys + +import torch + +from primus_turbo.pytorch.core.backend import BackendType +from primus_turbo.pytorch.core.low_precision import ( + Float8QuantConfig, + ScalingGranularity, + float8_e4m3, + float8_e5m2, +) +from primus_turbo.pytorch.core.quantized_tensor import QuantizedTensor +from primus_turbo.pytorch.ops.grouped_mlp_fp8 import grouped_mlp_fp8 +from primus_turbo.pytorch.ops.moe.moe_permute import moe_permute, moe_unpermute + +T, H, E, TOPK, PAD, I = 32768, 2048, 16, 8, 16, 768 +TURBO = BackendType.TURBO +dev = "cuda" +torch.manual_seed(0) + + +def make_inputs(even=True): + if even: + # all 8 experts of a token on this rank, round robin + base = (torch.arange(T, device=dev) * TOPK) % E + idx = (base[:, None] + torch.arange(TOPK, device=dev)[None]) % E + else: + idx = torch.argsort(torch.rand(T, E, device=dev), dim=1)[:, :TOPK] + idx = idx.to(torch.int64).contiguous() + tokens = torch.randn(T, H, device=dev, dtype=torch.bfloat16) + probs = torch.rand(T, TOPK, device=dev, dtype=torch.float32) + return tokens, idx, probs + + +def permute(tokens, idx, probs, qdtype=None): + return moe_permute( + tokens, + topk_indices=idx, + num_local_experts=E, + num_topk=TOPK, + pad_multiple=PAD, + num_permuted_tokens=-1, + probs=probs, + probs_layout="topk", + backend=TURBO, + quantize_dtype=qdtype, + ) + + +def bench(fn, iters=50): + for _ in range(5): + fn() + torch.cuda.synchronize() + s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + s.record() + for _ in range(iters): + fn() + e.record() + torch.cuda.synchronize() + return s.elapsed_time(e) / iters * 1e3 + + +def check_permute(even): + tokens, idx, probs = make_inputs(even) + # an unrouted outlier token must not reach the scale + idx[5] = -1 + tokens[5] = 3.0e4 + ref, ref_map, ref_tpe, _, _, ref_probs = permute(tokens, idx, probs) + ref_q = QuantizedTensor.quantize( + ref, float8_e4m3, ScalingGranularity.TENSORWISE, axis=-1, group_lens=ref_tpe, pad_align_last=128 + ) + out, row_map, tpe, _, _, out_probs = permute(tokens, idx, probs, float8_e4m3) + assert isinstance(out, QuantizedTensor) and out._is_grouped_tensor + ok = ( + torch.equal(out.qdata.view(torch.uint8), ref_q.qdata.view(torch.uint8)) + and torch.equal(out.scale_inv, ref_q.scale_inv) + and torch.equal(out.group_offs, ref_q.group_offs) + and torch.equal(tpe, ref_tpe) + and torch.equal(out_probs, ref_probs) + and out.shape == ref.shape + ) + print(f"permute fwd bit-exact (even={even}): {ok} scale_inv={out.scale_inv.item():.6e}") + return ok + + +class _Capture(torch.autograd.Function): + grads = [] + + @staticmethod + def forward(ctx, x): + return x.clone() + + @staticmethod + def backward(ctx, g): + _Capture.grads.append(g) + return torch.zeros(g.shape, dtype=torch.bfloat16, device=dev) + + +def check_unpermute_grad(even): + tokens, idx, probs = make_inputs(even) + idx[7] = -1 + permuted, row_map, tpe, _, _, _ = permute(tokens, idx, probs) + gout = torch.randn(T, H, device=dev, dtype=torch.bfloat16) + gout[7] = 3.0e4 + + x = permuted.detach().requires_grad_() + out, _ = moe_unpermute( + x, row_map, restore_shape=tokens.shape, num_local_experts=E, pad_multiple=PAD, backend=TURBO + ) + (ref_grad,) = torch.autograd.grad(out, x, gout) + ref_q = QuantizedTensor.quantize(ref_grad, float8_e5m2, ScalingGranularity.TENSORWISE, axis=-1, pad_align_last=128) + + _Capture.grads.clear() + x2 = permuted.detach().requires_grad_() + out2, _ = moe_unpermute( + _Capture.apply(x2), + row_map, + restore_shape=tokens.shape, + num_local_experts=E, + pad_multiple=PAD, + backend=TURBO, + grad_quantize_dtype=float8_e5m2, + ) + out2.backward(gout) + q = _Capture.grads[0] + ok = ( + isinstance(q, QuantizedTensor) + and torch.equal(q.qdata.view(torch.uint8), ref_q.qdata.view(torch.uint8)) + and torch.equal(q.scale_inv, ref_q.scale_inv) + and q.shape == ref_grad.shape + ) + print(f"unpermute grad bit-exact (even={even}): {ok}") + return ok + + +def check_mlp_e2e(fmt): + """permute -> grouped_mlp_fp8 -> unpermute: FP8 handoff vs bf16 handoff, all bitwise.""" + from primus_turbo.pytorch.ops.utils import _get_fp8_dtype + + tokens, idx, probs = make_inputs(True) + w1 = (torch.randn(E, 2 * I, H, device=dev) * 0.02).bfloat16() + w2 = (torch.randn(E, H, I, device=dev) * 0.02).bfloat16() + gout = torch.randn(T, H, device=dev, dtype=torch.bfloat16) + cfg = Float8QuantConfig(format=fmt) + x_dtype, g_dtype = _get_fp8_dtype(fmt, True), _get_fp8_dtype(fmt, False) + print(f" mlp e2e format={fmt} x={x_dtype} grad={g_dtype}") + + def run(fp8): + t = tokens.clone().requires_grad_() + p = probs.clone().requires_grad_() + a, b = w1.clone().requires_grad_(), w2.clone().requires_grad_() + x, row_map, tpe, _, _, pp = permute(t, idx, p, x_dtype if fp8 else None) + y = grouped_mlp_fp8(x, a, b, tpe, probs=pp, trans_w1=True, trans_w2=True, config=cfg, activation="silu") + o, _ = moe_unpermute( + y, + row_map, + restore_shape=t.shape, + num_local_experts=E, + pad_multiple=PAD, + backend=TURBO, + grad_quantize_dtype=g_dtype if fp8 else None, + ) + o.backward(gout) + return [o.detach(), t.grad, p.grad, a.grad, b.grad] + + ref, new = run(False), run(True) + names = ["out", "grad_tokens", "grad_probs", "grad_w1", "grad_w2"] + ok = True + for n, r, g in zip(names, ref, new): + same = torch.equal(r, g) + ok &= same + print(f" mlp e2e {n}: bit-exact={same}") + return ok + + +def timing(): + tokens, idx, probs = make_inputs(True) + permuted, row_map, tpe, _, _, _ = permute(tokens, idx, probs) + ndisp = torch.full((1,), T, dtype=torch.int32, device=dev) + + def ref_fwd(): + p, _, tpe_, _, _, _ = permute(tokens, idx, probs) + QuantizedTensor.quantize( + p, float8_e4m3, ScalingGranularity.TENSORWISE, axis=-1, group_lens=tpe_, pad_align_last=128 + ) + + def new_fwd(): + permute(tokens, idx, probs, float8_e4m3) + + def perm_only(): + permute(tokens, idx, probs) + + from primus_turbo.pytorch.kernels.moe.moe_permute_impl import ( + moe_permute_routed_amax_impl, + ) + + def amax_only(): + moe_permute_routed_amax_impl(tokens, row_map, ndisp, E) + + print(f"fwd bf16 permute only : {bench(perm_only):8.1f} us") + print(f"fwd bf16 permute + quantize : {bench(ref_fwd):8.1f} us") + print(f"fwd fp8 permute (new) : {bench(new_fwd):8.1f} us") + print(f" routed amax only : {bench(amax_only):8.1f} us") + + gout = torch.randn(T, H, device=dev, dtype=torch.bfloat16) + + def bwd(qd): + x = permuted.detach().requires_grad_() + y = _Capture.apply(x) if qd is not None else x + out, _ = moe_unpermute( + y, row_map, restore_shape=tokens.shape, num_local_experts=E, pad_multiple=PAD, backend=TURBO, + grad_quantize_dtype=qd, + ) + if qd is None: + (g,) = torch.autograd.grad(out, x, gout) + QuantizedTensor.quantize(g, float8_e5m2, ScalingGranularity.TENSORWISE, axis=-1, pad_align_last=128) + else: + out.backward(gout) + _Capture.grads.clear() + + from primus_turbo.pytorch.kernels.moe.moe_permute_impl import moe_permute_impl + from primus_turbo.pytorch.ops.moe.moe_permute import _permute_quantized_tensorwise + + num_permuted = permuted.shape[0] + + def ref_bwd(): + g, _, _ = moe_permute_impl( + TURBO, gout, row_map, None, num_permuted, E, H, PAD, None, None, 0, False, 0 + ) + QuantizedTensor.quantize(g, float8_e5m2, ScalingGranularity.TENSORWISE, axis=-1, pad_align_last=128) + + def new_bwd(): + _permute_quantized_tensorwise(gout, row_map, None, num_permuted, E, PAD, None, 0, float8_e5m2, None) + + print(f"bwd bf16 permute(grad) + quantize: {bench(ref_bwd):8.1f} us") + print(f"bwd fp8 permute(grad) (new) : {bench(new_bwd):8.1f} us") + + from torch.profiler import ProfilerActivity, profile + + for name, fn in (("ref_fwd", ref_fwd), ("new_fwd", new_fwd), ("ref_bwd", ref_bwd), ("new_bwd", new_bwd)): + fn() + torch.cuda.synchronize() + with profile(activities=[ProfilerActivity.CUDA]) as prof: + for _ in range(10): + fn() + torch.cuda.synchronize() + print(f"--- {name} per-call kernel times") + for ev in sorted(prof.key_averages(), key=lambda e: -e.device_time_total): + if ev.device_time_total > 0: + print(f" {ev.device_time_total / 10:8.1f} us {ev.key[:90]}") + + +if __name__ == "__main__": + ok = check_permute(True) & check_permute(False) + ok &= check_unpermute_grad(True) & check_unpermute_grad(False) + from primus_turbo.pytorch.core.low_precision import Format + + ok &= check_mlp_e2e(Format.E4M3) + ok &= check_mlp_e2e(Format.HYBRID) + print("ALL BIT-EXACT:", ok) + if "--time" in sys.argv: + timing() diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_grouped_gemm_qwen3.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_grouped_gemm_qwen3.py new file mode 100644 index 000000000..fd52a2474 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_grouped_gemm_qwen3.py @@ -0,0 +1,111 @@ +"""Item 2 microbench: FlyDSL FP8 tensorwise grouped GEMM at Qwen3-30B-A3B EP8 shapes. + +16 local experts, 262144 routed rows per rank per microbatch. Reports the GEMM +kernels only (fwd nt / dgrad nn / wgrad tn); quantization is excluded. +""" + +import argparse +import collections +import re + +import torch +from torch.profiler import ProfilerActivity, profile + +import primus_turbo.pytorch as turbo +from primus_turbo.pytorch.core.low_precision import ( + Float8QuantConfig, + Format, + ScalingGranularity, +) + +G, M_TOTAL = 16, 262144 +SHAPES = {"fc1": (1536, 2048), "fc2": (2048, 768)} # (N, K) with trans_b + + +def group_lens_for(dist, gen): + if dist == "even": + return torch.full((G,), M_TOTAL // G, dtype=torch.int64) + if dist == "uniform": # each row routed uniformly at random + idx = torch.randint(0, G, (M_TOTAL,), generator=gen) + return torch.bincount(idx, minlength=G).to(torch.int64) + if dist == "skew": # bench_grouped_gemm_turbo's unbalanced generator + d = 0.2 + 0.8 * torch.rand(G, generator=gen) + lens = (d / d.sum() * M_TOTAL).to(torch.int64) + lens[-1] += M_TOTAL - lens.sum() + return lens + if dist == "hot": # one expert takes half the rows, two experts empty + lens = torch.zeros(G, dtype=torch.int64) + lens[0] = M_TOTAL // 2 + rest = M_TOTAL - lens[0] + lens[1:14] = rest // 13 + lens[13] = rest - lens[1:13].sum() + return lens + raise ValueError(dist) + + +def classify(name): + if "grouped_nt" in name: + return "fwd" + if "grouped_nn" in name: + return "dgrad" + if "wgrad" in name or "grouped_tn" in name: + return "wgrad" + return None + + +def run_case(name, dist, iters, gen): + N, K = SHAPES[name] + lens = group_lens_for(dist, gen).cuda() + M = int(lens.sum()) + cfg = Float8QuantConfig(format=Format.E4M3, granularity=ScalingGranularity.TENSORWISE) + a = torch.randn(M, K, device="cuda", dtype=torch.bfloat16, requires_grad=True) + w = torch.nn.Parameter(torch.randn(G, N, K, device="cuda", dtype=torch.bfloat16) * 0.02) + w.main_grad = torch.zeros(G, N, K, device="cuda", dtype=torch.float32) + w.grad_added_to_main_grad = False + go = torch.randn(M, N, device="cuda", dtype=torch.bfloat16) + + def step(): + out = turbo.ops.grouped_gemm_fp8(a, w, lens, trans_b=True, config=cfg, fuse_bgrad_accum_pattern="megatron") + out.backward(go) + a.grad = None + + for _ in range(3): + step() + torch.cuda.synchronize() + with profile(activities=[ProfilerActivity.CUDA]) as prof: + for _ in range(iters): + step() + torch.cuda.synchronize() + t = collections.Counter() + names = {} + for e in prof.events(): + if e.device_type == torch.autograd.DeviceType.CUDA: + c = classify(e.name) + if c: + t[c] += e.device_time + names[c] = re.sub(r"\.kd$", "", e.name) + flops = 2 * M * N * K + res = {c: (t[c] / iters, flops / (t[c] / iters) / 1e6) for c in ("fwd", "dgrad", "wgrad")} + tot = sum(v[0] for v in res.values()) + return M, res, 3 * flops / tot / 1e6, names + + +def main(): + p = argparse.ArgumentParser() + p.add_argument("--iters", type=int, default=20) + p.add_argument("--dists", default="even,uniform,skew,hot") + args = p.parse_args() + gen = torch.Generator().manual_seed(1234) + print(f"{'case':14s} {'M':>7s} | {'fwd us':>7s} {'TF/s':>5s} | {'dgrad us':>8s} {'TF/s':>5s} | {'wgrad us':>8s} {'TF/s':>5s} | step TF/s") + for dist in args.dists.split(","): + for name in SHAPES: + M, r, step, names = run_case(name, dist, args.iters, gen) + print( + f"{name + '/' + dist:14s} {M:7d} | {r['fwd'][0]:7.1f} {r['fwd'][1]:5.0f} | {r['dgrad'][0]:8.1f} {r['dgrad'][1]:5.0f} |" + f" {r['wgrad'][0]:8.1f} {r['wgrad'][1]:5.0f} | {step:6.0f}" + ) + print("kernels:", names) + + +if __name__ == "__main__": + main() diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_permute_qwen3.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_permute_qwen3.py new file mode 100644 index 000000000..0fe5643e4 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_permute_qwen3.py @@ -0,0 +1,92 @@ +"""moe_permute / moe_unpermute at the Qwen3-30B-A3B EP8 dispatcher call (hidden 2048, 16 local +experts, top-8, pad_multiple 0, topk probs): HIP (TURBO) vs TRITON backends, even and random routing. +""" + +import os + +import torch + +import primus_turbo.pytorch as turbo +from primus_turbo.pytorch.core.backend import BackendType + +H = int(os.environ.get("H", "2048")) +E_LOCAL = int(os.environ.get("E_LOCAL", "16")) +TOPK, RANKS = 8, 8 + + +def timed(fn, iters=30, warmup=5): + for _ in range(warmup): + fn() + torch.cuda.synchronize() + ts = [] + for _ in range(iters): + s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + s.record() + fn() + e.record() + torch.cuda.synchronize() + ts.append(s.elapsed_time(e) * 1e3) + return sorted(ts)[len(ts) // 2] + + +def dispatched_indices(routing): + """Local expert ids per received token (-1 = expert on another rank), as DeepEP returns them.""" + if routing == "even": + # Each token goes to one rank with 8 consecutive experts. + n = 32768 + base = (torch.arange(n, device="cuda") % (E_LOCAL // TOPK)) * TOPK + return (base[:, None] + torch.arange(TOPK, device="cuda")[None, :]).int() + g = torch.Generator(device="cuda").manual_seed(0) + scores = torch.rand(32768 * RANKS, E_LOCAL * RANKS, device="cuda", generator=g) + experts = scores.topk(TOPK, dim=-1).indices + local = (experts // E_LOCAL) == 0 + keep = local.any(dim=-1) + idx = torch.where(local, experts % E_LOCAL, torch.full_like(experts, -1))[keep] + return idx.int() + + +for routing in ("even", "random"): + idx = dispatched_indices(routing) + n = idx.shape[0] + rows = int((idx >= 0).sum()) + x = torch.randn(n, H, device="cuda", dtype=torch.bfloat16) + probs = torch.rand(n, TOPK, device="cuda", dtype=torch.float32) + for backend in (BackendType.TURBO, BackendType.TRITON): + + def permute(x_=x, probs_=probs): + return turbo.ops.moe_permute( + x_, topk_indices=idx, num_local_experts=E_LOCAL, num_topk=TOPK, pad_multiple=0, + num_permuted_tokens=rows, probs=probs_, probs_layout="topk", backend=backend, + ) + + out, row_id_map, *_rest, pprobs = permute() + + def unpermute(y=out): + return turbo.ops.moe_unpermute( + y, row_id_map, restore_shape=x.shape, num_local_experts=E_LOCAL, pad_multiple=0, + backend=backend, + ) + + xg = x.clone().requires_grad_() + pg = probs.clone().requires_grad_() + yg = out.detach().clone().requires_grad_() + gy = torch.randn_like(out) + gp = torch.randn_like(pprobs) + gx = torch.randn_like(x) + + def permute_fb(): + o, _, *_r, pp = permute(xg, pg) + torch.autograd.backward((o, pp), (gy, gp)) + + def unpermute_fb(): + unpermute(yg)[0].backward(gx) + + t_p, t_u = timed(permute), timed(unpermute) + t_pfb, t_ufb = timed(permute_fb), timed(unpermute_fb) + # permute fwd: read n*H, write rows*H; unpermute fwd: read rows*H, write n*H (bf16) + gb = (n + rows) * H * 2 / 1e9 + print( + f"{routing:6s} {backend.name:6s} tokens {n:6d} rows {rows:6d} | permute {t_p:6.1f} us " + f"({gb / t_p * 1e3:4.2f} TB/s) bwd {t_pfb - t_p:6.1f} | unpermute {t_u:6.1f} us " + f"({gb / t_u * 1e3:4.2f} TB/s) bwd {t_ufb - t_u:6.1f}" + ) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_qk_rmsnorm_rope.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_qk_rmsnorm_rope.py new file mode 100644 index 000000000..e7c61f34d --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_qk_rmsnorm_rope.py @@ -0,0 +1,97 @@ +"""Item 5 microbench: Qwen3-30B-A3B q/k RMSNorm + RoPE at seq 4096, MBS 8, NG=4 KV groups, +NPG=8 q heads per group, head_dim 128 (packed qkv [S, B, NG, (NPG+2)*D]). + +unfused: Megatron's current path -- split the packed qkv, TE RMSNorm on q and k (with the + reshape copies the split views force), TE fused RoPE on q and k, v made contiguous. +fused: primus_turbo fused_qkv_rmsnorm_rope (one read of qkv, one write of q/k/v). +Reports fwd and fwd+bwd wall time per call (CUDA events) and fused HBM bandwidth. +""" + +import os + +import torch +import transformer_engine.pytorch as tep +from transformer_engine.pytorch.attention.rope import apply_rotary_pos_emb + +from primus_turbo.pytorch.ops.rope import fused_qkv_rmsnorm_rope + +S, B, NG, NPG, D = (int(v) for v in os.environ.get("SHAPE", "4096,8,4,8,128").split(",")) +EPS = 1e-6 +dev = "cuda" + + +def bench(fn, iters=30): + for _ in range(5): + fn() + s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + torch.cuda.synchronize() + s.record() + for _ in range(iters): + fn() + e.record() + torch.cuda.synchronize() + return s.elapsed_time(e) * 1e3 / iters + + +torch.manual_seed(0) +split = [NPG * D, D, D] +qkv = torch.randn(S, B, NG, sum(split), device=dev, dtype=torch.bfloat16, requires_grad=True) +half = torch.randn(S, D // 2, device=dev) +freqs = torch.cat((half, half), -1).reshape(S, 1, 1, D).contiguous() +q_norm = tep.RMSNorm(D, eps=EPS, params_dtype=torch.bfloat16).to(dev) +k_norm = tep.RMSNorm(D, eps=EPS, params_dtype=torch.bfloat16).to(dev) +with torch.no_grad(): + q_norm.weight.normal_() + k_norm.weight.normal_() +q_gamma = q_norm.weight.detach().clone().requires_grad_() +k_gamma = k_norm.weight.detach().clone().requires_grad_() + + +def unfused(): + q, k, v = torch.split(qkv, split, dim=3) + q = q.reshape(S, B, NG * NPG, D) + k = k.reshape(S, B, NG, D) + v = v.reshape(S, B, NG, D) + q = apply_rotary_pos_emb(q_norm(q), freqs, tensor_format="sbhd", fused=True) + k = apply_rotary_pos_emb(k_norm(k), freqs, tensor_format="sbhd", fused=True) + return q, k, v.contiguous() + + +def fused(): + return fused_qkv_rmsnorm_rope(qkv, q_gamma, k_gamma, freqs, split, EPS) + + +grads = None + + +def fwd_bwd(fn): + def run(): + outs = fn() + torch.autograd.backward(outs, grads) + qkv.grad = None + + return run + + +with torch.no_grad(): + ref = unfused() + out = fused() +grads = tuple(torch.randn_like(t) for t in ref) +diff = [(a.float() - b.float()).abs().max().item() for a, b in zip(out, ref)] + +nbytes_fwd = qkv.numel() * 2 * 2 # read packed qkv, write q/k/v of the same size +nbytes_bwd = qkv.numel() * 2 * 3 # read dq/dk/dv + packed qkv, write dqkv +res = {} +for name, fn in (("unfused", unfused), ("fused", fused)): + with torch.no_grad(): + t_f = bench(fn) + t_fb = bench(fwd_bwd(fn)) + res[name] = (t_f, t_fb) +print(f"shape S={S} B={B} NG={NG} NPG={NPG} D={D}; fused vs unfused max|diff| q/k/v = {diff}") +for name, (t_f, t_fb) in res.items(): + print(f"{name:8s} fwd {t_f:7.1f} us fwd+bwd {t_fb:7.1f} us bwd {t_fb - t_f:7.1f} us") +t_f, t_fb = res["fused"] +print( + f"fused fwd {nbytes_fwd / t_f / 1e6:.2f} TB/s, bwd {nbytes_bwd / (t_fb - t_f) / 1e6:.2f} TB/s; " + f"speedup fwd {res['unfused'][0] / t_f:.2f}x fwd+bwd {res['unfused'][1] / t_fb:.2f}x" +) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_quant_chain.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_quant_chain.py new file mode 100644 index 000000000..1767fe690 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_quant_chain.py @@ -0,0 +1,111 @@ +"""Item 1 microbench: tensorwise FP8 quant chain at Qwen3-30B-A3B (EP8, MBS8) shapes. + +Per-kernel time comes from torch.profiler so amax / scale / quant / transpose are +separated; bandwidth = bytes moved / kernel time. +""" + +import argparse +import collections +import re + +import torch +from torch.profiler import ProfilerActivity, profile + +from primus_turbo.pytorch.kernels.quantization.quantization_impl import ( + quantize_fp8_tensorwise_pad_impl, +) + +E4M3 = torch.float8_e4m3fn +E5M2 = torch.float8_e5m2 + +ROWS_GG = 262144 +TOK = 32768 +# (tag, rows, cols, fp8 dtype, calls per layer-microbatch) +QUANT_CASES = [ + ("gg_fwd_fc1_in", ROWS_GG, 2048, E4M3), + ("gg_fwd_fc2_in", ROWS_GG, 768, E4M3), + ("gg_bwd_fc2_gout", ROWS_GG, 2048, E5M2), + ("gg_bwd_fc1_gout", ROWS_GG, 1536, E5M2), + ("dense_fwd_qkv_in", TOK, 2048, E4M3), + ("dense_fwd_proj_in", TOK, 4096, E4M3), + ("dense_bwd_qkv_gout", TOK, 5120, E5M2), + ("dense_bwd_proj_gout", TOK, 2048, E5M2), +] +TRANSPOSE_CASES = [ + ("fwd_a_t_qkv", TOK, 2048), + ("fwd_a_t_proj", TOK, 4096), + ("fwd_w_t_qkv", 5120, 2048), + ("fwd_w_t_proj", 2048, 4096), + ("bwd_gout_t_qkv", TOK, 5120), + ("bwd_gout_t_proj", TOK, 2048), +] + + +def kernel_times(fn, iters): + for _ in range(3): + fn() + torch.cuda.synchronize() + with profile(activities=[ProfilerActivity.CUDA]) as prof: + for _ in range(iters): + fn() + torch.cuda.synchronize() + agg = collections.Counter() + for e in prof.events(): + if e.device_type == torch.autograd.DeviceType.CUDA: + name = re.sub(r"<.*", "", e.name).replace("void ", "").replace("primus_turbo::", "") + agg[name] += e.device_time + return {k: v / iters for k, v in agg.items()} # us per call + + +def short(name): + for k in ("amax_partial", "amax_scale", "pad_row", "transpose_2d", "elementwise", "reduce", "copy"): + if k in name: + return k + return name[:40] + + +def main(): + p = argparse.ArgumentParser() + p.add_argument("--iters", type=int, default=50) + args = p.parse_args() + torch.manual_seed(0) + dev = "cuda" + + print("== HBM reference (torch)") + x = torch.randn(ROWS_GG, 2048, device=dev, dtype=torch.bfloat16) + nbytes = x.numel() * 2 + t = kernel_times(lambda: x.clone(), args.iters) + tt = sum(t.values()) + print(f" clone 1.07GB bf16 : {tt:8.1f} us {2 * nbytes / tt / 1e6:6.2f} TB/s (r+w)") + t = kernel_times(lambda: x.abs().amax(), args.iters) + rd = max(t.values()) + print(f" amax (largest kernel): {rd:8.1f} us {nbytes / rd / 1e6:6.2f} TB/s (r)") + del x + + print("== quantize_fp8_tensorwise (amax partial + scale + quant)") + print(f" {'case':22s} {'shape':>14s} {'amax us':>8s} {'TB/s':>5s} {'quant us':>8s} {'TB/s':>5s} {'scale us':>8s} {'total us':>8s}") + for tag, m, k, dt in QUANT_CASES: + x = torch.randn(m, k, device=dev, dtype=torch.bfloat16) + t = kernel_times(lambda: quantize_fp8_tensorwise_pad_impl(x, dt, k_align=128), args.iters) + d = collections.Counter() + for n, v in t.items(): + d[short(n)] += v + n = m * k + amax_bw = n * 2 / d["amax_partial"] / 1e6 if d["amax_partial"] else 0 + q_bw = n * 3 / d["pad_row"] / 1e6 if d["pad_row"] else 0 + print( + f" {tag:22s} {m:7d}x{k:<6d} {d['amax_partial']:8.1f} {amax_bw:5.2f} {d['pad_row']:8.1f} {q_bw:5.2f}" + f" {d['amax_scale']:8.1f} {sum(t.values()):8.1f}" + ) + del x + + print("== transpose_2d") + for tag, m, k in TRANSPOSE_CASES: + x = torch.randint(0, 255, (m, k), device=dev, dtype=torch.uint8).view(E4M3) + t = kernel_times(lambda: torch.ops.primus_turbo_cpp_extension.transpose_2d(x, 0, 1), args.iters) + tt = sum(t.values()) + print(f" {tag:22s} {m:7d}x{k:<6d} {tt:8.1f} us {2 * m * k / tt / 1e6:5.2f} TB/s") + + +if __name__ == "__main__": + main() diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_aiter_fwd.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_aiter_fwd.py new file mode 100644 index 000000000..6a829b1a7 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_aiter_fwd.py @@ -0,0 +1,42 @@ +"""Why does Turbo's aiter fwd take 2x TE's time with the same fmha_fwd_hd128_bf16_causal kernel?""" + +import torch +from aiter.ops.mha import _flash_attn_forward + +S, B, HQ, HKV, D = 4096, 8, 32, 4, 128 + + +def timed(fn, iters=20): + for _ in range(5): + fn() + torch.cuda.synchronize() + ts = [] + for _ in range(iters): + s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + s.record() + fn() + e.record() + torch.cuda.synchronize() + ts.append(s.elapsed_time(e) * 1e3) + return sorted(ts)[len(ts) // 2] + + +torch.manual_seed(0) +sbhd = [torch.randn(S, B, h, D, device="cuda", dtype=torch.bfloat16) for h in (HQ, HKV, HKV)] +layouts = { + "sbhd-view": [t.permute(1, 0, 2, 3) for t in sbhd], + "bshd-contig": [t.permute(1, 0, 2, 3).contiguous() for t in sbhd], +} +q, k, v = layouts["sbhd-view"] +out = torch.empty(S, B, HQ, D, device="cuda", dtype=torch.bfloat16).permute(1, 0, 2, 3) +k32, v32 = (t.repeat_interleave(HQ // HKV, dim=2) for t in (k, v)) +for causal in (True, False): + for name, (kk, vv) in (("gqa 32/4", (k, v)), ("mha 32/32", (k32, v32))): + + def run(): + _flash_attn_forward( + q, kk, vv, 0.0, D**-0.5, causal, -1, -1, 0, None, None, None, None, None, + True, False, out=out, + ) + + print(f"causal={causal} {name}: {timed(run):8.1f} us") diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_permute_probs.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_permute_probs.py new file mode 100644 index 000000000..336685601 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_permute_probs.py @@ -0,0 +1,55 @@ +"""Does moe_permute write every permuted_probs row for top-k indices with -1 slots + top-k probs?""" + +import torch + +from primus_turbo.pytorch.core.backend import BackendType +from primus_turbo.pytorch.ops import moe_permute + +num_tokens, E, K, H = 333, 16, 8, 256 +g = torch.Generator(device="cuda").manual_seed(31) +scores = torch.rand((num_tokens, 4 * E), generator=g, device="cuda") +ge = scores.topk(K, dim=-1).indices +local = ge < E +topk = torch.where(local, ge, torch.full_like(ge, -1)).int() +import os # noqa: E402 + +if os.environ.get("PREFIX") == "1": + order = torch.argsort((~local.any(dim=1)).int(), stable=True) # routed tokens first, as DeepEP + topk, ge, local = topk[order], ge[order], local[order] +rows = int(local.sum()) +x = torch.randn((num_tokens, H), dtype=torch.bfloat16, device="cuda") +probs = torch.rand((num_tokens, K), dtype=torch.float32, device="cuda") + +# Reference: expert-major, token order inside an expert. +ref_tok, ref_prob = [], [] +for e in range(E): + t, k = (topk == e).nonzero(as_tuple=True) + ref_tok.append(t) + ref_prob.append(probs[t, k]) +ref_tok, ref_prob = torch.cat(ref_tok), torch.cat(ref_prob) + +for backend in (BackendType.TURBO, BackendType.TRITON): + for fill, idx_dtype in ((7.0, torch.int32), (7.0, torch.int64), (-3.0, torch.int64)): + torch.cuda.empty_cache() + junk = torch.full((rows * H,), fill, device="cuda") # noqa: F841 (poison the allocator) + del junk + out, _, tpe, _, _, pp = moe_permute( + x, topk_indices=topk.to(idx_dtype), num_local_experts=E, num_topk=K, num_permuted_tokens=rows, + probs=probs, probs_layout="topk", backend=backend, + ) + print(idx_dtype, end=" ") + bad_tok = (out != x[ref_tok]).any(dim=1).nonzero().flatten() + bad_prob = (pp != ref_prob).nonzero().flatten() + print( + f"{backend.name:6s} fill={fill} rows={rows} tokens bad={bad_tok.numel()} probs bad={bad_prob.numel()} " + f"first bad prob rows={bad_prob[:6].tolist()} vals={pp[bad_prob[:3]].tolist()}" + ) + if backend is BackendType.TURBO and idx_dtype is torch.int64 and fill > 0: + ref_cnt = torch.stack([(topk == e).sum() for e in range(E)]) + print(" ref tokens_per_expert", ref_cnt.tolist()) + print(" got tokens_per_expert", tpe.tolist()) + seg_end = ref_cnt.cumsum(0).tolist() + print(" segment ends", seg_end) + print(" bad rows", bad_prob.tolist()[:40]) + multi = (local.sum(dim=1) > 1).sum().item() + print(" tokens with >1 local expert", multi, "with 0 local", (local.sum(dim=1) == 0).sum().item()) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128.py new file mode 100644 index 000000000..9255cb92d --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128.py @@ -0,0 +1,39 @@ +import sys + +import torch + +sys.path.insert(0, "/home/xiaompen/turbo-opt") +from tests.pytorch.ops.test_qk_rmsnorm_rope import _inputs, _reference # noqa: E402 + +from primus_turbo.pytorch.ops.rope import fused_qkv_rmsnorm_rope # noqa: E402 + +S, B, NG, NPG, D = (int(v) for v in sys.argv[1:6]) +qkv, qg, kg, freqs, split = _inputs(S=S, B=B, NG=NG, NPG=NPG, seed=321, D=D) +qkv.requires_grad_() +ref_in = qkv.detach().clone().requires_grad_() +out = fused_qkv_rmsnorm_rope(qkv, qg, kg, freqs, split, 1e-5) +ref = _reference(ref_in, qg, kg, freqs, split, 1e-5) +grads = tuple(torch.randn_like(t) for t in out) +torch.autograd.backward(out, grads) +torch.autograd.backward(ref, grads) +g, r = qkv.grad.float(), ref_in.grad.float() +bad = ~torch.isclose(g, r, rtol=2e-2, atol=2e-2) +print("bad", bad.sum().item(), "of", bad.numel(), "nan", g.isnan().sum().item()) +if bad.any(): + idx = bad.nonzero() + print("seq range", idx[:, 0].min().item(), idx[:, 0].max().item()) + print("batch values", idx[:, 1].unique().tolist()) + print("group values", idx[:, 2].unique().tolist()) + col = idx[:, 3] + print("heads in packed group", (col // D).unique().tolist()) + print("elem-in-head values", (col % D).unique().tolist()[:80]) + s0, b0, g0, c0 = idx[0].tolist() + print("first bad", (s0, b0, g0, c0), "got", g[s0, b0, g0, c0 - 4 : c0 + 4].tolist()) + print(" ref", r[s0, b0, g0, c0 - 4 : c0 + 4].tolist()) + first = g.clone() + qkv.grad = None + torch.autograd.backward(fused_qkv_rmsnorm_rope(qkv, qg, kg, freqs, split, 1e-5), grads) + g2 = qkv.grad.float() + print("rerun identical:", torch.equal(first.nan_to_num(7.0), g2.nan_to_num(7.0))) + bad2 = ~torch.isclose(g2, r, rtol=2e-2, atol=2e-2) + print("bad sets identical:", torch.equal(bad, bad2)) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128_isa.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128_isa.py new file mode 100644 index 000000000..0dd1a90db --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128_isa.py @@ -0,0 +1,22 @@ +"""Compile only the qk_rmsnorm_rope backward kernel (for an ISA dump).""" + +import sys + +import torch + +sys.path.insert(0, "/home/xiaompen/turbo-opt") +import primus_turbo.flydsl.rope.qk_rmsnorm_rope_kernel as K # noqa: E402 +from tests.pytorch.ops.test_qk_rmsnorm_rope import _inputs # noqa: E402 + +S, B, NG, NPG, D = (int(v) for v in sys.argv[1:6]) +qkv, qg, kg, freqs, split = _inputs(S=S, B=B, NG=NG, NPG=NPG, seed=321, D=D) +q, k, v, qr, kr = K.flydsl_qkv_rmsnorm_rope_forward(qkv, qg, kg, freqs, split, 1e-5) +K._make_fold_kernel = None # stop before the fold kernel compiles over the dump +try: + K.flydsl_qkv_rmsnorm_rope_backward( + torch.randn_like(q), torch.randn_like(k), torch.randn_like(v), qkv, qg, kg, freqs, qr, kr, split + ) +except TypeError: + pass +torch.cuda.synchronize() +print("bwd compiled") diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128b.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128b.py new file mode 100644 index 000000000..d8adf72b3 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/dbg_rope128b.py @@ -0,0 +1,30 @@ +import sys + +import torch + +sys.path.insert(0, "/home/xiaompen/turbo-opt") +import primus_turbo.flydsl.rope.qk_rmsnorm_rope_kernel as K # noqa: E402 +from tests.pytorch.ops.test_qk_rmsnorm_rope import _inputs # noqa: E402 + +S, B, NG, NPG, D = (int(v) for v in sys.argv[1:6]) +qkv, qg, kg, freqs, split = _inputs(S=S, B=B, NG=NG, NPG=NPG, seed=321, D=D) +q, k, v, qr, kr = K.flydsl_qkv_rmsnorm_rope_forward(qkv, qg, kg, freqs, split, 1e-5) + +orig_empty_like = torch.empty_like +K.torch.empty_like = lambda t, **kw: torch.full_like(t, 123.0) +try: + dqkv, _, _ = K.flydsl_qkv_rmsnorm_rope_backward( + torch.zeros_like(q), torch.zeros_like(k), torch.zeros_like(v), qkv, qg, kg, freqs, qr, kr, split + ) +finally: + K.torch.empty_like = orig_empty_like +torch.cuda.synchronize() +g = dqkv.float() +bad = g != 0 +print("nonzero", bad.sum().item(), "==123:", (g == 123).sum().item()) +if bad.any(): + idx = bad.nonzero() + col = idx[:, 3] + print("heads", (col // D).unique().tolist(), "elems", (col % D).unique().tolist()[:40]) + vals = g[bad] + print("sample values", vals[:8].tolist()) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/debug_fp8_worst.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/debug_fp8_worst.py new file mode 100644 index 000000000..7a9f146e3 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/debug_fp8_worst.py @@ -0,0 +1,98 @@ +"""2-GPU debug: worst-tokens + pad 16 FP8 permute mismatch (mirrors test_fp8_permute).""" + +import os + +import torch +import torch.distributed as dist + +import primus_turbo.pytorch as turbo +from primus_turbo.pytorch.core.low_precision import ( + ScalingGranularity, + float8_e4m3, + float8_e5m2, +) +from primus_turbo.pytorch.core.quantized_tensor import QuantizedTensor + +rank = int(os.environ["RANK"]) +torch.cuda.set_device(rank) +dist.init_process_group("nccl") +T, H, E, K = 4096, 4096, 256, 8 + + +class Stub(torch.autograd.Function): + @staticmethod + def forward(ctx, x, out, record): + ctx.record = record + record["input"] = x + return out.clone() + + @staticmethod + def backward(ctx, grad): + ctx.record["grad"] = grad + return torch.zeros(grad.shape, dtype=torch.bfloat16, device=grad.device), None, None + + +def case(pad, worst, rep): + d = turbo.modules.DeepEPTokenDispatcher( + E, K, dist.group.WORLD, pad_multiple=pad, deepep_num_worst_tokens=worst, + deepep_use_cuda_num_tokens_per_expert=True, + ) + gen = torch.Generator(device="cuda").manual_seed(1234 + rank) + h0 = torch.randn((T, H), dtype=torch.bfloat16, device="cuda", generator=gen) + probs = torch.rand((T, E), dtype=torch.float32, device="cuda", generator=gen) + gout = torch.randn((T, H), dtype=torch.bfloat16, device="cuda", generator=gen) + + def run(q): + rec = {} + h = h0.clone().requires_grad_(True) + x, tp = d._pre_dispatch(h, probs) + x, dp = d._exec_dispatch(x, tp) + rec["recv"] = x.detach().clone() + rec["idx"] = d.dispatched_indices.clone() + p, tpe, _ = d._post_dispatch(x, dp, quantize_dtype=float8_e4m3 if q else None) + rec["rmap"] = d.row_id_map.clone() + out = torch.ones(p.shape, dtype=torch.bfloat16, device="cuda") + y = Stub.apply(p, out, rec) + y = d._pre_combine(y, grad_quantize_dtype=float8_e5m2 if q else None) + d._post_combine(d._exec_combine(y)).backward(gout) + return rec, tpe + + ref, tpe = run(False) + got, tpe2 = run(True) + want = QuantizedTensor.quantize( + ref["input"].detach(), float8_e4m3, ScalingGranularity.TENSORWISE, axis=-1, group_lens=tpe, pad_align_last=128 + ) + q = got["input"] + diff = q.qdata.view(torch.uint8) != want.qdata.view(torch.uint8) + rows = diff.any(1).nonzero().flatten() + recv_same = torch.equal(ref["recv"], got["recv"]) + idx = ref["idx"] + nloc = idx.shape[1] and d.num_local_experts + real = torch.stack([(idx == e).sum() for e in range(nloc)]).to(tpe.device) + offs = torch.cat([tpe.new_zeros(1), tpe.cumsum(0)]) + pad_mask = torch.zeros(q.shape[0], dtype=torch.bool, device="cuda") + for e in range(nloc): + pad_mask[int(offs[e] + real[e]) : int(offs[e + 1])] = True + ref_rows_nonzero = (ref["input"].detach() != 0).any(1) + got_rows_nonzero = (q.qdata.view(torch.uint8) != 0).any(1) + print( + f"rank{rank} pad={pad} worst={worst} rep={rep}: pad_rows={int(pad_mask.sum())} " + f"ref_nonzero_pad_rows={int((ref_rows_nonzero & pad_mask).sum())} " + f"fp8_nonzero_pad_rows={int((got_rows_nonzero & pad_mask).sum())} " + f"diff_rows_in_pad={int(pad_mask[rows].sum()) if rows.numel() else 0}/{rows.numel()} " + f"idx_tail_valid={int((idx[ref['rmap'].shape[0] - pad:] >= 0).sum()) if worst else -1}", + flush=True, + ) + print( + f"rank{rank} pad={pad} worst={worst} rep={rep}: tpe_eq={torch.equal(tpe, tpe2)} recv_eq={recv_same} " + f"scale got={q.scale_inv.item():.6e} want={want.scale_inv.item():.6e} ndiff_rows={rows.numel()} " + f"first={rows[:6].tolist()} nperm={q.shape[0]}", + flush=True, + ) + + +for rep in range(3): + for pad in (0, 16): + for worst in (0, T * 8): + case(pad, worst, rep) +dist.destroy_process_group() diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/amend_msg.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/amend_msg.py new file mode 100644 index 000000000..606d2dd45 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/amend_msg.py @@ -0,0 +1,49 @@ +"""Insert the end-to-end section into a commit message (before the Tests: paragraph). + +Usage: git log -1 --format=%B | python amend_msg.py ROW_KEY > msg.txt +ROW_KEY marks the row of this change and selects the before/after pair. +""" + +import sys + +ROWS = [ + ("cfg", "submitted FP8 tensorwise config (Turbo e2d9f1d7)", 8894.0, 670.6, 29474, 277, "11.33393"), + ("act", "+ use_turbo_fused_act_with_probs (Primus flag)", 7915.6, 753.5, 33117, 223, "11.33435"), + ("base", "Turbo 1103b2df, base of these changes", 7925.4, 752.5, 33077, 222, "11.33427"), + ("quant", "+ flat tensorwise FP8 quant kernel", 7870.9, 757.7, 33305, 222, "11.33415"), + ("layout", "+ per-token get_dispatch_layout", 7774.2, 767.2, 33720, 222, "11.33431"), + ("permute", "+ HIP permute in DeepEPTokenDispatcher", 7643.5, 780.3, 34296, 222, "11.33423"), + ("cu160", "+ turbo_deepep_num_cu 160 (Primus flag)", 7353.4, 811.1, 35649, 223, "11.33456"), + ("qknorm", "+ fused qk RMSNorm + RoPE, head_dim 128", 6987.1, 853.6, 37518, 224, "11.33373"), +] +PREV = {"quant": "base", "layout": "quant", "permute": "layout", "qknorm": "cu160"} + +key = sys.argv[1] +by_key = {r[0]: r for r in ROWS} +cur, prev = by_key[key], by_key[PREV[key]] + +lines = [ + "End-to-end: Qwen3-30B-A3B FP8 tensorwise pretrain on 8x MI355X (EP8, MBS 8,", + "GBS 512, seq 4096, even routing), mean of iterations 11-20 of 20-iteration", + "runs; each row adds one change on top of the row above:", + f" {'':49s} {'ms/iter':>7s} {'TFLOP/s':>7s} {'tok/s/GPU':>9s} {'mem':>6s} {'loss@20':>8s}", +] +for k, name, ms, tf, tok, mem, loss in ROWS: + mark = "*" if k == key else " " + lines.append(f"{mark} {name:49s} {ms:7.1f} {tf:7.1f} {tok:9,d} {mem:4d}GB {loss:>8s}") +d_ms = cur[2] - prev[2] +lines += [ + f"This change (*): {prev[2]:.1f} -> {cur[2]:.1f} ms/iter ({100 * d_ms / prev[2]:+.2f}%),", + f"{prev[4]:,d} -> {cur[4]:,d} tokens/s/GPU ({100 * (cur[4] / prev[4] - 1):+.2f}%).", + "Loss differences are within run-to-run noise (about 4e-4 at iteration 20).", +] +section = "\n".join(lines) + "\n" + +msg = sys.stdin.read().rstrip("\n") + "\n" +idx = msg.find("\nTests:") +if idx < 0: + idx = msg.find("\nRequires ") +if idx < 0: + sys.stdout.write(msg + "\n" + section) +else: + sys.stdout.write(msg[: idx + 1] + section + msg[idx:]) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_ab.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_ab.sh new file mode 100644 index 000000000..e5a977708 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_ab.sh @@ -0,0 +1,35 @@ +#!/bin/bash +# Sequential 20-iteration Qwen3-30B-A3B FP8 tensorwise runs (even routing, MBS 8, GBS 512). +# Usage: e2e_ab.sh NAME [NAME ...] (names below); logs: /tmp/q3_e2e_.log +ulimit -c 0 +E2E=/home/xiaompen/turbo-opt/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e +CFG=examples/megatron/configs/MI355X/qwen3_30B_A3B-FP8-pretrain.yaml +BASE_ARGS=(--train_iters 20 --use_turbo_fused_act_with_probs True) + +run() { + local name=$1 repo=$2 + shift 2 + echo "[e2e] $name start $(date +%T)" + (cd "$repo" && bash runner/primus-cli direct "$@" -- train pretrain --config "$CFG" "${BASE_ARGS[@]}" "${EXTRA[@]}") \ + > "/tmp/q3_e2e_${name}.log" 2>&1 + echo "[e2e] $name exit=$? $(date +%T)" +} + +for name in "$@"; do + EXTRA=() + case $name in + r0_base) run "$name" /home/xiaompen/Primus-qwen3 ;; + rb_base | b1_quant | b2_layout) run "$name" /home/xiaompen/Primus-qwen3 --env "$E2E/env_turbo_base.sh" ;; + r1_turbo) run "$name" /home/xiaompen/Primus-qwen3 --env "$E2E/env_turbo_opt.sh" ;; + r2_cu160) + EXTRA=(--turbo_deepep_num_cu 160) + run "$name" /home/xiaompen/Primus-qwen3 --env "$E2E/env_turbo_opt.sh" + ;; + r3_qknorm) + EXTRA=(--turbo_deepep_num_cu 160) + run "$name" /home/xiaompen/Primus-qwen3-qknorm --env "$E2E/env_turbo_opt_qknorm.sh" + ;; + *) echo "unknown run $name" ;; + esac +done +echo "[e2e] done" diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_chain.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_chain.sh new file mode 100644 index 000000000..9fa5afdf5 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_chain.sh @@ -0,0 +1,20 @@ +#!/bin/bash +# r3_qknorm, then per-PR e2e attribution on top of turbo-base (1103b2df): base -> +quant -> +layout. +# turbo-opt (r1_turbo) is +layout plus the dispatcher permute change. +ulimit -c 0 +E2E=/home/xiaompen/turbo-opt/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e +TB=/home/xiaompen/turbo-base +G="git -c safe.directory=* -C $TB" + +build() { + (cd $TB && GPU_ARCHS=gfx950 MAX_JOBS=64 python setup.py build_ext --inplace) > "/tmp/turbo_base_incr_$1.log" 2>&1 + local rc=$? + echo "[chain] build $1 exit=$rc $(date +%T)" + return $rc +} + +bash $E2E/e2e_ab.sh r3_qknorm rb_base +$G cherry-pick -n 1194eaa0 && build quant && bash $E2E/e2e_ab.sh b1_quant +$G cherry-pick -n 11b6d998 && build layout && bash $E2E/e2e_ab.sh b2_layout +chown -R 11768:11768 $TB /home/xiaompen/Primus/third_party/Turbo/.git/worktrees/turbo-base +echo "[chain] done $(date +%T)" diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue2.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue2.sh new file mode 100644 index 000000000..3a7bc2b01 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue2.sh @@ -0,0 +1,9 @@ +#!/bin/bash +# After the attribution chain: MegaMoE (mxfp8 / bf16 experts) and sync-free stage 2 on the R3 stack. +ulimit -c 0 +E2E=/home/xiaompen/turbo-opt/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e +while ! grep -q "^\[chain\] done" /tmp/q3_e2e_chain.log; do sleep 15; done +bash $E2E/e2e_run.sh m1_mega_mxfp8 --use_turbo_mega_moe True --turbo_mega_moe_precision mxfp8 +bash $E2E/e2e_run.sh m2_mega_bf16 --use_turbo_mega_moe True --turbo_mega_moe_precision bf16 +bash $E2E/e2e_run.sh s2_syncfree2 --turbo_sync_free_moe_stage 2 +echo "[queue2] done $(date +%T)" diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue3.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue3.sh new file mode 100644 index 000000000..472e584a6 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue3.sh @@ -0,0 +1,8 @@ +#!/bin/bash +# After queue2: one profiled step (rank 0, iteration 10) of the R3 stack. +ulimit -c 0 +E2E=/home/xiaompen/turbo-opt/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e +while ! grep -q "^\[queue2\] done" /tmp/q3_e2e_queue2.log; do sleep 15; done +bash $E2E/e2e_run.sh p3_profile --profile True --use_pytorch_profiler True \ + --profile_step_start 10 --profile_step_end 11 --train_iters 12 +echo "[queue3] done $(date +%T)" diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue4.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue4.sh new file mode 100644 index 000000000..2b3589175 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue4.sh @@ -0,0 +1,8 @@ +#!/bin/bash +# After queue3: DeepEP CU sweep on the R3 stack. +ulimit -c 0 +E2E=/home/xiaompen/turbo-opt/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e +while ! grep -q "^\[queue3\] done" /tmp/q3_e2e_queue3.log; do sleep 15; done +bash $E2E/e2e_run.sh c192 --turbo_deepep_num_cu 192 +bash $E2E/e2e_run.sh c256 --turbo_deepep_num_cu 256 +echo "[queue4] done $(date +%T)" diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue5.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue5.sh new file mode 100644 index 000000000..51a8ad85a --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_queue5.sh @@ -0,0 +1,8 @@ +#!/bin/bash +# After queue4: fused grouped MLP (fc1 + GLU + fc2 in Turbo), then also without explicit padding. +ulimit -c 0 +E2E=/home/xiaompen/turbo-opt/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e +while ! grep -q "^\[queue4\] done" /tmp/q3_e2e_queue4.log; do sleep 15; done +bash $E2E/e2e_run.sh f1_fused_mlp --turbo_fused_grouped_gemm True +bash $E2E/e2e_run.sh f2_fused_mlp_nopad --turbo_fused_grouped_gemm True --turbo_grouped_gemm_without_padding True +echo "[queue5] done $(date +%T)" diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run.sh new file mode 100644 index 000000000..e298496c5 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run.sh @@ -0,0 +1,13 @@ +#!/bin/bash +# One 20-iteration Qwen3-30B-A3B FP8 run on top of the R3 stack (turbo-opt, fused qk-norm, 160 DeepEP CUs). +# Usage: e2e_run.sh NAME [extra primus args ...] log: /tmp/q3_e2e_.log +ulimit -c 0 +E2E=/home/xiaompen/turbo-opt/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e +CFG=examples/megatron/configs/MI355X/qwen3_30B_A3B-FP8-pretrain.yaml +name=$1 +shift +echo "[e2e] $name start $(date +%T)" +(cd /home/xiaompen/Primus-qwen3-qknorm && bash runner/primus-cli direct --env "$E2E/env_turbo_opt_qknorm.sh" -- \ + train pretrain --config "$CFG" --train_iters 20 --use_turbo_fused_act_with_probs True \ + --turbo_deepep_num_cu 160 "$@") > "/tmp/q3_e2e_${name}.log" 2>&1 +echo "[e2e] $name exit=$? $(date +%T)" diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_base.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_base.sh new file mode 100644 index 000000000..a6fd80103 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_base.sh @@ -0,0 +1,2 @@ +export PYTHONPATH=/home/xiaompen/turbo-base:${PYTHONPATH} +export FLYDSL_RUNTIME_CACHE_DIR=/tmp/fly_cache_base diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt.sh new file mode 100644 index 000000000..16cd9f29a --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt.sh @@ -0,0 +1,2 @@ +export PYTHONPATH=/home/xiaompen/turbo-opt:${PYTHONPATH} +export FLYDSL_RUNTIME_CACHE_DIR=/tmp/fly_cache_opt diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt_qknorm.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt_qknorm.sh new file mode 100644 index 000000000..83e17d4e7 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt_qknorm.sh @@ -0,0 +1,3 @@ +export PYTHONPATH=/home/xiaompen/turbo-opt:/home/xiaompen/Primus-qwen3/third_party/Megatron-LM:${PYTHONPATH} +export FLYDSL_RUNTIME_CACHE_DIR=/tmp/fly_cache_opt +export PRIMUS_FUSED_QK_RMSNORM_ROPE=1 diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/parse_e2e.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/parse_e2e.py new file mode 100644 index 000000000..0ff4fc551 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/parse_e2e.py @@ -0,0 +1,40 @@ +"""Summarize /tmp/q3_e2e_.log runs over iterations 11-20. + +tokens/s/GPU = GBS 512 x seq 4096 / 8 GPUs / mean iteration time; memory is the max-rank ROCm usage. +""" + +import re +import statistics +import sys + +TOKENS_PER_GPU_ITER = 512 * 4096 // 8 +ITER = re.compile(r"iteration\s+(\d+)/\s*\d+.*?elapsed time per iteration \(ms\): ([\d.]+)") +TFLOPS = re.compile(r"per GPU \(TFLOP/s/GPU\): ([\d.]+)") +LOSS = re.compile(r"lm loss: ([\d.E+-]+)") +MEM = re.compile(r"rocm max mem usage/usage_ratio: ([\d.]+)GB") + +print("| run | ms/iter | sd | TFLOP/s/GPU | tokens/s/GPU | max mem GB | loss@20 |") +print("|---|---|---|---|---|---|---|") +for name in sys.argv[1:]: + rows = [] + for line in open(f"/tmp/q3_e2e_{name}.log", errors="ignore"): + m = ITER.search(line) + if m and int(m.group(1)) > 10: + t, loss, mem = TFLOPS.search(line), LOSS.search(line), MEM.search(line) + rows.append( + ( + float(m.group(2)), + float(t.group(1)) if t else 0.0, + loss.group(1) if loss else "-", + float(mem.group(1)) if mem else 0.0, + ) + ) + if not rows: + print(f"| {name} | no iterations |") + continue + ms = statistics.mean(r[0] for r in rows) + print( + f"| {name} | {ms:.1f} | {statistics.pstdev(r[0] for r in rows):.1f} | " + f"{statistics.mean(r[1] for r in rows):.1f} | {TOKENS_PER_GPU_ITER / ms * 1e3:,.0f} | " + f"{max(r[3] for r in rows):.1f} | {float(rows[-1][2]):.5f} |" + ) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-moe-permute-padding-rows.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-moe-permute-padding-rows.md new file mode 100644 index 000000000..84c80ebb2 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-moe-permute-padding-rows.md @@ -0,0 +1,54 @@ +# Description + +The TURBO (HIP) `moe_permute` leaves each expert's padding rows uninitialized when `pad_multiple > 0` and the dispatch buffer ends in unrouted rows. The padding rows then hold whatever `torch.empty` returned instead of zeros. + +**Root cause.** `permute_preprocessing` writes the `pad_multiple` padding rows of `row_id_map` **after all `max_num_dispatched_tokens` token rows**. `permute_kernel`, however, looked up padding token `i` at row `num_dispatched_tokens + i`, where `num_dispatched_tokens` is the *routed* count. The two positions coincide only when every buffer row is routed. With an unrouted tail, the kernel read the empty rows of unrouted tokens (`n_routed = 0`), so no padding row was ever written. + +**Who is affected.** Anyone combining `pad_multiple > 0` with a worst-case receive buffer: +- `DeepEPTokenDispatcher` with `deepep_num_worst_tokens > 0` and `pad_multiple > 0`. +- In Primus, sync-free MoE stage 2 and up (which set `num_worst_tokens`) together with padding for quantization. + +The grouped GEMM consumes the garbage rows. They enter both expert weight gradients and, under tensorwise FP8, the amax of the permuted activations, so they can change the scale. The gradient permute in `moe_unpermute`'s backward runs the same kernel and had the same bug. Paths where the routed count equals the buffer size (the default dispatch path, DeepEP sync-free stage 1) are not affected. Before this fix they behaved correctly, and they behave identically after it. + +**Repro** (op level, before the fix): a 32768-row buffer with 4000 routed tokens (hidden 4096, 32 local experts, top-2, `pad_multiple=16`) had about 220 of its 224 padding rows nonzero. After the fix, all are zero, and the FP8 and bf16 permute outputs are bit-identical. + +**Performance.** The fix only changes which `row_id_map` row a padding token reads. `permute_kernel` time is unchanged at the Qwen3-30B-A3B EP8 shape (32768 tokens, hidden 2048, 16 local experts, top-8, `pad_multiple=16`): bf16 232.3 µs (232 before), FP8 97.3 µs (95 before). + +## Type of change + +- [ ] Documentation change (change only to the documentation, either a fix or a new content) +- [x] Bug fix (non-breaking change which fixes an issue) +- [ ] New feature (non-breaking change which adds functionality) +- [ ] Breaking change (fix or feature that would cause existing functionality to not work as expected) +- [ ] Infra/Build change +- [ ] Code refactoring + +## Changes + +- `csrc/kernels/moe_permute/moe_permute.cu`: + - `permute_kernel` takes `max_num_dispatched_tokens` (= `row_id_map.size(0) - pad_multiple`). + - It reads padding token `i` from row `max_num_dispatched_tokens + i`; the row offset is now computed in `int64`. + - When every buffer row is routed, the row index is unchanged. +- `tests/pytorch/ops/test_moe_permute.py`: new `test_moe_permute_padding_rows_zero_with_unrouted_tail` (`pad_multiple` 8 and 16). + - It routes only the first 300 of 1024 tokens. + - It hands the output allocation NaN-filled memory, so an unwritten row cannot pass as zero. + - It checks that the padding rows are zero and the real rows match the source tokens. + +Not changed: TURBO `unpermute` still leaves output rows past the routed count unwritten. DeepEP combine reads only the received rows, so those rows are never consumed. + +## Tests + +- The new test **fails on main (1103b2df) and passes with this fix** (2/2). +- `tests/pytorch/ops/test_moe_permute.py`: 144 passed, 19 skipped (MI355X / gfx950). +- `tests/pytorch/modules/test_token_dispatcher.py` (2 GPUs): 25 passed. This includes a worst-token + `pad_multiple=16` case that failed intermittently before the fix, and passed 5/5 repeated runs with it. + +The last two suites ran on a build that also contains the follow-up branch `perf/moe/fp8-permute-tensorwise`, which builds on this fix. + +# Checklist: + +- [x] The functionality is complete +- [x] I have commented my code, particularly in hard-to-understand areas +- [ ] I have made corresponding changes to the documentation +- [ ] My changes generate no new warnings +- [x] I have added tests that prove my fix is effective or that my feature works +- [x] New and existing unit tests pass locally with my changes diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-fp8-permute-tensorwise.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-fp8-permute-tensorwise.md new file mode 100644 index 000000000..15378f153 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-fp8-permute-tensorwise.md @@ -0,0 +1,82 @@ +# Description + +> Stacked on # (`perf/moe/permute-default-hip`): the dispatcher must run the TURBO permute. I will rebase onto main once that merges; until then, review the last two commits (`perf(moe): emit tensorwise FP8 ...` and the format fix). + +Under FP8 tensorwise, the fused grouped MLP (`grouped_mlp_fp8`) quantizes: +- its input right after `DeepEPTokenDispatcher` permutes it, and +- its output gradient right after `moe_unpermute`'s backward permutes that. + +Each time, permute writes bf16 rows, an amax pass re-reads them, and the quant kernel re-reads them and writes FP8. All three run at the permuted size (top-k × the received tokens). + +A tensorwise scale commutes with a row copy. So the quantize can run on the received tokens, and the permute can move FP8 bytes instead: +- A new `permute_routed_amax` op computes per-block abs-max partials over exactly the tokens the permute copies. The unrouted tail of a worst-case buffer never reaches the scale. +- `quantize_fp8_tensorwise_pad_impl` takes those partials, so there is no separate amax pass, and pads the columns to 128. +- The TURBO permute copies the FP8 rows; padding rows stay zero. + +The result is **bit-identical** to quantizing the bf16 permute output. + +**New opt-in API** (off by default; TURBO permute backend only): +- `moe_permute(..., quantize_dtype=)` returns the permuted tokens as a grouped tensorwise `QuantizedTensor` (`group_lens = tokens_per_expert`). +- `moe_unpermute(..., grad_quantize_dtype=)` makes its backward return the gradient as a tensorwise `QuantizedTensor`. +- `DeepEPTokenDispatcher._post_dispatch` / `_pre_combine` pass both options through. +- `grouped_mlp_fp8` backward accepts a `QuantizedTensor` `grad_out`, as its forward already accepted a `QuantizedTensor` input. + +Primus turns it on with `turbo_fp8_permute` (companion PR from `AMD-AGI/Primus` branch `perf/megatron/turbo-fp8-permute`). + +**Microbenchmark.** Qwen3-30B-A3B EP8 dispatcher call (32768 received tokens, hidden 2048, 16 local experts, top-8, `pad_multiple=16`, about 262k permuted rows), kernel time per call: + +| | before | after | +|---|---|---| +| forward | bf16 permute 232 + amax 154 + quant 255 + scale 4 = **645 µs** | routed amax 36 + quant 32 + scale 4 + FP8 permute 95 = **167 µs** | +| backward (gradient permute + quantize) | **687 µs** | **174 µs** | + +The FP8 permute and the quant kernel run at about 6.4 and 6.3 TB/s. + +**End-to-end.** Qwen3-30B-A3B FP8 tensorwise pretrain on 8× MI355X (EP8, MBS 8, GBS 512, seq 4096, even routing), mean of iterations 11–20 of 20-iteration runs. The base has every earlier change of this series plus Primus `turbo_fused_grouped_gemm`: + +| run | runs (ms/iter) | mean ms/iter | tokens/s/GPU | mem | loss@20 | +|---|---|---|---|---|---| +| `turbo_fused_grouped_gemm` | 6820.2, 6815.1 | 6817.7 | 38,451 | 224.5 GB | 11.33412 | +| + `turbo_fp8_permute` | 6442.4, 6515.0, 6523.8, 6549.1 | 6507.6 | 40,283 | 223.3 GB | 11.33404 | + +That is −4.55% ms/iter and +4.76% tokens/s/GPU. The later FP8-permute runs show more per-iteration jitter than the first one; I have not identified the cause. Loss tracks the baseline at every iteration, within run-to-run noise (about 4e-4 at iteration 20). + +## Type of change + +- [ ] Documentation change (change only to the documentation, either a fix or a new content) +- [ ] Bug fix (non-breaking change which fixes an issue) +- [x] New feature (non-breaking change which adds functionality) +- [ ] Breaking change (fix or feature that would cause existing functionality to not work as expected) +- [ ] Infra/Build change +- [ ] Code refactoring + +## Changes + +- `csrc/kernels/moe_permute/moe_permute.cu`, `csrc/include/primus_turbo/moe_permute.h`: `permute_routed_amax_kernel` / `permute_routed_amax_impl` (bf16 and fp16). +- `csrc/pytorch/moe_permute/moe_permute.cpp`, `extensions.h`, `bindings_pytorch.cpp`: `permute_routed_amax` op with CUDA and Meta implementations. +- `primus_turbo/pytorch/kernels/moe/moe_permute_impl.py`: `moe_permute_routed_amax_impl`. +- `primus_turbo/pytorch/ops/moe/moe_permute.py`: + - `quantize_dtype` / `grad_quantize_dtype` options, through a shared helper. + - Either option raises `ValueError` with a non-TURBO backend, and `quantize_dtype` also does with `use_fp8` / `scaling_factor`. +- `primus_turbo/pytorch/ops/grouped_mlp_fp8.py`: tensorwise backward accepts a `QuantizedTensor` `grad_out`, after checking its dtype and 128-column padding. +- `primus_turbo/pytorch/modules/moe/token_dispatcher.py`: `_post_dispatch(quantize_dtype=...)`, `_pre_combine(grad_quantize_dtype=...)`. +- Tests, described below. + +## Tests + +- `tests/pytorch/ops/test_moe_permute.py`: 144 passed, 19 skipped. + - New `test_moe_permute_quantize_tensorwise` and `test_moe_unpermute_grad_quantize_tensorwise`, over e4m3/e5m2, hidden 2048 and 1000, and `pad_multiple` 0/16. Outliers sit in the unrouted tail, and the results are bit-identical to quantizing the bf16 result. + - New `test_moe_permute_quantize_requires_turbo` and `test_moe_permute_quantize_empty_input`. +- `tests/pytorch/ops/test_grouped_mlp_fp8.py`: new `test_grouped_mlp_fp8_prequantized_input_and_grad`, over E4M3 and HYBRID. The output and every gradient are bit-identical to the bf16-input path. All tensorwise tests pass (14 passed). +- `tests/pytorch/modules/test_token_dispatcher.py` (2 GPUs): 25 passed. New `test_fp8_permute` covers `pad_multiple` 0/16 × `num_worst_tokens` 0/32768 and passed 5/5 repeated runs. + +The MXFP8 grouped-MLP path (`FP8GroupedMLPMXFunc`) is untouched. + +# Checklist: + +- [x] The functionality is complete +- [x] I have commented my code, particularly in hard-to-understand areas +- [ ] I have made corresponding changes to the documentation +- [ ] My changes generate no new warnings +- [x] I have added tests that prove my fix is effective or that my feature works +- [x] New and existing unit tests pass locally with my changes diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-permute-default-hip.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-permute-default-hip.md new file mode 100644 index 000000000..ece71245b --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-permute-default-hip.md @@ -0,0 +1,60 @@ +# Description + +`DeepEPTokenDispatcher` now runs `moe_permute` / `moe_unpermute` on the HIP (TURBO) backend by default, through a new `permute_backend` argument (TRITON is still selectable). + +**Motivation.** #455 made TRITON the generic `moe_permute` default, and the dispatcher inherited it because it never passed a backend. On MI355X the TURBO scatter / gather is faster at every shape measured. The gain is largest in the forward permute and in unpermute's backward, which are the two biggest permute costs in an MoE step. + +**Why only the dispatcher.** The generic `moe_permute` default stays TRITON. The TURBO kernel bounds its rows by the routed-token count, so it requires routed tokens to come first. DeepEP's receive buffer always satisfies this: every received token has a local expert, and worst-case buffer rows trail. Arbitrary callers may not satisfy it. The worst-case buffer + `pad_multiple > 0` padding-row bug in the TURBO permute was fixed in #553, so all dispatcher modes are covered. + +**Microbenchmark.** Qwen3-30B-A3B EP8 dispatcher call (hidden 2048, 16 local experts, top-8, topk probs, bf16), median µs, TURBO vs TRITON: + +| shape | op | TURBO | TRITON | +|---|---|---|---| +| even routing, 32768 tokens → 262144 rows | permute fwd | 255 | 390 | +| | permute bwd | 345 | 442 | +| | unpermute fwd | 253 | 257 | +| | unpermute bwd | 815 | 932 | +| random routing | permute fwd | 452 | 792 | +| hidden 7168, 32 local experts, even | permute fwd | 955 | 1011 | +| hidden 4096, 16 local experts, even | permute fwd | 465 | 579 | + +That is about 0.35 ms per MoE layer (fwd + bwd) at the Qwen3 shape. + +**End-to-end.** Qwen3-30B-A3B FP8 tensorwise pretrain on 8× MI355X (EP8, MBS 8, GBS 512, seq 4096, even routing), mean of iterations 11–20 of 20-iteration runs: + +| run | ms/iter | TFLOP/s/GPU | tokens/s/GPU | loss@20 | +|---|---|---|---|---| +| before (TRITON permute in the dispatcher) | 7774.2 | 767.2 | 33,720 | 11.33431 | +| this change (TURBO permute) | 7643.5 | 780.3 | 34,296 | 11.33423 | + +That is −1.68% ms/iter and +1.71% tokens/s/GPU. The loss difference is within run-to-run noise (about 4e-4 at iteration 20). The base of both runs is Turbo 1103b2df with the flat tensorwise FP8 quant kernel and the per-token `get_dispatch_layout` change applied, plus Primus `use_turbo_fused_act_with_probs`. + +## Type of change + +- [ ] Documentation change (change only to the documentation, either a fix or a new content) +- [ ] Bug fix (non-breaking change which fixes an issue) +- [x] New feature (non-breaking change which adds functionality) +- [ ] Breaking change (fix or feature that would cause existing functionality to not work as expected) +- [ ] Infra/Build change +- [ ] Code refactoring + +## Changes + +- `primus_turbo/pytorch/modules/moe/token_dispatcher.py`: `DeepEPTokenDispatcher(..., permute_backend=BackendType.TURBO)`. The backend is passed to both `moe_permute` (`_post_dispatch`) and `moe_unpermute` (`_pre_combine`). The docstring states the routed-prefix requirement. +- `tests/pytorch/modules/test_token_dispatcher.py`: `test_basic` and `test_worst_tokens` run on both permute backends (TURBO and TRITON). + +Behavior change: dispatcher users who relied on the implicit TRITON default now get TURBO. Pass `permute_backend=BackendType.TRITON` to keep the old path. With `pad_multiple > 0` the dispatcher already resolved to TURBO, because TRITON does not support padding. + +## Tests + +- `tests/pytorch/modules/test_token_dispatcher.py` (2 GPUs, MI355X / gfx950): `test_basic` and `test_worst_tokens` on both permute backends, 21 passed. +- Re-run today on a build that also contains #553 and the follow-up FP8 permute branch: 25 passed. + +# Checklist: + +- [x] The functionality is complete +- [x] I have commented my code, particularly in hard-to-understand areas +- [ ] I have made corresponding changes to the documentation +- [ ] My changes generate no new warnings +- [x] I have added tests that prove my fix is effective or that my feature works +- [x] New and existing unit tests pass locally with my changes diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/prompt.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/prompt.md new file mode 100644 index 000000000..b3e0ae9d1 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/prompt.md @@ -0,0 +1,104 @@ +# 任务：为 Qwen3-30B-A3B FP8 tensorwise 训练优化 Primus-Turbo kernel（MI355X / gfx950） + +## 背景 + +Primus Megatron 后端，Qwen3-30B-A3B FP8 预训练，单节点 8×MI355X。当前最好成绩：8894 ms/iter，670.6 TFLOP/s/GPU，29,474 tokens/s/GPU。 + +- config：`examples/megatron/configs/MI355X/qwen3_30B_A3B-FP8-pretrain.yaml`（Primus 分支 `perf/qwen3-30b-a3b-tuning`） +- 量化与 GEMM：`fp8: hybrid`，`fp8_recipe: tensorwise`，开启 `use_turbo_gemm` 和 `use_turbo_grouped_gemm`（FlyDSL grouped GEMM） +- MoE：`use_turbo_deepep`（`turbo_deepep_num_cu: 80`），`turbo_sync_free_moe_stage: 1` +- attention：TE CK v3（`NVTE_CK_USES_FWD_V3=1`，`NVTE_CK_USES_BWD_V3=1`，`NVTE_CK_IS_V3_ATOMIC_FP32=0`） +- Turbo 版本：e2d9f1d7 + +模型和并行形状： + +- 结构：48 层，hidden 2048，32 个 Q head、4 个 KV head（GQA），head_dim 128，开启 qk_layernorm +- MoE：128 个 expert，top-8，moe_ffn_hidden 768（SwiGLU，所以 fc1 输出宽 1536） +- 并行与 batch：TP1 / PP1 / EP8，每个 rank 16 个本地 expert；seq 4096，MBS 8，GBS 512（每步 8 个 microbatch） +- 每个 microbatch、每个 rank：32768 个 token，路由后约 262k 行（even routing 下每个 expert 约 16k 行） +- attention 输入布局：sbhd，causal + +## profile 结论 + +数据来自 rank 0 一个训练步（8 个 microbatch × 48 层），计算流的 GPU 忙碌时间占 97%。所以瓶颈完全在 kernel 本身，CPU 和 launch 不是问题。 + +| 类别 | 时间 | 占比 | 主要 kernel | +| --- | --- | --- | --- | +| attention | 1578 ms | 18% | `aiter::fmha_bwd_hd128_bf16_causal_a16` 1014 ms；fwd 390 ms；`ck_fused_attn::dk_dv_reduce` 81 ms | +| FP8 量化 | 1310 ms | 15% | `tensorwise_amax_partial_kernel` 252 ms（3264 次）；`quantize_tensorwise_pad_row_kernel` e5m2 220 ms + e4m3 206 ms；`transpose_2d_kernel` 83 ms（2304 次） | +| grouped GEMM | 1267 ms | 14.5% | `grouped_nt_persistent` 452 ms；`grouped_nn_persistent` 445 ms；`grouped_tn_wgrad_4wave` 361 ms | +| DeepEP | 1248 ms | 14% | `intranode::dispatch` 591 ms；`intranode::combine` 351 ms；`layout::get_dispatch_layout` 132 ms | +| dense GEMM（hipBLASLt） | 987 ms | 11% | QKV、proj 和 lm head 的 FP8/BF16 GEMM | +| norm / rope | 529 ms | 6% | TE `fused_rope_fwd` + `fused_rope_bwd` 204 ms；rmsnorm | +| permute | 404 ms | 4.6% | `unpermute_kernel` 189 ms；`permute_kernel` 157 ms | + +另外有约 1.1 s 的 elementwise 开销（SwiGLU 拆成 silu、mul、cat，再加 probs 乘法），Primus 侧打开 `use_turbo_fused_act_with_probs` 就能解决，不在本任务范围内。 + +## 需要优化的 kernel（按预期收益排序） + +### 1. FP8 tensorwise 量化链路（目标：省 0.6 s 以上 / step） + +现在每个 tensor 要先跑一遍 amax，再跑一遍 quantize（带 pad），有的还要额外跑一个 uint8 transpose 给 wgrad 用。至少 3 次读写 HBM。 + +- 把 amax 融合进上游 producer 的 epilogue（unpermute、rmsnorm、activation、DeepEP combine 输出），或者让 quantize 一次 kernel 同时输出 rowwise 和 colwise（transposed）的 FP8，把 `transpose_2d` 消掉。 +- 先算出 `tensorwise_amax_partial` 和 `quantize_tensorwise_pad_row` 在上述形状下的实际带宽（bytes / time），对照 MI355X 的 HBM 峰值，说明现在离峰值还差多少。 +- 评估 grouped GEMM 能不能直接吃 unpadded 的行（配合 Primus 的 `turbo_grouped_gemm_without_padding`），把 pad 的拷贝去掉。 + +### 2. FlyDSL grouped GEMM（FP8 tensorwise） + +按 `fwd + dgrad + wgrad` 估算，算力利用率大约 2.2 PFLOP/s，约为 FP8 峰值的 45%。 + +- 涉及的形状（每个 rank 16 个 expert，每个 expert 约 16k 行）： + - fc1：K=2048，N=1536 + - fc2：K=768，N=2048 + - 以及两者对应的 dgrad（nn）和 wgrad（tn） +- fc2 的 K=768 偏小；wgrad 的 tn 4wave kernel 尤其值得看。 +- 针对这些形状调 tile / split-K / persistent 调度，并给出每个形状 before/after 的 TFLOP/s。 +- 同时测 even routing 和 uniform routing 两种 expert 负载分布。 + +### 3. DeepEP intranode dispatch / combine + +- dispatch 目前传的是 bf16 hidden（2048），评估 FP8 dispatch：先按 tensorwise e4m3 量化再传，接收端直接把 FP8 和 scale 交给 fc1。这样通信量减半，还顺带省掉 fc1 前的那次 quantize。 +- `get_dispatch_layout` 有 132 ms，看能不能和 router 的 topk 融合，或者在 fwd/bwd 之间复用。 +- 测 `num_cu` 在 64、80、96 时的 kernel 时间。 + +### 4. attention backward（CK 路径，head_dim 128，sbhd，causal，GQA 32/4） + +- bwd/fwd 时间比约 2.6。分析 `dk_dv_reduce`（GQA 归约）能不能融进主 kernel。 +- 对比 FlyDSL attention：之前端到端测过，开 FlyDSL 比 CK 慢约 2%。说明 FlyDSL 在这个形状下的 fwd 和 bwd 分别差在哪里。 +- 注意：现有 `suachong/attn-bwd-q-cache-*` 分支只支持 head_dim 64，这里用不上。 + +### 5. qk_layernorm + RoPE 融合 + +Qwen3 有 q/k RMSNorm 后接 RoPE，现在是 TE 的 rmsnorm 加 `fused_rope` 分开跑。参考 `suachong/gptoss-fused-qk-rmsnorm-rope` 分支，给 head_dim 128 提供融合的 fwd 和 bwd。 + +### 6. permute / unpermute + +- `unpermute_kernel` 189 ms，`permute_kernel` 157 ms。 +- 看带宽利用率；如果 unpermute 能和 probs 加权、FP8 量化融合，就和第 1 项合并处理。 + +## 要求 + +- 每一项先写 microbench，复现上面的形状和当前耗时，再动 kernel。每项都要给出 before/after 的 kernel 时间，以及带宽或 TFLOP/s 利用率。 +- 正确性： + - 新增或扩展 `tests/pytorch/ops` 下的单测，覆盖 fwd 和 bwd、even 和 uniform 负载、非对齐 token 数。 + - FP8 结果和现有实现对比，误差不能变大。 + - attention 和 GEMM 不能引入新的非确定性。 +- 不能对现有路径造成回退：默认行为要么不变，要么新路径在全部现有测试下都通过。新功能可以先做成 opt-in 开关，并在 PR 里写清楚开关名。 +- 不要改 MXFP4 / MXFP8 / MegaMoE 的现有行为。 +- 每项优化单独一个分支、单独一个 PR，commit message 用英文，写明形状、收益和测试结果。 +- 完成后告诉我需要在 Primus 侧打开什么开关或改什么接口，我来做端到端验证（基准：8894 ms/iter，even routing，MBS 8，20 iter）。 + +## 复现 + +在 Primus 仓库下运行： + +```bash +./primus-cli direct -- train pretrain \ + --config examples/megatron/configs/MI355X/qwen3_30B_A3B-FP8-pretrain.yaml \ + --train_iters 14 --profile True --use_pytorch_profiler True \ + --profile_step_start 12 --profile_step_end 13 \ + --torch_profiler_with_stack False --torch_profiler_record_shapes True +``` + +trace 输出在 `output/amd/root/qwen3_30B_A3B-pretrain/tensorboard/*.pt.trace.json.gz`。开启 `record_shapes` 可以拿到每个 kernel 的实际输入形状。 diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_bitexact.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_bitexact.py new file mode 100644 index 000000000..cf8dbf333 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_bitexact.py @@ -0,0 +1,48 @@ +"""Dump (mode=dump) or compare (mode=check) tensorwise FP8 quant outputs byte-for-byte.""" + +import sys +import zlib + +import torch + +from primus_turbo.pytorch.kernels.quantization.quantization_impl import ( + quantize_fp8_tensorwise_pad_impl, +) + +PATH = "/tmp/quant_bitexact_ref.pt" +CASES = [] +for dt in (torch.bfloat16, torch.float16, torch.float32): + for shape in [(262144, 768), (4096, 2048), (1000, 1536), (37, 100), (3, 5), (17,), (8, 64, 96), (513, 1000)]: + for qdt in (torch.float8_e4m3fn, torch.float8_e5m2): + for k_align, pad_n in ((1, False), (128, False), (128, True)): + CASES.append((dt, shape, qdt, k_align, pad_n, 0)) +# misaligned input: a view starting one element in +for dt in (torch.bfloat16, torch.float32): + CASES.append((dt, (4096, 768), torch.float8_e4m3fn, 1, False, 1)) + + +def run(case): + dt, shape, qdt, k_align, pad_n, offset = case + g = torch.Generator(device="cuda").manual_seed(zlib.crc32(repr((dt, shape, qdt)).encode())) + n = 1 + for s in shape: + n *= s + base = torch.randn(n + offset, device="cuda", generator=g, dtype=torch.float32) * 4 + x = base.to(dt)[offset:].view(shape) # offset > 0 leaves data_ptr off 16-byte alignment + y, s = quantize_fp8_tensorwise_pad_impl(x, qdt, pad_n=pad_n, k_align=k_align) + return y.view(torch.uint8).cpu(), s.cpu() + + +mode = sys.argv[1] +if mode == "dump": + torch.save([run(c) for c in CASES], PATH) + print(f"dumped {len(CASES)} cases") +else: + ref = torch.load(PATH) + bad = 0 + for c, (ry, rs) in zip(CASES, ref): + y, s = run(c) + if not (torch.equal(y, ry) and torch.equal(s, rs)): + bad += 1 + print("MISMATCH", c) + print(f"checked {len(CASES)} cases, {bad} mismatches") diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/quant_variants.hip b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/quant_variants.hip new file mode 100644 index 000000000..32f1837cb --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/quant_variants.hip @@ -0,0 +1,160 @@ +// Standalone A/B harness for tensorwise FP8 quantize (no-pad case). +#include +#include +#include +#include +#include + +using u32x4 = __attribute__((vector_size(16))) uint32_t; +using u32x2 = __attribute__((vector_size(8))) uint32_t; + +__device__ __forceinline__ float bf16_to_f32(uint16_t v) { + return __uint_as_float(static_cast(v) << 16); +} + +template +__device__ __forceinline__ uint32_t cvt_pk(float a, float b, uint32_t acc) { + if constexpr (E5M2) { + return static_cast(__builtin_amdgcn_cvt_pk_bf8_f32(a, b, static_cast(acc), WORD_SEL)); + } else { + return static_cast(__builtin_amdgcn_cvt_pk_fp8_f32(a, b, static_cast(acc), WORD_SEL)); + } +} + +template +__device__ __forceinline__ float qclip(float x, float scale) { + constexpr float qmax = E5M2 ? 57344.0f : 448.0f; + const float v = x * scale; + return fmaxf(fminf(v, qmax), -qmax); +} + +// 8 bf16 (one u32x4) -> 2 dwords of fp8 +template +__device__ __forceinline__ void quant8(const u32x4 v, float scale, uint32_t &w0, uint32_t &w1) { + float f[8]; +#pragma unroll + for (int i = 0; i < 4; ++i) { + f[2 * i] = bf16_to_f32(static_cast(v[i] & 0xffffu)); + f[2 * i + 1] = bf16_to_f32(static_cast(v[i] >> 16)); + } + uint32_t lo = cvt_pk(qclip(f[0], scale), qclip(f[1], scale), 0u); + w0 = cvt_pk(qclip(f[2], scale), qclip(f[3], scale), lo); + lo = cvt_pk(qclip(f[4], scale), qclip(f[5], scale), 0u); + w1 = cvt_pk(qclip(f[6], scale), qclip(f[7], scale), lo); +} + +// ---- V0: copy of the current row-per-block kernel (UNROLL=8, BLOCK=256, no pad) ---- +template +__launch_bounds__(256) __global__ void v0_row_kernel(const uint16_t *__restrict__ x, uint8_t *__restrict__ y, + const float *scale_ptr, int32_t N, int32_t K) { + const float scale = scale_ptr[0]; + const int32_t packs_per_row = K / 8; + for (int32_t r = blockIdx.x; r < N; r += gridDim.x) { + const uint16_t *xr = x + static_cast(r) * K; + uint8_t *yr = y + static_cast(r) * K; + for (int32_t p = threadIdx.x; p < packs_per_row; p += 256) { + const u32x4 v = __builtin_nontemporal_load(reinterpret_cast(xr + p * 8)); + uint32_t w0, w1; + quant8(v, scale, w0, w1); + *reinterpret_cast(yr + p * 8) = u32x2{w0, w1}; + } + } +} + +// ---- V1/V2/V3: flat kernel, PACKS x 16 elements per thread, packs strided by BLOCK ---- +template +__launch_bounds__(BLOCK) __global__ void flat16_kernel(const uint16_t *__restrict__ x, uint8_t *__restrict__ y, + const float *scale_ptr, int64_t n16) { + const float scale = scale_ptr[0]; + const int64_t base = static_cast(blockIdx.x) * BLOCK * PACKS + threadIdx.x; + u32x4 a[PACKS], b[PACKS]; +#pragma unroll + for (int u = 0; u < PACKS; ++u) { + const int64_t i = base + static_cast(u) * BLOCK; + if (i < n16) { + const u32x4 *src = reinterpret_cast(x + i * 16); + if constexpr (NT) { + a[u] = __builtin_nontemporal_load(src); + b[u] = __builtin_nontemporal_load(src + 1); + } else { + a[u] = src[0]; + b[u] = src[1]; + } + } + } +#pragma unroll + for (int u = 0; u < PACKS; ++u) { + const int64_t i = base + static_cast(u) * BLOCK; + if (i < n16) { + uint32_t w0, w1, w2, w3; + quant8(a[u], scale, w0, w1); + quant8(b[u], scale, w2, w3); + *reinterpret_cast(y + i * 16) = u32x4{w0, w1, w2, w3}; + } + } +} + +// ---- V4: grid-stride persistent flat kernel ---- +template +__launch_bounds__(BLOCK) __global__ void gs16_kernel(const uint16_t *__restrict__ x, uint8_t *__restrict__ y, + const float *scale_ptr, int64_t n16) { + const float scale = scale_ptr[0]; + const int64_t stride = static_cast(gridDim.x) * BLOCK * PACKS; + for (int64_t base = static_cast(blockIdx.x) * BLOCK * PACKS + threadIdx.x; base < n16; + base += stride) { + u32x4 a[PACKS], b[PACKS]; +#pragma unroll + for (int u = 0; u < PACKS; ++u) { + const int64_t i = base + static_cast(u) * BLOCK; + if (i < n16) { + const u32x4 *src = reinterpret_cast(x + i * 16); + a[u] = __builtin_nontemporal_load(src); + b[u] = __builtin_nontemporal_load(src + 1); + } + } +#pragma unroll + for (int u = 0; u < PACKS; ++u) { + const int64_t i = base + static_cast(u) * BLOCK; + if (i < n16) { + uint32_t w0, w1, w2, w3; + quant8(a[u], scale, w0, w1); + quant8(b[u], scale, w2, w3); + *reinterpret_cast(y + i * 16) = u32x4{w0, w1, w2, w3}; + } + } + } +} + +template +void launch(int variant, const uint16_t *x, uint8_t *y, const float *s, int64_t N, int64_t K, hipStream_t st) { + const int64_t n16 = N * K / 16; + auto blocks = [&](int64_t per) { return static_cast((n16 + per - 1) / per); }; + const unsigned g0 = static_cast(std::min(N, 1048576)); + switch (variant) { + case 0: v0_row_kernel<<>>(x, y, s, N, K); break; + case 1: flat16_kernel<512, 1, false, E5M2><<>>(x, y, s, n16); break; + case 2: flat16_kernel<512, 1, true, E5M2><<>>(x, y, s, n16); break; + case 3: flat16_kernel<256, 2, true, E5M2><<>>(x, y, s, n16); break; + case 4: flat16_kernel<256, 4, true, E5M2><<>>(x, y, s, n16); break; + case 5: flat16_kernel<256, 1, true, E5M2><<>>(x, y, s, n16); break; + case 6: gs16_kernel<256, 2, E5M2><<<256 * 8, 256, 0, st>>>(x, y, s, n16); break; + case 7: gs16_kernel<512, 2, E5M2><<<256 * 4, 512, 0, st>>>(x, y, s, n16); break; + case 8: flat16_kernel<1024, 1, true, E5M2><<>>(x, y, s, n16); break; + default: TORCH_CHECK(false, "bad variant"); + } +} + +void run(int64_t variant, torch::Tensor x, torch::Tensor scale, torch::Tensor y) { + TORCH_CHECK(x.scalar_type() == at::kBFloat16 && x.is_contiguous() && x.dim() == 2); + TORCH_CHECK(x.size(1) % 16 == 0); + auto st = at::hip::getCurrentHIPStream(); + const bool e5 = y.scalar_type() == at::kFloat8_e5m2; + auto *xp = reinterpret_cast(x.data_ptr()); + auto *yp = reinterpret_cast(y.data_ptr()); + if (e5) + launch(variant, xp, yp, scale.data_ptr(), x.size(0), x.size(1), st); + else + launch(variant, xp, yp, scale.data_ptr(), x.size(0), x.size(1), st); +} + +PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { m.def("run", &run); } diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/quant_variants_hip.hip b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/quant_variants_hip.hip new file mode 100644 index 000000000..b2b826327 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/quant_variants_hip.hip @@ -0,0 +1,160 @@ +// Standalone A/B harness for tensorwise FP8 quantize (no-pad case). +#include +#include +#include +#include +#include + +using u32x4 = __attribute__((vector_size(16))) uint32_t; +using u32x2 = __attribute__((vector_size(8))) uint32_t; + +__device__ __forceinline__ float bf16_to_f32(uint16_t v) { + return __uint_as_float(static_cast(v) << 16); +} + +template +__device__ __forceinline__ uint32_t cvt_pk(float a, float b, uint32_t acc) { + if constexpr (E5M2) { + return static_cast(__builtin_amdgcn_cvt_pk_bf8_f32(a, b, static_cast(acc), WORD_SEL)); + } else { + return static_cast(__builtin_amdgcn_cvt_pk_fp8_f32(a, b, static_cast(acc), WORD_SEL)); + } +} + +template +__device__ __forceinline__ float qclip(float x, float scale) { + constexpr float qmax = E5M2 ? 57344.0f : 448.0f; + const float v = x * scale; + return fmaxf(fminf(v, qmax), -qmax); +} + +// 8 bf16 (one u32x4) -> 2 dwords of fp8 +template +__device__ __forceinline__ void quant8(const u32x4 v, float scale, uint32_t &w0, uint32_t &w1) { + float f[8]; +#pragma unroll + for (int i = 0; i < 4; ++i) { + f[2 * i] = bf16_to_f32(static_cast(v[i] & 0xffffu)); + f[2 * i + 1] = bf16_to_f32(static_cast(v[i] >> 16)); + } + uint32_t lo = cvt_pk(qclip(f[0], scale), qclip(f[1], scale), 0u); + w0 = cvt_pk(qclip(f[2], scale), qclip(f[3], scale), lo); + lo = cvt_pk(qclip(f[4], scale), qclip(f[5], scale), 0u); + w1 = cvt_pk(qclip(f[6], scale), qclip(f[7], scale), lo); +} + +// ---- V0: copy of the current row-per-block kernel (UNROLL=8, BLOCK=256, no pad) ---- +template +__launch_bounds__(256) __global__ void v0_row_kernel(const uint16_t *__restrict__ x, uint8_t *__restrict__ y, + const float *scale_ptr, int32_t N, int32_t K) { + const float scale = scale_ptr[0]; + const int32_t packs_per_row = K / 8; + for (int32_t r = blockIdx.x; r < N; r += gridDim.x) { + const uint16_t *xr = x + static_cast(r) * K; + uint8_t *yr = y + static_cast(r) * K; + for (int32_t p = threadIdx.x; p < packs_per_row; p += 256) { + const u32x4 v = __builtin_nontemporal_load(reinterpret_cast(xr + p * 8)); + uint32_t w0, w1; + quant8(v, scale, w0, w1); + *reinterpret_cast(yr + p * 8) = u32x2{w0, w1}; + } + } +} + +// ---- V1/V2/V3: flat kernel, PACKS x 16 elements per thread, packs strided by BLOCK ---- +template +__launch_bounds__(BLOCK) __global__ void flat16_kernel(const uint16_t *__restrict__ x, uint8_t *__restrict__ y, + const float *scale_ptr, int64_t n16) { + const float scale = scale_ptr[0]; + const int64_t base = static_cast(blockIdx.x) * BLOCK * PACKS + threadIdx.x; + u32x4 a[PACKS], b[PACKS]; +#pragma unroll + for (int u = 0; u < PACKS; ++u) { + const int64_t i = base + static_cast(u) * BLOCK; + if (i < n16) { + const u32x4 *src = reinterpret_cast(x + i * 16); + if constexpr (NT) { + a[u] = __builtin_nontemporal_load(src); + b[u] = __builtin_nontemporal_load(src + 1); + } else { + a[u] = src[0]; + b[u] = src[1]; + } + } + } +#pragma unroll + for (int u = 0; u < PACKS; ++u) { + const int64_t i = base + static_cast(u) * BLOCK; + if (i < n16) { + uint32_t w0, w1, w2, w3; + quant8(a[u], scale, w0, w1); + quant8(b[u], scale, w2, w3); + *reinterpret_cast(y + i * 16) = u32x4{w0, w1, w2, w3}; + } + } +} + +// ---- V4: grid-stride persistent flat kernel ---- +template +__launch_bounds__(BLOCK) __global__ void gs16_kernel(const uint16_t *__restrict__ x, uint8_t *__restrict__ y, + const float *scale_ptr, int64_t n16) { + const float scale = scale_ptr[0]; + const int64_t stride = static_cast(gridDim.x) * BLOCK * PACKS; + for (int64_t base = static_cast(blockIdx.x) * BLOCK * PACKS + threadIdx.x; base < n16; + base += stride) { + u32x4 a[PACKS], b[PACKS]; +#pragma unroll + for (int u = 0; u < PACKS; ++u) { + const int64_t i = base + static_cast(u) * BLOCK; + if (i < n16) { + const u32x4 *src = reinterpret_cast(x + i * 16); + a[u] = __builtin_nontemporal_load(src); + b[u] = __builtin_nontemporal_load(src + 1); + } + } +#pragma unroll + for (int u = 0; u < PACKS; ++u) { + const int64_t i = base + static_cast(u) * BLOCK; + if (i < n16) { + uint32_t w0, w1, w2, w3; + quant8(a[u], scale, w0, w1); + quant8(b[u], scale, w2, w3); + *reinterpret_cast(y + i * 16) = u32x4{w0, w1, w2, w3}; + } + } + } +} + +template +void launch(int variant, const uint16_t *x, uint8_t *y, const float *s, int64_t N, int64_t K, hipStream_t st) { + const int64_t n16 = N * K / 16; + auto blocks = [&](int64_t per) { return static_cast((n16 + per - 1) / per); }; + const unsigned g0 = static_cast(std::min(N, 1048576)); + switch (variant) { + case 0:hipLaunchKernelGGL(( v0_row_kernel), dim3(g0), dim3(256), 0, st, x, y, s, N, K); break; + case 1:hipLaunchKernelGGL(( flat16_kernel<512, 1, false, E5M2>), dim3(blocks(512)), dim3(512), 0, st, x, y, s, n16); break; + case 2:hipLaunchKernelGGL(( flat16_kernel<512, 1, true, E5M2>), dim3(blocks(512)), dim3(512), 0, st, x, y, s, n16); break; + case 3:hipLaunchKernelGGL(( flat16_kernel<256, 2, true, E5M2>), dim3(blocks(512)), dim3(256), 0, st, x, y, s, n16); break; + case 4:hipLaunchKernelGGL(( flat16_kernel<256, 4, true, E5M2>), dim3(blocks(1024)), dim3(256), 0, st, x, y, s, n16); break; + case 5:hipLaunchKernelGGL(( flat16_kernel<256, 1, true, E5M2>), dim3(blocks(256)), dim3(256), 0, st, x, y, s, n16); break; + case 6:hipLaunchKernelGGL(( gs16_kernel<256, 2, E5M2>), dim3(256 * 8), dim3(256), 0, st, x, y, s, n16); break; + case 7:hipLaunchKernelGGL(( gs16_kernel<512, 2, E5M2>), dim3(256 * 4), dim3(512), 0, st, x, y, s, n16); break; + case 8:hipLaunchKernelGGL(( flat16_kernel<1024, 1, true, E5M2>), dim3(blocks(1024)), dim3(1024), 0, st, x, y, s, n16); break; + default: TORCH_CHECK(false, "bad variant"); + } +} + +void run(int64_t variant, torch::Tensor x, torch::Tensor scale, torch::Tensor y) { + TORCH_CHECK(x.scalar_type() == at::kBFloat16 && x.is_contiguous() && x.dim() == 2); + TORCH_CHECK(x.size(1) % 16 == 0); + auto st = at::hip::getCurrentHIPStream(); + const bool e5 = y.scalar_type() == at::kFloat8_e5m2; + auto *xp = reinterpret_cast(x.data_ptr()); + auto *yp = reinterpret_cast(y.data_ptr()); + if (e5) + launch(variant, xp, yp, scale.data_ptr(), x.size(0), x.size(1), st); + else + launch(variant, xp, yp, scale.data_ptr(), x.size(0), x.size(1), st); +} + +PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { m.def("run", &run); } diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/run_variants.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/run_variants.py new file mode 100644 index 000000000..4c3bd6c74 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/quant_variants/run_variants.py @@ -0,0 +1,53 @@ +import os +import sys + +import torch +from torch.utils.cpp_extension import load + +HERE = os.path.dirname(os.path.abspath(__file__)) +ext = load( + name="quant_variants", + sources=[os.path.join(HERE, "quant_variants.hip")], + extra_cflags=["-O3"], + extra_cuda_cflags=["-O3", "--offload-arch=gfx950", "-std=c++20"], + build_directory=os.path.join(HERE, "build"), + verbose=False, +) + +from primus_turbo.pytorch.kernels.quantization.quantization_impl import ( # noqa: E402 + quantize_fp8_tensorwise_pad_impl, +) + + +def bench(fn, iters=50): + for _ in range(5): + fn() + s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + torch.cuda.synchronize() + s.record() + for _ in range(iters): + fn() + e.record() + torch.cuda.synchronize() + return s.elapsed_time(e) * 1e3 / iters + + +SHAPES = [(262144, 2048), (262144, 768), (262144, 1536), (32768, 2048), (32768, 4096), (32768, 5120), (4096, 768)] +VARIANTS = [int(v) for v in sys.argv[1].split(",")] if len(sys.argv) > 1 else list(range(9)) +for dt in (torch.float8_e4m3fn, torch.float8_e5m2): + print(f"== {dt}") + print(f"{'shape':>14s} " + " ".join(f"{'v' + str(v):>12s}" for v in VARIANTS)) + for m, k in SHAPES: + x = torch.randn(m, k, device="cuda", dtype=torch.bfloat16) * 3 + ref, sinv = quantize_fp8_tensorwise_pad_impl(x, dt, k_align=128) + scale = (1.0 / sinv).reshape(1).float().contiguous() + cells = [] + for v in VARIANTS: + y = torch.empty(m, k, device="cuda", dtype=dt) + ext.run(v, x, scale, y) + ok = torch.equal(y.view(torch.uint8), ref.view(torch.uint8)) + us = bench(lambda: ext.run(v, x, scale, y)) + bw = m * k * 3 / us / 1e6 + cells.append(f"{us:6.1f}/{bw:4.2f}{'' if ok else '!'}") + print(f"{m:7d}x{k:<6d} " + " ".join(f"{c:>12s}" for c in cells)) + del x, ref diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/qv_real.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/qv_real.py new file mode 100644 index 000000000..5afd5e993 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/qv_real.py @@ -0,0 +1,15 @@ +import torch +from primus_turbo.pytorch.kernels.quantization.quantization_impl import quantize_fp8_tensorwise_pad_impl +def bench(fn, iters=50): + for _ in range(5): fn() + s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + torch.cuda.synchronize(); s.record() + for _ in range(iters): fn() + e.record(); torch.cuda.synchronize() + return s.elapsed_time(e) * 1e3 / iters +for m, k in [(262144, 2048), (262144, 768), (262144, 1536), (32768, 2048), (32768, 4096), (32768, 5120)]: + x = torch.randn(m, k, device="cuda", dtype=torch.bfloat16) + p = x.abs().amax().float().reshape(1) + full = bench(lambda: quantize_fp8_tensorwise_pad_impl(x, torch.float8_e4m3fn, k_align=128)) + qonly = bench(lambda: quantize_fp8_tensorwise_pad_impl(x, torch.float8_e4m3fn, k_align=128, amax_partials=p)) + print(f"{m}x{k}: full {full:6.1f} us quant-only(+scale) {qonly:6.1f} us {m*k*3/qonly/1e6:4.2f} TB/s") diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_attrib.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_attrib.py new file mode 100644 index 000000000..e1f8917b5 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_attrib.py @@ -0,0 +1,59 @@ +"""Attribute GPU kernels to the CPU op stack that launched them (torch profiler trace). + +Usage: python trace_attrib.py TRACE.json.gz REGEX [REGEX ...] +For each kernel matching a regex, prints total time grouped by the enclosing op chain. +""" + +import bisect +import collections +import gzip +import json +import re +import sys + +path, patterns = sys.argv[1], [re.compile(p) for p in sys.argv[2:]] +with gzip.open(path, "rt") as f: + events = json.load(f)["traceEvents"] + +ops_by_tid = collections.defaultdict(list) +launch_by_corr = {} +kernels = [] +for e in events: + if e.get("ph") != "X": + continue + cat = e.get("cat") + if cat in ("cpu_op", "user_annotation", "python_function"): + ops_by_tid[e["tid"]].append((e["ts"], e["ts"] + e["dur"], e["name"])) + elif cat == "cuda_runtime": + corr = e.get("args", {}).get("correlation") + if corr is not None: + launch_by_corr[corr] = (e["tid"], e["ts"]) + elif cat == "kernel": + kernels.append(e) + +for tid in ops_by_tid: + ops_by_tid[tid].sort() +starts = {tid: [o[0] for o in ops] for tid, ops in ops_by_tid.items()} + + +def stack_at(tid, ts, depth=4): + ops = ops_by_tid.get(tid, []) + i = bisect.bisect_right(starts.get(tid, []), ts) + chain = [name for (s, e, name) in ops[max(0, i - 400) : i] if s <= ts <= e] + keep = [n for n in chain if not n.startswith(("cudaLaunch", "hipLaunch", "aten::empty"))] + return " < ".join(reversed(keep[-depth:])) + + +for rx in patterns: + groups = collections.defaultdict(lambda: [0, 0.0]) + for k in kernels: + if not rx.search(k["name"]): + continue + corr = k.get("args", {}).get("correlation") + tid, ts = launch_by_corr.get(corr, (None, None)) + key = stack_at(tid, ts) if tid is not None else "?" + groups[key][0] += 1 + groups[key][1] += k["dur"] + print(f"\n== /{rx.pattern}/") + for key, (n, us) in sorted(groups.items(), key=lambda kv: -kv[1][1])[:8]: + print(f"{us / 1e3:8.1f} ms {n:5d} x {us / n:7.1f} us {key[:260]}") diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_kernels.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_kernels.py new file mode 100644 index 000000000..fca101cec --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_kernels.py @@ -0,0 +1,37 @@ +"""Aggregate GPU kernel time per name from a torch profiler trace (one training step). + +Usage: python trace_kernels.py TRACE.json.gz [regex ...] +Prints the top kernels overall, then the kernels matching each regex. +""" + +import collections +import gzip +import json +import re +import sys + +path = sys.argv[1] +patterns = sys.argv[2:] +with gzip.open(path, "rt") as f: + trace = json.load(f) + +agg = collections.defaultdict(lambda: [0, 0.0]) +total = 0.0 +for e in trace["traceEvents"]: + if e.get("ph") != "X" or e.get("cat") != "kernel": + continue + agg[e["name"]][0] += 1 + agg[e["name"]][1] += e["dur"] + total += e["dur"] + +rows = sorted(agg.items(), key=lambda kv: -kv[1][1]) +print(f"total kernel time {total / 1e3:.1f} ms, {len(rows)} distinct kernels") +for name, (n, us) in rows[:40]: + print(f"{us / 1e3:9.1f} ms {n:6d} x {us / n:8.1f} us {name[:150]}") +for p in patterns: + rx = re.compile(p, re.I) + sel = [(k, v) for k, v in rows if rx.search(k)] + s = sum(v[1] for _, v in sel) + print(f"\n== /{p}/ : {s / 1e3:.1f} ms in {len(sel)} kernels") + for name, (n, us) in sel[:25]: + print(f"{us / 1e3:9.1f} ms {n:6d} x {us / n:8.1f} us {name[:170]}") diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_shapes.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_shapes.py new file mode 100644 index 000000000..cf5eda4d2 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/trace_shapes.py @@ -0,0 +1,60 @@ +"""Group GPU kernel time by the input shapes of the enclosing CPU op (torch profiler trace). + +Usage: python trace_shapes.py TRACE.json.gz OP_REGEX [KERNEL_REGEX] +For every kernel launched inside an op matching OP_REGEX (and matching KERNEL_REGEX), +prints total time per (op input dims, kernel name). +""" + +import bisect +import collections +import gzip +import json +import re +import sys + +path, op_rx = sys.argv[1], re.compile(sys.argv[2]) +k_rx = re.compile(sys.argv[3]) if len(sys.argv) > 3 else re.compile(".") +with gzip.open(path, "rt") as f: + events = json.load(f)["traceEvents"] + +ops_by_tid = collections.defaultdict(list) +launch_by_corr = {} +kernels = [] +for e in events: + if e.get("ph") != "X": + continue + cat = e.get("cat") + if cat == "cpu_op" and op_rx.search(e["name"]): + dims = e.get("args", {}).get("Input Dims") + ops_by_tid[e["tid"]].append((e["ts"], e["ts"] + e["dur"], str(dims))) + elif cat == "cuda_runtime": + corr = e.get("args", {}).get("correlation") + if corr is not None: + launch_by_corr[corr] = (e["tid"], e["ts"]) + elif cat == "kernel" and k_rx.search(e["name"]): + kernels.append(e) + +for tid in ops_by_tid: + ops_by_tid[tid].sort() +starts = {tid: [o[0] for o in ops] for tid, ops in ops_by_tid.items()} + +agg = collections.defaultdict(lambda: [0, 0.0]) +for k in kernels: + launch = launch_by_corr.get(k.get("args", {}).get("correlation")) + if launch is None: + continue + tid, ts = launch + ops = ops_by_tid.get(tid) + if not ops: + continue + i = bisect.bisect_right(starts[tid], ts) - 1 + while i >= 0 and ops[i][1] < ts: + i -= 1 + if i < 0: + continue + key = (ops[i][2][:90], re.sub(r"\(.*", "", k["name"])[:70]) + agg[key][0] += 1 + agg[key][1] += k["dur"] + +for (dims, name), (n, us) in sorted(agg.items(), key=lambda kv: -kv[1][1])[:30]: + print(f"{us / 1e3:8.1f} ms {n:5d} x {us / n:7.1f} us {dims:90s} {name}") From afb5e2f1322c6bfb78f5d38b38f81a3bc59a4499 Mon Sep 17 00:00:00 2001 From: x Date: Sat, 10 Oct 2026 15:00:04 +0000 Subject: [PATCH 02/14] docs(agent): record #554 and the rebased FP8 permute branch in the Qwen3 campaign notes --- agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md | 9 +++------ .../pr_bodies/perf-moe-fp8-permute-tensorwise.md | 2 +- 2 files changed, 4 insertions(+), 7 deletions(-) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md index 4512253d7..0a6ac6b2f 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -45,17 +45,14 @@ Turbo (`AMD-AGI/Primus-Turbo`): | branch | head | PR | notes | |---|---|---|---| | `fix/moe/permute-padding-rows` | f1a59fb9 | #553 merged | TURBO permute padding rows with worst-case buffers + pad | -| `perf/moe/permute-default-hip` | 30b6cc7c | open next | body: [pr_bodies/perf-moe-permute-default-hip.md](pr_bodies/perf-moe-permute-default-hip.md) | -| `perf/moe/fp8-permute-tensorwise` | 83e712a9 | after permute-default | stacked on it; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | +| `perf/moe/permute-default-hip` | 30b6cc7c | #554 open | body: [pr_bodies/perf-moe-permute-default-hip.md](pr_bodies/perf-moe-permute-default-hip.md) | +| `perf/moe/fp8-permute-tensorwise` | f4485b88 | opening | on main + #554's commit; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | | `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | not opened | independent; flat no-pad quant kernel | | `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | not opened | independent; `get_dispatch_layout` 340 -> 45 us | | `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | not opened | independent; needed by the Primus qk-norm patch | | `fix/fp8-attn-strided-layout` | c825b13e | not opened | FP8 attention with sbhd/bhsd storage | -Before opening the FP8 permute PR, rebase `perf/moe/fp8-permute-tensorwise` onto main once permute-default has merged: -- Drop a04d011b, its duplicate of #553. -- Put the permute-default PR number into the body. -- The commit message quotes the single-run gain (-5.54%); the PR body uses the 4-run mean (-4.55%). +`perf/moe/fp8-permute-tensorwise` is main (with #553) plus #554's commit plus the FP8 permute commit, whose message quotes the 4-run mean (-4.55%). The pre-rebase history is kept locally as `backup/fp8-permute-pre-rebase`. Once #554 merges, rebase the branch onto main and drop #554's commit. Primus (`AMD-AGI/Primus`): diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-fp8-permute-tensorwise.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-fp8-permute-tensorwise.md index 15378f153..093506fb9 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-fp8-permute-tensorwise.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-moe-fp8-permute-tensorwise.md @@ -1,6 +1,6 @@ # Description -> Stacked on # (`perf/moe/permute-default-hip`): the dispatcher must run the TURBO permute. I will rebase onto main once that merges; until then, review the last two commits (`perf(moe): emit tensorwise FP8 ...` and the format fix). +> Stacked on #554 (`perf/moe/permute-default-hip`): the dispatcher must run the TURBO permute. The first commit is #554's; please review the second one (`perf(moe): emit tensorwise FP8 from the TURBO permute for the grouped MLP`). I will rebase onto main once #554 merges. Under FP8 tensorwise, the fused grouped MLP (`grouped_mlp_fp8`) quantizes: - its input right after `DeepEPTokenDispatcher` permutes it, and From 6c00fe43af99244331e7f93bd1cf96ce58af48fc Mon Sep 17 00:00:00 2001 From: x Date: Sat, 10 Oct 2026 15:05:21 +0000 Subject: [PATCH 03/14] docs(agent): PR bodies for the remaining Qwen3 campaign Turbo branches --- .../qwen3_fp8_tw_gfx950_20261010/README.md | 10 +-- .../pr_bodies/fix-fp8-attn-strided-layout.md | 46 ++++++++++++++ .../perf-deep-ep-dispatch-layout-per-token.md | 62 +++++++++++++++++++ .../perf-flydsl-qk-rmsnorm-rope-hd128.md | 62 +++++++++++++++++++ .../perf-quantization-tensorwise-fp8-qwen3.md | 57 +++++++++++++++++ 5 files changed, 232 insertions(+), 5 deletions(-) create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-fp8-attn-strided-layout.md create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-deep-ep-dispatch-layout-per-token.md create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md index 0a6ac6b2f..73533899e 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -46,11 +46,11 @@ Turbo (`AMD-AGI/Primus-Turbo`): |---|---|---|---| | `fix/moe/permute-padding-rows` | f1a59fb9 | #553 merged | TURBO permute padding rows with worst-case buffers + pad | | `perf/moe/permute-default-hip` | 30b6cc7c | #554 open | body: [pr_bodies/perf-moe-permute-default-hip.md](pr_bodies/perf-moe-permute-default-hip.md) | -| `perf/moe/fp8-permute-tensorwise` | f4485b88 | opening | on main + #554's commit; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | -| `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | not opened | independent; flat no-pad quant kernel | -| `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | not opened | independent; `get_dispatch_layout` 340 -> 45 us | -| `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | not opened | independent; needed by the Primus qk-norm patch | -| `fix/fp8-attn-strided-layout` | c825b13e | not opened | FP8 attention with sbhd/bhsd storage | +| `perf/moe/fp8-permute-tensorwise` | f4485b88 | #555 open | on main + #554's commit; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | +| `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | not opened | independent; body: [pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md](pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md) | +| `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | not opened | independent; body: [pr_bodies/perf-deep-ep-dispatch-layout-per-token.md](pr_bodies/perf-deep-ep-dispatch-layout-per-token.md) | +| `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | not opened | independent, open next; needed by the Primus qk-norm patch; body: [pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md](pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md) | +| `fix/fp8-attn-strided-layout` | c825b13e | not opened | FP8 attention with sbhd/bhsd storage; body: [pr_bodies/fix-fp8-attn-strided-layout.md](pr_bodies/fix-fp8-attn-strided-layout.md) | `perf/moe/fp8-permute-tensorwise` is main (with #553) plus #554's commit plus the FP8 permute commit, whose message quotes the 4-run mean (-4.55%). The pre-rebase history is kept locally as `backup/fp8-permute-pre-rebase`. Once #554 merges, rebase the branch onto main and drop #554's commit. diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-fp8-attn-strided-layout.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-fp8-attn-strided-layout.md new file mode 100644 index 000000000..f761c1e07 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-fp8-attn-strided-layout.md @@ -0,0 +1,46 @@ +# Description + +`flash_attn_fp8_func` and `flash_attn_fp8_usp_func` mishandle q/k/v whose memory layout is sbhd or bhsd, which is what Megatron passes. This PR fixes both. + +**Bug.** +- Both functions take q/k/v whose **logical** shape is always `[b, s, h, d]`. `_infer_qkv_format` only reports how that shape is laid out in memory. +- Both functions nevertheless permuted the tensors as if the inferred format were the logical one. +- So a `[b, s, h, d]` view of sbhd storage became `[s, b, h, d]`. +- Block scaling then failed with `shape '[4096, 32, 0, 64, 128]' is invalid`, or, when the shapes happened to fit, silently attended over the wrong axis. +- bhsd storage hit the same bug. + +**Fix.** Make q/k/v bshd-contiguous instead, and hand the output back in the caller's storage layout with its logical `[b, s, h, d]` shape, as the BF16 path already does. + +## Type of change + +- [ ] Documentation change (change only to the documentation, either a fix or a new content) +- [x] Bug fix (non-breaking change which fixes an issue) +- [ ] New feature (non-breaking change which adds functionality) +- [ ] Breaking change (fix or feature that would cause existing functionality to not work as expected) +- [ ] Infra/Build change +- [ ] Code refactoring + +## Changes + +- `primus_turbo/pytorch/ops/attention/attention_utils.py`: new `_restore_qkv_storage(o, qkv_format)`, which gives the bshd-contiguous output the caller's storage layout and keeps its logical `[b, s, h, d]` shape. +- `primus_turbo/pytorch/ops/attention/flash_attn_interface.py`, `flash_attn_usp_interface.py`: + - Make q/k/v `.contiguous()` instead of permuting them by the inferred format. + - Restore the output with `_restore_qkv_storage`. +- `tests/pytorch/ops/test_attention.py`: new `test_attention_fp8_strided_layout`. + +## Tests + +- New `test_attention_fp8_strided_layout` covers bshd / sbhd / bhsd storage × causal / non-causal (GQA 16/4 heads, head_dim 128, seq 1024). It checks: + - the output shape; + - that the output keeps the caller's storage layout; + - output, dQ, dK and dV SNR > 20 dB vs the PyTorch reference. +- `tests/pytorch/ops/test_attention.py -k test_attention_fp8`: 78 passed (MI355X / gfx950). + +# Checklist: + +- [x] The functionality is complete +- [x] I have commented my code, particularly in hard-to-understand areas +- [ ] I have made corresponding changes to the documentation +- [ ] My changes generate no new warnings +- [x] I have added tests that prove my fix is effective or that my feature works +- [x] New and existing unit tests pass locally with my changes diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-deep-ep-dispatch-layout-per-token.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-deep-ep-dispatch-layout-per-token.md new file mode 100644 index 000000000..078b32df2 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-deep-ep-dispatch-layout-per-token.md @@ -0,0 +1,62 @@ +# Description + +`get_dispatch_layout` (intranode) is about 7.5× faster with a one-thread-per-token kernel. + +**Why it was slow.** The existing layout kernel gives each block 4 experts (or 8 ranks) and has it scan every token. So only about num_experts / 4 + 1 blocks run, and the rank block walks all tokens alone. + +**New kernel.** It applies to intranode buffers (no RDMA counters) with top-k ≤ 16. It: +- reads each `topk_idx` row once, with one thread per token; +- gathers expert / rank counts with shared-memory integer atomics, which are order-independent, so results stay deterministic; +- adds the counts once per block to zeroed global counters; +- writes `is_token_in_rank` with 64-bit stores when `num_ranks % 8 == 0`. + +Internode buffers, top-k > 16 and very large expert counts keep the previous kernel. + +**Microbenchmark.** Qwen3-30B-A3B EP8 on 8× MI355X (32768 tokens, top-8, 128 experts), rank-max median: + +| routing | before | after | +|---|---|---| +| random | 340 µs | 44 µs | +| even | 340 µs | 45 µs | + +The training step calls it 384 times (48 layers × 8 micro-batches), about −114 ms per iteration. Dispatch / combine times are unchanged. + +**End-to-end.** Qwen3-30B-A3B FP8 tensorwise pretrain on 8× MI355X (EP8, MBS 8, GBS 512, seq 4096, even routing), mean of iterations 11–20 of 20-iteration runs: + +| run | ms/iter | TFLOP/s/GPU | tokens/s/GPU | loss@20 | +|---|---|---|---|---| +| before | 7870.9 | 757.7 | 33,305 | 11.33415 | +| per-token `get_dispatch_layout` | 7774.2 | 767.2 | 33,720 | 11.33431 | + +That is −1.23% ms/iter and +1.25% tokens/s/GPU. The loss difference is within run-to-run noise (about 4e-4 at iteration 20). Both runs use Turbo 1103b2df plus the flat tensorwise quant kernel, and Primus `use_turbo_fused_act_with_probs`. + +## Type of change + +- [ ] Documentation change (change only to the documentation, either a fix or a new content) +- [ ] Bug fix (non-breaking change which fixes an issue) +- [x] New feature (non-breaking change which adds functionality) +- [ ] Breaking change (fix or feature that would cause existing functionality to not work as expected) +- [ ] Infra/Build change +- [ ] Code refactoring + +## Changes + +- `csrc/kernels/deep_ep/layout.cu`: new per-token intranode layout kernel, and dispatch between it and the previous kernel. +- `tests/pytorch/deep_ep/test_intranode.py`: new `test_dispatch_layout`. + +## Tests + +- New `test_dispatch_layout` compares rank / expert counts and `is_token_in_rank` with `get_dispatch_layout_ref` (`torch.equal`): + - routing: random, masked (-1) and even; + - 0 / 1 / 777 / 4099 / 32768 tokens; + - top-k 1 / 8 / 9, and top-k 20 for the fallback path. +- `tests/pytorch/deep_ep/test_intranode.py` on 8× MI355X: 2 passed (`test_dispatch_layout` and the existing intranode dispatch / combine test). + +# Checklist: + +- [x] The functionality is complete +- [x] I have commented my code, particularly in hard-to-understand areas +- [ ] I have made corresponding changes to the documentation +- [ ] My changes generate no new warnings +- [x] I have added tests that prove my fix is effective or that my feature works +- [x] New and existing unit tests pass locally with my changes diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md new file mode 100644 index 000000000..b42ba5779 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md @@ -0,0 +1,62 @@ +# Description + +The FlyDSL fused split + Q/K RMSNorm + RoPE kernels (forward, backward, gamma-grad fold) only accepted head_dim 64 (GPT-OSS). Qwen3 uses head_dim 128 with `qk_layernorm`, so on Qwen3 Primus's `PRIMUS_FUSED_QK_RMSNORM_ROPE` patch always fell back to split + TE RMSNorm + TE fused RoPE. This PR adds head_dim 128. + +**How.** head_dim is now a compile-time parameter taken from the k width: +- Each lane keeps its 8-element (dwordx4) chunk per half. +- A D=128 row therefore spans 8 lanes, and a wave carries 8 rows instead of 16. +- Row reductions, backward slot counts and the gamma fold loop follow from that. + +**gfx950 hazard.** At D=128 a row's hi half sits 128 bytes past its lo half, which no longer fits an inline `soffset` constant. On gfx950, a dwordx4 buffer store with an SGPR `soffset` followed by a VALU write of its data VGPRs is a hazard that LLVM does not pad: the last lanes of each 16-lane group stored clobbered data. D=128 therefore folds the hi-half offset into the VGPR offset instead. D=64 code generation is unchanged (identical ISA). + +**API.** `QK_RMSNORM_ROPE_HEAD_DIM` (an int) becomes `QK_RMSNORM_ROPE_HEAD_DIMS = (64, 128)`. Nothing else in Turbo or in Primus main imports the old name. The companion Primus patch reads the tuple and falls back to `(64,)` on older Turbo. + +**Microbenchmark.** Qwen3-30B-A3B attention shape on MI355X (S=4096, B=8, 4 KV groups × 8 Q heads per group, D=128), compared with split + TE RMSNorm + TE fused RoPE: + +| | unfused | fused | | +|---|---|---|---| +| forward | 519.5 µs | 126.8 µs | 5.29 TB/s | +| forward + backward | 1243.6 µs | 369.9 µs | 3.36× | + +The maximum difference vs the unfused path is 1 bf16 ulp. + +**End-to-end.** Qwen3-30B-A3B FP8 tensorwise pretrain on 8× MI355X (EP8, MBS 8, GBS 512, seq 4096, even routing), mean of iterations 11–20 of 20-iteration runs. The patch is enabled with `PRIMUS_FUSED_QK_RMSNORM_ROPE=1` and the companion Primus change (`AMD-AGI/Primus` branch `perf/qwen3-qk-rmsnorm-rope-hd128`): + +| run | ms/iter | TFLOP/s/GPU | tokens/s/GPU | loss@20 | +|---|---|---|---|---| +| before (unfused) | 7353.4 | 811.1 | 35,649 | 11.33456 | +| fused qk RMSNorm + RoPE, head_dim 128 | 6987.1 | 853.6 | 37,518 | 11.33373 | + +That is −4.98% ms/iter and +5.24% tokens/s/GPU. The loss difference is within run-to-run noise (about 4e-4 at iteration 20). Both runs use Turbo 1103b2df plus the flat tensorwise quant kernel, per-token `get_dispatch_layout` and the HIP permute in the dispatcher, and Primus `use_turbo_fused_act_with_probs` and `turbo_deepep_num_cu 160`. + +## Type of change + +- [ ] Documentation change (change only to the documentation, either a fix or a new content) +- [ ] Bug fix (non-breaking change which fixes an issue) +- [x] New feature (non-breaking change which adds functionality) +- [ ] Breaking change (fix or feature that would cause existing functionality to not work as expected) +- [ ] Infra/Build change +- [ ] Code refactoring + +## Changes + +- `primus_turbo/flydsl/rope/qk_rmsnorm_rope_kernel.py`: + - The forward, backward and gamma-fold kernels are parametric in head_dim. + - New `QK_RMSNORM_ROPE_HEAD_DIMS`; `_check_row_tileable` takes D. + - For D=128, the hi-half store offset goes in the VGPR offset. +- `primus_turbo/pytorch/kernels/rope/qk_rmsnorm_rope_impl.py`: the shape check takes D from the k width and accepts any head_dim in `QK_RMSNORM_ROPE_HEAD_DIMS`. +- `tests/pytorch/ops/test_qk_rmsnorm_rope.py`: D=128 cases. + +## Tests + +- `tests/pytorch/ops/test_qk_rmsnorm_rope.py` adds D=128 forward / backward cases, including an unaligned sequence length (S=129) and NG=1 / NPG=1. The unsupported-head-dim and untileable-shape errors are checked for both head dims. +- `test_qk_rmsnorm_rope.py` + `test_rope.py`: 40 passed (MI355X / gfx950). + +# Checklist: + +- [x] The functionality is complete +- [x] I have commented my code, particularly in hard-to-understand areas +- [ ] I have made corresponding changes to the documentation +- [ ] My changes generate no new warnings +- [x] I have added tests that prove my fix is effective or that my feature works +- [x] New and existing unit tests pass locally with my changes diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md new file mode 100644 index 000000000..4743c3635 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md @@ -0,0 +1,57 @@ +# Description + +This PR adds a flat tensorwise FP8 quant kernel for unpadded rows shorter than 4 KB. At the Qwen3 MoE fc2 width (K=768) the quant runs 39% faster, and the output is byte-identical. + +**Why.** The row-per-block tensorwise quant kernel gives each row one 256-lane block, with a 16-byte pack per lane. A row shorter than 4 KB therefore leaves lanes idle: only 96 of 256 lanes work at K=768 bf16. When nothing is padded (K == Kp, no N pad), the output keeps the input's flat layout, so a flat kernel with a fixed 16-element pack per thread keeps every lane busy. + +**When it applies.** Only for rows shorter than 4 KB. Longer rows keep the row kernel, which is as fast or faster there. Misaligned pointers fall back as before. + +**Microbenchmark.** Qwen3-30B-A3B MoE shapes on MI355X, quant only, bf16 → e4m3: + +| shape | before | after | +|---|---|---| +| 262144 × 768 | 156.5 µs | 95.3 µs (−39%, 6.34 TB/s) | +| 262144 × 1536 | 221 µs | 207.5 µs (−6%) | +| 262144 × 2048, 32768 × {2048, 4096, 5120} | unchanged (row kernel) | | + +The full `quantize_fp8` at 262144 × 768 goes from 217 µs to 156 µs. The output is byte-identical to the previous kernel on 146 shape/dtype cases. + +**End-to-end.** Qwen3-30B-A3B FP8 tensorwise pretrain on 8× MI355X (EP8, MBS 8, GBS 512, seq 4096, even routing), mean of iterations 11–20 of 20-iteration runs: + +| run | ms/iter | TFLOP/s/GPU | tokens/s/GPU | loss@20 | +|---|---|---|---|---| +| Turbo 1103b2df | 7925.4 | 752.5 | 33,077 | 11.33427 | +| flat tensorwise FP8 quant kernel | 7870.9 | 757.7 | 33,305 | 11.33415 | + +That is −0.69% ms/iter and +0.69% tokens/s/GPU. The loss difference is within run-to-run noise (about 4e-4 at iteration 20). Both runs use Primus `use_turbo_fused_act_with_probs`. + +## Type of change + +- [ ] Documentation change (change only to the documentation, either a fix or a new content) +- [ ] Bug fix (non-breaking change which fixes an issue) +- [x] New feature (non-breaking change which adds functionality) +- [ ] Breaking change (fix or feature that would cause existing functionality to not work as expected) +- [ ] Infra/Build change +- [ ] Code refactoring + +## Changes + +- `csrc/kernels/quantization/quantization_tensorwise.cu`: flat no-pad kernel, selected when nothing is padded and rows are shorter than 4 KB. +- `tests/pytorch/ops/test_quantization.py`: new `test_quantize_fp8_tensorwise_no_pad_rows`. + +## Tests + +- New `test_quantize_fp8_tensorwise_no_pad_rows`: + - covers bf16 / fp16 / fp32 × e4m3 / e5m2, unaligned row counts, and 1-row and odd-K shapes; + - checks against the reference, and byte equality with the misaligned fallback path. +- `tests/pytorch/ops/test_quantization.py -k tensorwise`: 116 passed (MI355X / gfx950). +- The tensorwise FP8 GEMM / grouped GEMM tests pass. + +# Checklist: + +- [x] The functionality is complete +- [x] I have commented my code, particularly in hard-to-understand areas +- [ ] I have made corresponding changes to the documentation +- [ ] My changes generate no new warnings +- [x] I have added tests that prove my fix is effective or that my feature works +- [x] New and existing unit tests pass locally with my changes From 66c15f9fb86c7a5fe8e887428ae4d45f79792be8 Mon Sep 17 00:00:00 2001 From: x Date: Sat, 10 Oct 2026 15:07:23 +0000 Subject: [PATCH 04/14] docs(agent): rope PR opened as #556 --- agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md index 73533899e..6f74ef153 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -48,8 +48,8 @@ Turbo (`AMD-AGI/Primus-Turbo`): | `perf/moe/permute-default-hip` | 30b6cc7c | #554 open | body: [pr_bodies/perf-moe-permute-default-hip.md](pr_bodies/perf-moe-permute-default-hip.md) | | `perf/moe/fp8-permute-tensorwise` | f4485b88 | #555 open | on main + #554's commit; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | | `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | not opened | independent; body: [pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md](pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md) | -| `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | not opened | independent; body: [pr_bodies/perf-deep-ep-dispatch-layout-per-token.md](pr_bodies/perf-deep-ep-dispatch-layout-per-token.md) | -| `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | not opened | independent, open next; needed by the Primus qk-norm patch; body: [pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md](pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md) | +| `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | opening | independent; body: [pr_bodies/perf-deep-ep-dispatch-layout-per-token.md](pr_bodies/perf-deep-ep-dispatch-layout-per-token.md) | +| `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | #556 open | independent; needed by the Primus qk-norm patch; body: [pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md](pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md) | | `fix/fp8-attn-strided-layout` | c825b13e | not opened | FP8 attention with sbhd/bhsd storage; body: [pr_bodies/fix-fp8-attn-strided-layout.md](pr_bodies/fix-fp8-attn-strided-layout.md) | `perf/moe/fp8-permute-tensorwise` is main (with #553) plus #554's commit plus the FP8 permute commit, whose message quotes the 4-run mean (-4.55%). The pre-rebase history is kept locally as `backup/fp8-permute-pre-rebase`. Once #554 merges, rebase the branch onto main and drop #554's commit. From fd3f9a0622d8af84d9339ada0a10d7c8150e6b83 Mon Sep 17 00:00:00 2001 From: x Date: Sat, 10 Oct 2026 15:09:25 +0000 Subject: [PATCH 05/14] docs(agent): layout PR opened as #557; quant PR body test line --- agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md | 4 ++-- .../pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md index 6f74ef153..178fe32d6 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -47,8 +47,8 @@ Turbo (`AMD-AGI/Primus-Turbo`): | `fix/moe/permute-padding-rows` | f1a59fb9 | #553 merged | TURBO permute padding rows with worst-case buffers + pad | | `perf/moe/permute-default-hip` | 30b6cc7c | #554 open | body: [pr_bodies/perf-moe-permute-default-hip.md](pr_bodies/perf-moe-permute-default-hip.md) | | `perf/moe/fp8-permute-tensorwise` | f4485b88 | #555 open | on main + #554's commit; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | -| `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | not opened | independent; body: [pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md](pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md) | -| `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | opening | independent; body: [pr_bodies/perf-deep-ep-dispatch-layout-per-token.md](pr_bodies/perf-deep-ep-dispatch-layout-per-token.md) | +| `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | opening | independent; body: [pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md](pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md) | +| `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | #557 open | independent; body: [pr_bodies/perf-deep-ep-dispatch-layout-per-token.md](pr_bodies/perf-deep-ep-dispatch-layout-per-token.md) | | `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | #556 open | independent; needed by the Primus qk-norm patch; body: [pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md](pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md) | | `fix/fp8-attn-strided-layout` | c825b13e | not opened | FP8 attention with sbhd/bhsd storage; body: [pr_bodies/fix-fp8-attn-strided-layout.md](pr_bodies/fix-fp8-attn-strided-layout.md) | diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md index 4743c3635..bc739086b 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md @@ -45,7 +45,7 @@ That is −0.69% ms/iter and +0.69% tokens/s/GPU. The loss difference is within - covers bf16 / fp16 / fp32 × e4m3 / e5m2, unaligned row counts, and 1-row and odd-K shapes; - checks against the reference, and byte equality with the misaligned fallback path. - `tests/pytorch/ops/test_quantization.py -k tensorwise`: 116 passed (MI355X / gfx950). -- The tensorwise FP8 GEMM / grouped GEMM tests pass. +- Tensorwise FP8 GEMM / grouped GEMM tests (which quantize through this path): about 1400 passed. # Checklist: From 8ded5aac7187121f91bdd4fb8270c1aa595a32a8 Mon Sep 17 00:00:00 2001 From: x Date: Sat, 10 Oct 2026 15:11:05 +0000 Subject: [PATCH 06/14] docs(agent): quant PR opened as #558; fp8-attn PR body test line --- agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md | 4 ++-- .../pr_bodies/fix-fp8-attn-strided-layout.md | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md index 178fe32d6..bf01ee76b 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -47,10 +47,10 @@ Turbo (`AMD-AGI/Primus-Turbo`): | `fix/moe/permute-padding-rows` | f1a59fb9 | #553 merged | TURBO permute padding rows with worst-case buffers + pad | | `perf/moe/permute-default-hip` | 30b6cc7c | #554 open | body: [pr_bodies/perf-moe-permute-default-hip.md](pr_bodies/perf-moe-permute-default-hip.md) | | `perf/moe/fp8-permute-tensorwise` | f4485b88 | #555 open | on main + #554's commit; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | -| `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | opening | independent; body: [pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md](pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md) | +| `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | #558 open | independent; body: [pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md](pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md) | | `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | #557 open | independent; body: [pr_bodies/perf-deep-ep-dispatch-layout-per-token.md](pr_bodies/perf-deep-ep-dispatch-layout-per-token.md) | | `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | #556 open | independent; needed by the Primus qk-norm patch; body: [pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md](pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md) | -| `fix/fp8-attn-strided-layout` | c825b13e | not opened | FP8 attention with sbhd/bhsd storage; body: [pr_bodies/fix-fp8-attn-strided-layout.md](pr_bodies/fix-fp8-attn-strided-layout.md) | +| `fix/fp8-attn-strided-layout` | c825b13e | opening | FP8 attention with sbhd/bhsd storage; body: [pr_bodies/fix-fp8-attn-strided-layout.md](pr_bodies/fix-fp8-attn-strided-layout.md) | `perf/moe/fp8-permute-tensorwise` is main (with #553) plus #554's commit plus the FP8 permute commit, whose message quotes the 4-run mean (-4.55%). The pre-rebase history is kept locally as `backup/fp8-permute-pre-rebase`. Once #554 merges, rebase the branch onto main and drop #554's commit. diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-fp8-attn-strided-layout.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-fp8-attn-strided-layout.md index f761c1e07..c682b2605 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-fp8-attn-strided-layout.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-fp8-attn-strided-layout.md @@ -34,7 +34,7 @@ - the output shape; - that the output keeps the caller's storage layout; - output, dQ, dK and dV SNR > 20 dB vs the PyTorch reference. -- `tests/pytorch/ops/test_attention.py -k test_attention_fp8`: 78 passed (MI355X / gfx950). +- `tests/pytorch/ops/test_attention.py -k test_attention_fp8`: 78 passed, including the 6 new cases (MI355X / gfx950). # Checklist: From 31c71fa8627ae5a4b4b0c24d0b0965ab872e37b4 Mon Sep 17 00:00:00 2001 From: sim Date: Sat, 10 Oct 2026 15:39:48 +0000 Subject: [PATCH 07/14] docs(agent): one Primus PR (perf/qwen3-30b-a3b-fp8-turbo); intermittent NaN open item --- .../qwen3_fp8_tw_gfx950_20261010/README.md | 39 ++++++---- .../e2e/e2e_run_cfg.sh | 13 ++++ .../e2e/e2e_run_pr.sh | 15 ++++ .../e2e/env_turbo_opt_cfg.sh | 2 + .../primus-perf-qwen3-30b-a3b-fp8-turbo.md | 74 +++++++++++++++++++ 5 files changed, 127 insertions(+), 16 deletions(-) create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run_cfg.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run_pr.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt_cfg.sh create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/primus-perf-qwen3-30b-a3b-fp8-turbo.md diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md index bf01ee76b..91593b5e2 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -25,18 +25,20 @@ Run-to-run loss noise is about 4e-4 at iteration 20. ## Settings for the last row -The final configuration is `examples/megatron/configs/MI355X/qwen3_30B_A3B-FP8-pretrain.yaml` plus: -- `--use_turbo_fused_act_with_probs True` -- `--turbo_deepep_num_cu 160` -- `--turbo_fused_grouped_gemm True` -- `--turbo_fp8_permute True` -- env `PRIMUS_FUSED_QK_RMSNORM_ROPE=1` +On the Primus branch `perf/qwen3-30b-a3b-fp8-turbo`, `examples/megatron/configs/MI355X/qwen3_30B_A3B-FP8-pretrain.yaml` turns all of these on: +- `use_turbo_fused_act_with_probs: true` +- `turbo_deepep_num_cu: 160` +- `turbo_fused_grouped_gemm: true` +- `turbo_fp8_permute: true` +- top-level `env: PRIMUS_FUSED_QK_RMSNORM_ROPE: "1"` + +On that branch, the config alone gave 6410.6 and 6459.6 ms/iter (40,892 / 40,582 tokens/s/GPU). One more config-only run hit the intermittent NaN (see open items). Code it needs: - Turbo: all five Turbo branches below. Building on `perf/moe/fp8-permute-tensorwise` also brings in the permute-default change. -- Primus: `perf/qwen3-qk-rmsnorm-rope-hd128`, which sits on `perf/qwen3-30b-a3b-tuning`, plus `perf/megatron/turbo-fp8-permute`. +- Primus: `perf/qwen3-30b-a3b-fp8-turbo`. -See [e2e/e2e_run.sh](e2e/e2e_run.sh) and [e2e/env_turbo_opt_qknorm.sh](e2e/env_turbo_opt_qknorm.sh). Their paths are for the old machine. +See [e2e/e2e_run_cfg.sh](e2e/e2e_run_cfg.sh) (config only) and [e2e/env_turbo_opt_cfg.sh](e2e/env_turbo_opt_cfg.sh). The older [e2e/e2e_run.sh](e2e/e2e_run.sh) passes the switches on the command line. All paths are for the old machine. ## Branches and PRs @@ -58,9 +60,9 @@ Primus (`AMD-AGI/Primus`): | branch | head | PR | notes | |---|---|---|---| -| `perf/qwen3-30b-a3b-tuning` | 82ebf666 | not opened | Qwen3 FP8 config tuned to Turbo tensorwise | -| `perf/qwen3-qk-rmsnorm-rope-hd128` | bfa227f2 | not opened | on top of the tuning branch; fused qk-norm patch for head_dim 128 | -| `perf/megatron/turbo-fp8-permute` | 69a32f64 | not opened | on main; `turbo_fp8_permute` flag | +| `perf/qwen3-30b-a3b-fp8-turbo` | bef41e8a | opening | the one Primus PR, on main 0a68e5cd; body: [pr_bodies/primus-perf-qwen3-30b-a3b-fp8-turbo.md](pr_bodies/primus-perf-qwen3-30b-a3b-fp8-turbo.md) | + +It holds five commits: the FP8 attention fix, the tensorwise config, the head_dim 128 qk-norm patch, `turbo_fp8_permute` (with new unit tests), and the config turning every switch on. It replaces `perf/qwen3-30b-a3b-tuning`, `perf/qwen3-qk-rmsnorm-rope-hd128` and `perf/megatron/turbo-fp8-permute`. Their force-"even" routing fix is already on main (#1243). ## Rebuilding the stack on a new machine @@ -80,15 +82,14 @@ GPU_ARCHS=gfx950 MAX_JOBS=64 python setup.py build_ext --inplace Primus: ```bash -git fetch origin && git checkout -b integ/qwen3-fp8 origin/perf/qwen3-qk-rmsnorm-rope-hd128 -git merge --no-edit origin/perf/megatron/turbo-fp8-permute +git fetch origin && git checkout -b integ/qwen3-fp8 origin/perf/qwen3-30b-a3b-fp8-turbo git submodule update --init third_party/Megatron-LM ``` -Then edit the paths in `e2e/env_turbo_opt_qknorm.sh` and `e2e/e2e_run.sh` for the new machine, and run, for example: +Then edit the paths in `e2e/env_turbo_opt_cfg.sh` and `e2e/e2e_run_cfg.sh` for the new machine, and run, for example: ```bash -e2e/e2e_run.sh g1 --turbo_fused_grouped_gemm True --turbo_fp8_permute True +e2e/e2e_run_cfg.sh cfg1 ``` Summarize the runs with `e2e/parse_e2e.py `, which reads `/tmp/q3_e2e_.log`. @@ -99,7 +100,7 @@ Run Python from the campaign directory with `PYTHONPATH=`. Other - **MegaMoE** (bf16 and mxfp8) runs out of memory at MBS 8. When the worst-case dispatch pool allocation fails, 261 GB (bf16, failing on 3 GiB) or 264 GB (mxfp8, failing on 1.5 GiB) is already allocated. Using it would need a MegaMoE change, a smaller MBS, or recompute. - **`turbo_deepep_num_cu`**: - - 192 gives a NaN forward loss on rank 2 at iteration 5, a DeepEP bug at high channel counts. + - 192 gave a NaN forward loss on rank 2 at iteration 5 in its only run. 160 also hit it once (see open items), so this is not specific to 192. - 256 is slower (7028 ms). - 160 is best. - **Sync-free MoE stage 2** costs +300 ms over stage 1. `num_worst_tokens` makes every post-dispatch step scan 262144 receive rows instead of 32768 tokens. @@ -111,6 +112,12 @@ Run Python from the campaign directory with `PYTHONPATH=`. Other ## Open items +- **Intermittent NaN forward loss.** Run `cfg1` of the final config (`perf/qwen3-30b-a3b-fp8-turbo`, config only) died at iteration 3 with `found NaN in local forward loss calculation` on rank 4. The per-rank traceback is in `output/.../logs/pre_trainer/rank-4/error.log`. + - Its 847 parsed arguments match the passing run `pr1` exactly, and the wgrad autotune winners match too. + - At `turbo_deepep_num_cu 160`, 1 of 13 runs hit it. The only run at 192 hit it (rank 2, iteration 5). The only run at 256 did not. + - None of the 5 runs at 80 CUs hit it, but those predate the fused qk-norm, fused grouped MLP and FP8 permute, so the cause is not isolated. + - The `turbo-opt` build includes the #553 padding fix. + - Next step: a DeepEP intranode dispatch/combine stress loop with exact checks at `num_sms` 80 / 160 / 192. It runs much faster than a 1-in-13 end-to-end bisection. - **Unexplained e2e jitter with FP8 permute.** The first run took 6442 ms; three later runs took 6515-6549 ms with more per-iteration jitter. These causes are ruled out: - the Primus change: a same-code rerun gave 6515 ms; - the padding fix: permute kernel times are unchanged; diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run_cfg.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run_cfg.sh new file mode 100644 index 000000000..e20bb3411 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run_cfg.sh @@ -0,0 +1,13 @@ +#!/bin/bash +# One 20-iteration Qwen3-30B-A3B FP8 run on the consolidated Primus PR branch using only +# the example config: no switch on the command line, no PRIMUS_FUSED_QK_RMSNORM_ROPE in the env. +# Usage: e2e_run_cfg.sh NAME [extra primus args ...] log: /tmp/q3_e2e_.log +ulimit -c 0 +E2E=/home/xiaompen/turbo-opt/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e +CFG=examples/megatron/configs/MI355X/qwen3_30B_A3B-FP8-pretrain.yaml +name=$1 +shift +echo "[e2e] $name start $(date +%T)" +(cd /home/xiaompen/Primus-qwen3-pr && bash runner/primus-cli direct --env "$E2E/env_turbo_opt_cfg.sh" -- \ + train pretrain --config "$CFG" --train_iters 20 "$@") > "/tmp/q3_e2e_${name}.log" 2>&1 +echo "[e2e] $name exit=$? $(date +%T)" diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run_pr.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run_pr.sh new file mode 100644 index 000000000..ad19276bd --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/e2e_run_pr.sh @@ -0,0 +1,15 @@ +#!/bin/bash +# One 20-iteration Qwen3-30B-A3B FP8 run on the consolidated Primus PR branch +# (perf/qwen3-30b-a3b-fp8-turbo) with every switch of the 40k row turned on. +# Usage: e2e_run_pr.sh NAME [extra primus args ...] log: /tmp/q3_e2e_.log +ulimit -c 0 +E2E=/home/xiaompen/turbo-opt/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e +CFG=examples/megatron/configs/MI355X/qwen3_30B_A3B-FP8-pretrain.yaml +name=$1 +shift +echo "[e2e] $name start $(date +%T)" +(cd /home/xiaompen/Primus-qwen3-pr && bash runner/primus-cli direct --env "$E2E/env_turbo_opt_qknorm.sh" -- \ + train pretrain --config "$CFG" --train_iters 20 --use_turbo_fused_act_with_probs True \ + --turbo_deepep_num_cu 160 --turbo_fused_grouped_gemm True --turbo_fp8_permute True "$@") \ + > "/tmp/q3_e2e_${name}.log" 2>&1 +echo "[e2e] $name exit=$? $(date +%T)" diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt_cfg.sh b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt_cfg.sh new file mode 100644 index 000000000..42f91b731 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/env_turbo_opt_cfg.sh @@ -0,0 +1,2 @@ +export PYTHONPATH=/home/xiaompen/turbo-opt:/home/xiaompen/Primus-qwen3/third_party/Megatron-LM:${PYTHONPATH} +export FLYDSL_RUNTIME_CACHE_DIR=/tmp/fly_cache_opt diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/primus-perf-qwen3-30b-a3b-fp8-turbo.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/primus-perf-qwen3-30b-a3b-fp8-turbo.md new file mode 100644 index 000000000..d1c8faf3a --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/primus-perf-qwen3-30b-a3b-fp8-turbo.md @@ -0,0 +1,74 @@ +# Description + +Qwen3-30B-A3B FP8 pretraining on 8× MI355X goes from a memory fault at iteration 3 to **40,283 tokens/s/GPU** (6507.6 ms/iter), 47% above the BF16 config. This PR switches the FP8 example config to Primus-Turbo tensorwise FP8 and turns on every Turbo fusion that got there. Two of those fusions are new in this PR: +- **head_dim 128 for the fused QKV split + Q/K RMSNorm + RoPE patch.** The patch was GPT-OSS-only before. +- **`turbo_fp8_permute`**: the DeepEP permute hands the fused grouped MLP FP8 tokens, and takes its gradient back in FP8. + +The PR also fixes `enable_turbo_attention_float8`, which failed on its first call. + +**Primus-Turbo requirements.** +- `turbo_fp8_permute` needs a Primus-Turbo whose `moe_permute` takes `quantize_dtype` ([AMD-AGI/Primus-Turbo#555](https://github.com/AMD-AGI/Primus-Turbo/pull/555)). With an older Turbo, the first dispatch fails. +- The head_dim 128 fusion needs [AMD-AGI/Primus-Turbo#556](https://github.com/AMD-AGI/Primus-Turbo/pull/556). Without it, the patch keeps the unfused path. +- The other Turbo rows in the table below come from [#554](https://github.com/AMD-AGI/Primus-Turbo/pull/554), [#557](https://github.com/AMD-AGI/Primus-Turbo/pull/557) and [#558](https://github.com/AMD-AGI/Primus-Turbo/pull/558). They need no Primus change. + +**End-to-end.** Qwen3-30B-A3B FP8 pretrain on 8× MI355X (EP8, MBS 8, GBS 512, seq 4096, even routing), mean of iterations 11–20 of 20-iteration runs. Each row adds one change on top of the row above: + +| run | from | ms/iter | TFLOP/s/GPU | tokens/s/GPU | mem | loss@20 | +|---|---|---|---|---|---|---| +| BF16 config (legacy grouped GEMM) | | 9587 | 622.1 | 27,344 | | | +| FP8 config before this PR | | memory fault at iteration 3 | | | | | +| FP8 tensorwise config (Turbo e2d9f1d7) | this PR | 8894.0 | 670.6 | 29,474 | 277 GB | 11.33393 | +| + `use_turbo_fused_act_with_probs` | this PR | 7915.6 | 753.5 | 33,117 | 223 GB | 11.33435 | +| Turbo 1103b2df | Turbo main | 7925.4 | 752.5 | 33,077 | 222 GB | 11.33427 | +| + flat tensorwise FP8 quant kernel | Turbo #558 | 7870.9 | 757.7 | 33,305 | 222 GB | 11.33415 | +| + per-token `get_dispatch_layout` | Turbo #557 | 7774.2 | 767.2 | 33,720 | 222 GB | 11.33431 | +| + HIP permute in `DeepEPTokenDispatcher` | Turbo #554 | 7643.5 | 780.3 | 34,296 | 222 GB | 11.33423 | +| + `turbo_deepep_num_cu` 80 → 160 | this PR | 7353.4 | 811.1 | 35,649 | 223 GB | 11.33456 | +| + fused QKV split + Q/K RMSNorm + RoPE, head_dim 128 | this PR + Turbo #556 | 6987.1 | 853.6 | 37,518 | 224 GB | 11.33373 | +| + `turbo_fused_grouped_gemm` (mean of 2 runs) | this PR | 6817.7 | 874.9 | 38,451 | 225 GB | 11.33412 | +| + `turbo_fp8_permute` (mean of 4 runs) | this PR + Turbo #555 | 6507.6 | 916.7 | 40,283 | 223 GB | 11.33404 | + +From the FP8 tensorwise config to the last row: −26.8% ms/iter, +36.7% tokens/s/GPU. Loss differences are within run-to-run noise (about 4e-4 at iteration 20). This branch, rebased on main 0a68e5cd, with the switches on: + +| run | ms/iter | tokens/s/GPU | loss@20 | +|---|---|---|---| +| switches on the command line | 6518.7 | 40,214 | 11.33415 | +| this config alone | 6410.6 | 40,892 | 11.33433 | +| this config alone | 6459.6 | 40,582 | 11.33445 | + +**Known issue: intermittent NaN.** +- One more config-only run of this branch hit a NaN forward loss on rank 4 at iteration 3. Its parsed arguments match a passing run exactly. +- Across this work, 1 of 13 20-iteration runs at `turbo_deepep_num_cu 160` hit this, as did the only run at 192 (rank 2, iteration 5). +- None of the 5 runs at 80 CUs did. Those runs also predate the fused paths, though, so the cause is not isolated yet. + +## Changes + +- `fix(turbo): make enable_turbo_attention_float8 runnable again` + - `primus/backends/megatron/core/extensions/primus_turbo.py`: + - Pass `sink=` only when a sink tensor exists; `flash_attn_fp8_func` has no `sink` parameter. + - Make q/k/v bshd-contiguous on the FP8 path. `flash_attn_fp8_func` permutes sbhd storage as if it were the logical layout (fixed in Turbo by `fix/fp8-attn-strided-layout`). +- `perf(qwen3): tune Qwen3-30B-A3B FP8 MI355X config to Turbo tensorwise` + - `examples/megatron/configs/MI355X/qwen3_30B_A3B-FP8-pretrain.yaml`: `fp8_recipe: tensorwise`, `use_turbo_gemm` / `use_turbo_grouped_gemm`, and the TE grouped-MLP path instead of the legacy grouped GEMM. +- `perf(megatron): enable the fused packed-QKV RMSNorm + RoPE patch for head_dim 128` + - `primus/backends/megatron/patches/turbo/qk_rmsnorm_rope_patches.py`: + - Accept any head_dim in Turbo's `QK_RMSNORM_ROPE_HEAD_DIMS`, falling back to `(64,)` on older Turbo. + - Accept TE RMSNorm as well as `PrimusTurboRMSNorm`. The fused path only reads weight / eps; `zero_centered_gamma` is still rejected. + - Fused QKV split + Q/K RMSNorm + RoPE at the Qwen3 attention shape: fwd+bwd 1243.6 → 369.9 µs (3.36×). +- `perf(megatron): hand the Turbo fused grouped MLP FP8 tokens from the DeepEP permute` + - New opt-in flag `turbo_fp8_permute` (default false in `primus_turbo.yaml`). + - Under Turbo FP8 tensorwise current scaling, `PrimusTurboDeepEPTokenDispatcher` passes `quantize_dtype` to `_post_dispatch` and `grad_quantize_dtype` to `_pre_combine`. The dtypes follow the FP8 format (HYBRID: e4m3 input, e5m2 gradient), and results are bit-identical to the default path. Any other recipe, or a layer with Turbo FP8 off, keeps the bf16 path. + - Argument validation requires `enable_primus_turbo`, `use_turbo_deepep` and `turbo_fused_grouped_gemm`, and rejects selective recompute of `moe_act`. + - Dispatcher kernel time per call at the Qwen3 EP8 shape: forward 645 → 167 µs, backward 687 → 174 µs. +- `perf(qwen3): turn on the Turbo fusions in the Qwen3-30B-A3B FP8 MI355X config` + - Top-level `env: PRIMUS_FUSED_QK_RMSNORM_ROPE: "1"`. + - `turbo_fused_grouped_gemm`, `use_turbo_fused_act_with_probs` and `turbo_fp8_permute` set to true. + - `turbo_deepep_num_cu` 80 → 160. 192 gave a NaN loss on one rank at iteration 5; 256 is slower (7028 ms). + +The force-"even" routing fix that this work also needed is already on main (#1243), so it is not part of this PR. + +## Tests + +- New in `tests/unit_tests/backends/megatron/test_rocm_arg_validation.py`: `validate_turbo_fp8_permute` (each required flag, selective recompute of `moe_act`, disabled no-op). +- New `tests/unit_tests/backends/megatron/test_turbo_fp8_permute_dtypes.py`: `_fp8_permute_dtypes` returns (e4m3, e5m2) for HYBRID and (e4m3, e4m3) for E4M3 tensorwise, and (None, None) with the flag off, Turbo FP8 off or blockwise scaling. +- `test_qk_rmsnorm_rope_patches.py`, `test_rocm_arg_validation.py`, `test_validate_args_patches.py`, `test_router_force_even_routing.py`, `test_turbo_fp8_permute_dtypes.py`: all pass on MI355X. +- End-to-end runs above. From cbff238790d769e95ee510b55ee96c6b4f77d4aa Mon Sep 17 00:00:00 2001 From: sim Date: Sat, 10 Oct 2026 15:47:04 +0000 Subject: [PATCH 08/14] docs(agent): add the round log (was gitignored), all e2e runs, container image --- .../qwen3_fp8_tw_gfx950_20261010/README.md | 6 +- .../e2e/results.md | 37 +++++++++ .../logs/optimize.md | 81 +++++++++++++++++++ 3 files changed, 122 insertions(+), 2 deletions(-) create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/results.md create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/logs/optimize.md diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md index 91593b5e2..bc88f5153 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -1,8 +1,8 @@ # Qwen3-30B-A3B FP8 tensorwise on 8x MI355X: Turbo kernel campaign (handoff) -Task: [prompt.md](prompt.md). Log of every round: [logs/optimize.md](logs/optimize.md). +Task: [prompt.md](prompt.md). Log of every round: [logs/optimize.md](logs/optimize.md). Every end-to-end run: [e2e/results.md](e2e/results.md). -Status (2026-10-10): the 40,000 tokens/s/GPU target is reached. The baseline was 29,474 tokens/s/GPU and the best stack runs at 40,283 (mean of 4 runs). +Status (2026-10-10): the 40,000 tokens/s/GPU target is reached. The baseline was 29,474 tokens/s/GPU and the best stack runs at 40,283 (mean of 4 runs). One open correctness issue: an intermittent NaN forward loss (see open items). ## End-to-end @@ -68,6 +68,8 @@ It holds five commits: the FP8 attention fix, the tensorwise config, the head_di All of these merges are clean as of 2026-10-10. +Container: every build, test and run used the image `tasimage/primus:pr-1236` (container `xiaoming-dev`, started with `--network=host --ipc=host --privileged`, `/dev/kfd`, `/dev/dri` and `/dev/infiniband`, home directory mounted). Run everything with `ulimit -c 0`, because core dumps fill the disk. + Turbo: ```bash diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/results.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/results.md new file mode 100644 index 000000000..19ca8a293 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/results.md @@ -0,0 +1,37 @@ +# All end-to-end runs + +Qwen3-30B-A3B FP8 tensorwise pretrain on 8x MI355X (EP8, MBS 8, GBS 512, seq 4096, even routing). +`parse_e2e.py` output: mean and standard deviation of iterations 11-20, tokens/s/GPU = 262144 / (s/iter), +max-rank ROCm memory. The raw logs were in the container's `/tmp/q3_e2e_.log` on the old machine. +They are not kept because they contain internal host names. + +`cu` is `turbo_deepep_num_cu`. "Fused act" is `use_turbo_fused_act_with_probs`. + +| run | what | cu | ms/iter | sd | TFLOP/s/GPU | tokens/s/GPU | max mem GB | loss@20 | +|---|---|---|---|---|---|---|---|---| +| r0_base | installed Turbo e2d9f1d7, fused act | 80 | 7915.6 | 47.8 | 753.5 | 33,117 | 222.7 | 11.33435 | +| rb_base | Turbo 1103b2df, fused act | 80 | 7925.4 | 8.5 | 752.5 | 33,077 | 222.3 | 11.33427 | +| b1_quant | + flat tensorwise quant kernel | 80 | 7870.9 | 20.9 | 757.7 | 33,305 | 222.3 | 11.33415 | +| b2_layout | + per-token `get_dispatch_layout` | 80 | 7774.2 | 63.2 | 767.2 | 33,720 | 222.3 | 11.33431 | +| r1_turbo | + HIP permute in the dispatcher (turbo-opt) | 80 | 7643.5 | 12.3 | 780.3 | 34,296 | 222.2 | 11.33423 | +| r2_cu160 | + `turbo_deepep_num_cu 160` | 160 | 7353.4 | 14.5 | 811.1 | 35,649 | 222.6 | 11.33456 | +| r3_qknorm | + fused qk RMSNorm + RoPE, head_dim 128 | 160 | 6987.1 | 22.0 | 853.6 | 37,518 | 224.0 | 11.33373 | +| s2_syncfree2 | r3 + sync-free MoE stage 2 | 160 | 7287.8 | 43.6 | 818.4 | 35,970 | 227.6 | 11.33385 | +| c192 | r3 with 192 CUs | 192 | NaN forward loss, rank 2, iteration 5 | | | | | | +| c256 | r3 with 256 CUs | 256 | 7028.3 | 41.8 | 848.6 | 37,299 | 224.5 | 11.33399 | +| p3_profile | r3 under the profiler (timings not representative) | 160 | 28309.8 | 21272.2 | 483.9 | 9,260 | 224.0 | 11.44972 | +| m1_mega_mxfp8 | MegaMoE mxfp8 | 160 | OOM at MBS 8 | | | | | | +| m2_mega_bf16 | MegaMoE bf16 | 160 | OOM at MBS 8 | | | | | | +| f1_fused_mlp | r3 + `turbo_fused_grouped_gemm` | 160 | 6820.2 | 34.9 | 874.5 | 38,437 | 224.5 | 11.33412 | +| f1b_fused_mlp | repeat of f1 | 160 | 6815.1 | 28.0 | 875.2 | 38,465 | 224.5 | 11.33421 | +| f2_fused_mlp_nopad | f1 + `turbo_grouped_gemm_without_padding` | 160 | argument validation exit (needs padding off; sync-free stage 1 forces it on) | | | | | | +| g1_fp8perm | f1 + `turbo_fp8_permute` (pre-padding-fix build, stateful Primus) | 160 | 6442.4 | 17.2 | 925.8 | 40,690 | 223.3 | 11.33404 | +| g1b_fp8perm_stateful | g1's Primus code, rebuilt turbo-opt (padding fix) | 160 | 6515.0 | 95.3 | 915.6 | 40,237 | 223.3 | 11.33411 | +| g2_fp8perm_stateless | stateless Primus dispatcher | 160 | 6549.1 | 136.4 | 911.1 | 40,028 | 223.3 | 11.33445 | +| g3_fp8perm_stateless | repeat of g2 | 160 | 6523.8 | 70.9 | 914.3 | 40,183 | 223.3 | 11.33420 | +| pr1 | Primus `perf/qwen3-30b-a3b-fp8-turbo`, switches on the command line | 160 | 6518.7 | 86.1 | 915.1 | 40,214 | 223.3 | 11.33415 | +| cfg1 | Primus `perf/qwen3-30b-a3b-fp8-turbo`, config only | 160 | NaN forward loss, rank 4, iteration 3 | | | | | | +| cfg2 | repeat of cfg1 | 160 | 6410.6 | 20.5 | 930.4 | 40,892 | 223.3 | 11.33433 | +| cfg3 | repeat of cfg1 | 160 | 6459.6 | 45.4 | 923.3 | 40,582 | 223.3 | 11.33445 | + +The cfg1 and pr1 runs parsed the same 847 arguments, and their wgrad autotune winners match. diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/logs/optimize.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/logs/optimize.md new file mode 100644 index 000000000..947910724 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/logs/optimize.md @@ -0,0 +1,81 @@ +# Qwen3-30B-A3B FP8 tensorwise pretrain on 8x MI355X — Turbo kernel campaign + +Baseline (prompt): 8894 ms/iter, 670.6 TFLOP/s/GPU, Turbo e2d9f1d7, Primus perf/qwen3-30b-a3b-tuning, +examples/megatron/configs/MI355X/qwen3_30B_A3B-FP8-pretrain.yaml. e2e: even routing, MBS 8, GBS 512, 20 iters. +All PR branches are based on Turbo 1103b2df (post-#455). + +## Item 1 — tensorwise FP8 quant chain +- Branch perf/quantization/tensorwise-fp8-qwen3 (1194eaa0): flat no-pad kernel for rows < 4 KB. +- 262144x768 quant: 156.5 -> 95.3 us (-39%); 262144x1536 -6%; larger K unchanged (row kernel). Byte-identical. +- Tests: test_quantize_fp8_tensorwise_no_pad_rows; tensorwise GEMM / grouped GEMM FP8 shards ~1400 passed. +- Not done: dense fused quant+transpose, amax fused into permute. + +## Item 2 — FlyDSL grouped GEMM tensorwise +- bench_grouped_gemm_qwen3.py (even / uniform / skew / hot). No kernel change committed. + +## Item 3 — DeepEP +- Branch perf/deep_ep/dispatch-layout-per-token (11b6d998): get_dispatch_layout 340 -> 44-46 us, 384 calls/iter. +- num_cu sweep (even): 160 CUs best, dispatch 920 -> 612 us, combine 536 -> 484 us. Primus flag only. +- FP8 dispatch: not adopted. + +## Item 4 — attention bwd +- TE CK is best overall. FlyDSL bwd main kernel +26% slower. TE bwd per call: main 2680 us, dk_dv_reduce 216, + odo 152, fill 107, dq_shuffle 94. Turbo AITER fwd 2x slower (causal kernel without tile skipping). +- Recommendation: keep use_turbo_attention false. + +## Item 5 — qk RMSNorm + RoPE head_dim 128 +- Turbo perf/flydsl/qk-rmsnorm-rope-hd128 (a895fdea): fwd+bwd 1243.6 -> 369.9 us (3.36x). +- Primus perf/qwen3-qk-rmsnorm-rope-hd128 (fb63615f): patch accepts head_dim 128 and TE RMSNorm. +- gfx950 hazard: dwordx4 buffer_store with SGPR soffset + VALU write of data VGPRs is not padded by LLVM. + +## Item 6 — permute / unpermute +- #455 TRITON default slower than HIP at every measured shape (Qwen3 even permute fwd 390 vs 255 us). +- TURBO kernel bounds rows by the routed-token count: routed tokens must come first. DeepEP guarantees it, + arbitrary callers do not, so the global default stays TRITON. +- Branch perf/moe/permute-default-hip (893d22f4): DeepEPTokenDispatcher(permute_backend=TURBO). + +## e2e attribution runs +| run | Turbo | flags | ms/iter | +|---|---|---|---| +| r0_base | installed e2d9f1d7 | fused_act_with_probs | 7835-7867 (iters 19-20) | +| rb_base | 1103b2df | fused act | 7925.4 | +| b1_quant | + quant | fused act | 7870.9 | +| b2_layout | + layout | fused act | 7774.2 | +| r1_turbo | turbo-opt (+ permute) | fused act | 7643.5 | +| r2_cu160 | turbo-opt | + num_cu 160 | 7353.4 | +| r3_qknorm | turbo-opt | + fused qk-norm/RoPE | 6987.1 (37,518 tok/s/GPU) | + +## Round 2 — toward 40,000 tok/s/GPU (<= 6553.6 ms/iter) +- m1 MegaMoE mxfp8 / m2 bf16: OOM at MBS 8 (264 / 261 GB allocated when the worst-case dispatch pool + allocation fails); MegaMoE pools are sized for the worst case and kept per layer for backward. +- s2 sync-free stage 2: 7287.8 ms (+300 ms vs R3): num_worst_tokens makes every post-dispatch step scan + 262144 receive rows instead of 32768 tokens. Keep stage 1. +- c192 (turbo_deepep_num_cu 192): NaN forward loss on rank 2 at iteration 5. DeepEP bug at high channel counts. +- R3 profile (rank 0, one step): kernels 6941 ms of 6987 ms; attention 1583, grouped GEMM 1268, hipBLASLt 1003, + DeepEP 876, quant chain 739 (amax 253), permute/unpermute 346, SwiGLU 262. +- c256: 7028 ms (slower than 160 CUs). +- f1 turbo_fused_grouped_gemm True: 6820.2 ms (38,437 tok/s/GPU). f2 (+ turbo_grouped_gemm_without_padding) + exits: needs moe_router_padding_for_quantization False, which sync-free stage 1 forces True. +- g1 turbo_fp8_permute True: 6442.4 ms, 925.8 TFLOP/s, 40,690 tok/s/GPU, 223.3 GB, loss@20 11.33404. Target met. + Per call (32768 tokens, H 2048, 16 local experts, top-8, pad 16): fwd kernels 645 -> 167 us, bwd 687 -> 174 us. + Turbo perf/moe/fp8-permute-tensorwise (82cf4373, on top of the permute branch + the padding fix). +- Pre-existing TURBO permute bug found by the new worst-token + pad dispatcher test: padding rows live after all + max_num_dispatched_tokens rows of row_id_map, but the kernel read them right after the routed count, so with an + unrouted tail (num_worst_tokens > 0) padding rows kept torch.empty garbage (bf16 path too). Fixed in + fix/moe/permute-padding-rows (941b9525); regression test fails on main, passes with the fix. +- TURBO unpermute leaves rows past the routed count unwritten (harmless: DeepEP combine reads only received rows). +- Primus perf/megatron/turbo-fp8-permute: turbo_fp8_permute flag, grad dtype recomputed in combine_preprocess + (no dispatcher state, safe under interleaved micro-batches). Re-validated as g2. + +## Round 3 — PRs and one Primus PR +- Turbo PRs: #553 merged (padding fix); #554 permute default, #555 FP8 permute (stacked on #554), #556 rope + head_dim 128, #557 dispatch layout, #558 quant, and fix/fp8-attn-strided-layout opened. All pass pre-commit. +- Primus: one branch perf/qwen3-30b-a3b-fp8-turbo on main 0a68e5cd (bef41e8a). The force-even router fix was + already on main (#1243) and is dropped. Added unit tests for validate_turbo_fp8_permute and _fp8_permute_dtypes. + The last commit turns every switch on in the Qwen3 FP8 config (env PRIMUS_FUSED_QK_RMSNORM_ROPE=1 included). +- e2e on that branch: pr1 (switches on the CLI) 6518.7 ms; cfg2 / cfg3 (config only) 6410.6 / 6459.6 ms + (40,892 / 40,582 tok/s/GPU). +- cfg1 (config only) died at iteration 3: NaN forward loss on rank 4. Arguments identical to pr1. At 160 CUs, + 1 of 13 20-iteration runs hit it; c192 hit it too. Not isolated; next step is a DeepEP intranode + dispatch/combine stress loop with exact checks at num_sms 80 / 160 / 192. +- All runs: e2e/results.md. From d70a72feceea9a7d9563879acedb613cbfcf07de Mon Sep 17 00:00:00 2001 From: xmpeng-dev Date: Sun, 11 Oct 2026 05:26:58 +0000 Subject: [PATCH 09/14] docs(agent): #554 merged; #555 rebased to cf630aad --- agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md index bc88f5153..76506c6f5 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -48,7 +48,7 @@ Turbo (`AMD-AGI/Primus-Turbo`): |---|---|---|---| | `fix/moe/permute-padding-rows` | f1a59fb9 | #553 merged | TURBO permute padding rows with worst-case buffers + pad | | `perf/moe/permute-default-hip` | 30b6cc7c | #554 open | body: [pr_bodies/perf-moe-permute-default-hip.md](pr_bodies/perf-moe-permute-default-hip.md) | -| `perf/moe/fp8-permute-tensorwise` | f4485b88 | #555 open | on main + #554's commit; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | +| `perf/moe/fp8-permute-tensorwise` | cf630aad | #555 open | rebased onto main after #554 squash; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | | `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | #558 open | independent; body: [pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md](pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md) | | `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | #557 open | independent; body: [pr_bodies/perf-deep-ep-dispatch-layout-per-token.md](pr_bodies/perf-deep-ep-dispatch-layout-per-token.md) | | `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | #556 open | independent; needed by the Primus qk-norm patch; body: [pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md](pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md) | From 726c5a7b1310991305b76cd7f2422776d0aace2c Mon Sep 17 00:00:00 2001 From: xmpeng-dev Date: Sun, 11 Oct 2026 05:27:23 +0000 Subject: [PATCH 10/14] docs(agent): mark #554 merged; drop stale #555 rebase note --- agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md index 76506c6f5..ca92d25c3 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -35,7 +35,7 @@ On the Primus branch `perf/qwen3-30b-a3b-fp8-turbo`, `examples/megatron/configs/ On that branch, the config alone gave 6410.6 and 6459.6 ms/iter (40,892 / 40,582 tokens/s/GPU). One more config-only run hit the intermittent NaN (see open items). Code it needs: -- Turbo: all five Turbo branches below. Building on `perf/moe/fp8-permute-tensorwise` also brings in the permute-default change. +- Turbo: all five Turbo branches below. `#554` (HIP permute default) is on main; `#555` is rebased onto it. - Primus: `perf/qwen3-30b-a3b-fp8-turbo`. See [e2e/e2e_run_cfg.sh](e2e/e2e_run_cfg.sh) (config only) and [e2e/env_turbo_opt_cfg.sh](e2e/env_turbo_opt_cfg.sh). The older [e2e/e2e_run.sh](e2e/e2e_run.sh) passes the switches on the command line. All paths are for the old machine. @@ -47,14 +47,14 @@ Turbo (`AMD-AGI/Primus-Turbo`): | branch | head | PR | notes | |---|---|---|---| | `fix/moe/permute-padding-rows` | f1a59fb9 | #553 merged | TURBO permute padding rows with worst-case buffers + pad | -| `perf/moe/permute-default-hip` | 30b6cc7c | #554 open | body: [pr_bodies/perf-moe-permute-default-hip.md](pr_bodies/perf-moe-permute-default-hip.md) | +| `perf/moe/permute-default-hip` | 30b6cc7c | #554 merged | body: [pr_bodies/perf-moe-permute-default-hip.md](pr_bodies/perf-moe-permute-default-hip.md) | | `perf/moe/fp8-permute-tensorwise` | cf630aad | #555 open | rebased onto main after #554 squash; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | | `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | #558 open | independent; body: [pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md](pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md) | | `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | #557 open | independent; body: [pr_bodies/perf-deep-ep-dispatch-layout-per-token.md](pr_bodies/perf-deep-ep-dispatch-layout-per-token.md) | | `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | #556 open | independent; needed by the Primus qk-norm patch; body: [pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md](pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md) | | `fix/fp8-attn-strided-layout` | c825b13e | opening | FP8 attention with sbhd/bhsd storage; body: [pr_bodies/fix-fp8-attn-strided-layout.md](pr_bodies/fix-fp8-attn-strided-layout.md) | -`perf/moe/fp8-permute-tensorwise` is main (with #553) plus #554's commit plus the FP8 permute commit, whose message quotes the 4-run mean (-4.55%). The pre-rebase history is kept locally as `backup/fp8-permute-pre-rebase`. Once #554 merges, rebase the branch onto main and drop #554's commit. +`perf/moe/fp8-permute-tensorwise` is now one commit on main (`cf630aad`); rebased after #554's squash merge. Primus (`AMD-AGI/Primus`): From f6cd6419ac7dafb2824bbc8156f450a95440a5e5 Mon Sep 17 00:00:00 2001 From: xmpeng-dev Date: Sun, 11 Oct 2026 06:50:02 +0000 Subject: [PATCH 11/14] =?UTF-8?q?docs(agent):=20#556=20cross-model=20micro?= =?UTF-8?q?bench=20=E2=80=94=20D=3D64=20no=20regression?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../bench_qk_rmsnorm_rope_ab.py | 98 +++++++++++++++++++ .../logs/optimize.md | 26 +++++ 2 files changed, 124 insertions(+) create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_qk_rmsnorm_rope_ab.py diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_qk_rmsnorm_rope_ab.py b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_qk_rmsnorm_rope_ab.py new file mode 100644 index 000000000..c7c748be0 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/bench_qk_rmsnorm_rope_ab.py @@ -0,0 +1,98 @@ +"""Compare fused QKV RMSNorm+RoPE: baseline Turbo (D=64 only) vs #556 (D=64+128). + +Shapes (S,B,NG,NPG,D): + gpt_oss_20B : 4096,8,8,8,64 — the existing fused path (PRIMUS_FUSED_QK_RMSNORM_ROPE) + qwen3_30B : 4096,8,4,8,128 — newly enabled by #556 + minimax_m3 : 4096,8,4,16,128 — D=128, NPG=16 (64 heads / 4 groups) + lfm2_8B : 4096,8,8,4,128 — D=128, NPG=4 (32 heads / 8 groups) + glm5 : 4096,4,1,64,128 — D=128; no GQA in yaml (treat as NG=1,NPG=64); B=4 to fit + +Usage (in container): + PYTHONPATH= FLYDSL_RUNTIME_CACHE_DIR=/tmp/fly_ \ + python bench_qk_rmsnorm_rope_ab.py [tag] [shape_name ...] +""" + +from __future__ import annotations + +import os +import sys + +import torch + +from primus_turbo.pytorch.ops.rope import fused_qkv_rmsnorm_rope + +SHAPES = { + "gpt_oss_20B": (4096, 8, 8, 8, 64), + "qwen3_30B": (4096, 8, 4, 8, 128), + "minimax_m3": (4096, 8, 4, 16, 128), + "lfm2_8B": (4096, 8, 8, 4, 128), + # B*NG must be a multiple of rows_per_wave (8 at D=128); B=8 keeps that. + "glm5": (4096, 8, 1, 64, 128), +} + +EPS = 1e-6 +tag = sys.argv[1] if len(sys.argv) > 1 else "run" +names = sys.argv[2:] or list(SHAPES) + + +def bench(fn, iters=40, warmup=8): + for _ in range(warmup): + fn() + s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True) + torch.cuda.synchronize() + s.record() + for _ in range(iters): + fn() + e.record() + torch.cuda.synchronize() + return s.elapsed_time(e) * 1e3 / iters + + +def run_shape(name: str): + S, B, NG, NPG, D = SHAPES[name] + try: + from primus_turbo.flydsl.rope.qk_rmsnorm_rope_kernel import QK_RMSNORM_ROPE_HEAD_DIMS + + supported = QK_RMSNORM_ROPE_HEAD_DIMS + except ImportError: + from primus_turbo.flydsl.rope.qk_rmsnorm_rope_kernel import QK_RMSNORM_ROPE_HEAD_DIM + + supported = (QK_RMSNORM_ROPE_HEAD_DIM,) + if D not in supported: + print(f"{tag:8s} {name:12s} D={D} unsupported (HEAD_DIMS={supported})") + return + + torch.manual_seed(0) + split = [NPG * D, D, D] + qkv = torch.randn(S, B, NG, sum(split), device="cuda", dtype=torch.bfloat16, requires_grad=True) + half = torch.randn(S, D // 2, device="cuda") + freqs = torch.cat((half, half), -1).reshape(S, 1, 1, D).contiguous() + q_gamma = torch.randn(D, device="cuda", dtype=torch.bfloat16, requires_grad=True) + k_gamma = torch.randn(D, device="cuda", dtype=torch.bfloat16, requires_grad=True) + + def fused(): + return fused_qkv_rmsnorm_rope(qkv, q_gamma, k_gamma, freqs, split, EPS) + + with torch.no_grad(): + outs = fused() + grads = tuple(torch.randn_like(t) for t in outs) + + def fwd_bwd(): + o = fused() + torch.autograd.backward(o, grads) + qkv.grad = None + q_gamma.grad = None + k_gamma.grad = None + + with torch.no_grad(): + t_f = bench(fused) + t_fb = bench(fwd_bwd) + print( + f"{tag:8s} {name:12s} S={S} B={B} NG={NG} NPG={NPG} D={D} " + f"fwd {t_f:7.1f} us fwd+bwd {t_fb:7.1f} us bwd {t_fb - t_f:7.1f} us" + ) + + +print(f"# tag={tag} pythonpath={os.environ.get('PYTHONPATH','').split(':')[0]}") +for n in names: + run_shape(n) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/logs/optimize.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/logs/optimize.md index 947910724..68e402ada 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/logs/optimize.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/logs/optimize.md @@ -79,3 +79,29 @@ All PR branches are based on Turbo 1103b2df (post-#455). 1 of 13 20-iteration runs hit it; c192 hit it too. Not isolated; next step is a DeepEP intranode dispatch/combine stress loop with exact checks at num_sms 80 / 160 / 192. - All runs: e2e/results.md. + +## #556 cross-model check (2026-10-11) +Microbench fused QKV RMSNorm+RoPE, main (turbo-base / D=64 only) vs #556 (turbo-pr-rope). +S=4096, MBS matching typical EP8 configs. Unit tests: 20 passed. + +GPT-OSS-20B shape (D=64, the existing fused path) — no regression: +| tree | fwd us | fwd+bwd us | +|---|---|---| +| main run1 | 140.5 | 361.2 | +| main run2 | 142.9 | 363.2 | +| #556 run1 | 140.7 | 359.8 | +| #556 run2 | 140.9 | 360.1 | + +Other D=128 shapes — only reachable after Primus enables the patch for head_dim 128 +(Turbo #556 alone does not change those training paths). fused vs unfused on #556: +| model shape | D | fused fwd+bwd | unfused fwd+bwd | speedup | +|---|---|---|---|---| +| gpt_oss_20B NG=8 NPG=8 | 64 | 367 | 2044 | 5.56x | +| qwen3_30B NG=4 NPG=8 | 128 | 357 | (prior) 1244 | ~3.5x | +| minimax_m3 NG=4 NPG=16 | 128 | 626 | 2228 | 3.56x | +| lfm2_8B NG=8 NPG=4 | 128 | 419 | 1395 | 3.33x | +| glm5-like NG=1 NPG=64 | 128 | 571 | 2112 | 3.70x | + +Models that do not hit this kernel with or without #556: no qk_layernorm, MLA +(DeepSeek), head_dim not in {64,128} (e.g. Qwen3.5 kv_channels 256), or +PRIMUS_FUSED_QK_RMSNORM_ROPE unset. From ddc475dff552148cbccc6b15ba3d451d382652d0 Mon Sep 17 00:00:00 2001 From: xmpeng-dev Date: Sun, 11 Oct 2026 06:50:04 +0000 Subject: [PATCH 12/14] docs(agent): mark #555 merged; note #556 D=64 check --- agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md index ca92d25c3..897e13691 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -48,10 +48,10 @@ Turbo (`AMD-AGI/Primus-Turbo`): |---|---|---|---| | `fix/moe/permute-padding-rows` | f1a59fb9 | #553 merged | TURBO permute padding rows with worst-case buffers + pad | | `perf/moe/permute-default-hip` | 30b6cc7c | #554 merged | body: [pr_bodies/perf-moe-permute-default-hip.md](pr_bodies/perf-moe-permute-default-hip.md) | -| `perf/moe/fp8-permute-tensorwise` | cf630aad | #555 open | rebased onto main after #554 squash; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | +| `perf/moe/fp8-permute-tensorwise` | cf630aad | #555 merged | rebased onto main after #554 squash; body: [pr_bodies/perf-moe-fp8-permute-tensorwise.md](pr_bodies/perf-moe-fp8-permute-tensorwise.md) | | `perf/quantization/tensorwise-fp8-qwen3` | 6351abf4 | #558 open | independent; body: [pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md](pr_bodies/perf-quantization-tensorwise-fp8-qwen3.md) | | `perf/deep_ep/dispatch-layout-per-token` | 0931ba1e | #557 open | independent; body: [pr_bodies/perf-deep-ep-dispatch-layout-per-token.md](pr_bodies/perf-deep-ep-dispatch-layout-per-token.md) | -| `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | #556 open | independent; needed by the Primus qk-norm patch; body: [pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md](pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md) | +| `perf/flydsl/qk-rmsnorm-rope-hd128` | 1b920afe | #556 open (D=64 no reg) | independent; needed by the Primus qk-norm patch; body: [pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md](pr_bodies/perf-flydsl-qk-rmsnorm-rope-hd128.md) | | `fix/fp8-attn-strided-layout` | c825b13e | opening | FP8 attention with sbhd/bhsd storage; body: [pr_bodies/fix-fp8-attn-strided-layout.md](pr_bodies/fix-fp8-attn-strided-layout.md) | `perf/moe/fp8-permute-tensorwise` is now one commit on main (`cf630aad`); rebased after #554's squash merge. From cdee943f3b2ab35ca217faa6b26e0218f358fcdf Mon Sep 17 00:00:00 2001 From: xmpeng-dev Date: Sun, 11 Oct 2026 09:06:47 +0000 Subject: [PATCH 13/14] notes(qwen3): document DeepEP full-fence NaN mitigation Record the Primus DISABLE_CHEAP_FENCE config fix and the Turbo high-CU fence default branch for the intermittent forward-loss NaN. Co-authored-by: Cursor --- .../qwen3_fp8_tw_gfx950_20261010/README.md | 10 ++++------ .../fix-deep-ep-cheap-fence-high-cu.md | 18 ++++++++++++++++++ .../primus-perf-qwen3-30b-a3b-fp8-turbo.md | 7 +++---- 3 files changed, 25 insertions(+), 10 deletions(-) create mode 100644 agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-deep-ep-cheap-fence-high-cu.md diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md index 897e13691..54d5d7ba4 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/README.md @@ -114,12 +114,10 @@ Run Python from the campaign directory with `PYTHONPATH=`. Other ## Open items -- **Intermittent NaN forward loss.** Run `cfg1` of the final config (`perf/qwen3-30b-a3b-fp8-turbo`, config only) died at iteration 3 with `found NaN in local forward loss calculation` on rank 4. The per-rank traceback is in `output/.../logs/pre_trainer/rank-4/error.log`. - - Its 847 parsed arguments match the passing run `pr1` exactly, and the wgrad autotune winners match too. - - At `turbo_deepep_num_cu 160`, 1 of 13 runs hit it. The only run at 192 hit it (rank 2, iteration 5). The only run at 256 did not. - - None of the 5 runs at 80 CUs hit it, but those predate the fused qk-norm, fused grouped MLP and FP8 permute, so the cause is not isolated. - - The `turbo-opt` build includes the #553 padding fix. - - Next step: a DeepEP intranode dispatch/combine stress loop with exact checks at `num_sms` 80 / 160 / 192. It runs much faster than a 1-in-13 end-to-end bisection. +- **Intermittent NaN forward loss (mitigated).** Run `cfg1` of the final config died at iteration 3 with `found NaN in local forward loss calculation` on rank 4. Rate: 1/13 at `turbo_deepep_num_cu 160`, 1/1 at 192; none at 80 (those also predated the fused paths). + - Idle DeepEP dispatch/combine stress at 80/160/192 (200 iters, exact checks) passed under both fences; concurrent-GEMM stress at 160 also stayed numerically clean until a launch-resource abort. + - Mitigation landed: Primus config sets `PRIMUS_TURBO_DEEPEP_DISABLE_CHEAP_FENCE=1` (full release/acquire fence; ~2–3%/step historically). Turbo branch `fix/deep_ep/cheap-fence-high-cu` defaults to the full fence when `num_sms > 80` unless the env forces otherwise (#539 only validated ≤80 CUs). + - e2e confirmation: `fence1`–`fence4` queue under `/tmp/q3_e2e_fence*.log`. - **Unexplained e2e jitter with FP8 permute.** The first run took 6442 ms; three later runs took 6515-6549 ms with more per-iteration jitter. These causes are ruled out: - the Primus change: a same-code rerun gave 6515 ms; - the padding fix: permute kernel times are unchanged; diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-deep-ep-cheap-fence-high-cu.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-deep-ep-cheap-fence-high-cu.md new file mode 100644 index 000000000..484813175 --- /dev/null +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/fix-deep-ep-cheap-fence-high-cu.md @@ -0,0 +1,18 @@ +# fix(deep_ep): default to the full fence above 80 CUs + +Branch: `fix/deep_ep/cheap-fence-high-cu` @ `394131b3` +Compare: https://github.com/AMD-AGI/Primus-Turbo/compare/main...fix/deep_ep/cheap-fence-high-cu?expand=1 + +## Summary + +- Default DeepEP intranode cheap fence only at `num_sms <= 80` (the #539 validated range). +- Above that, use the full release/acquire fence unless `PRIMUS_TURBO_DEEPEP_DISABLE_CHEAP_FENCE` forces a choice (`1` = full, `0` = cheap). +- Addresses intermittent NaN forward loss on the Qwen3-30B-A3B FP8 full-opt stack at `turbo_deepep_num_cu` 160/192 (~1/13 runs). +- Primus counterpart: `PRIMUS_TURBO_DEEPEP_DISABLE_CHEAP_FENCE=1` in the Qwen3 FP8 MI355X config (works with older Turbo that lacks this default). + +## Test plan + +- [ ] DeepEP intranode stress at 80 CUs still uses cheap fence when env unset +- [ ] At 160 CUs with env unset, dispatch/combine select the full fence +- [ ] `DISABLE=0` still forces cheap at 160 for A/B +- [ ] Qwen3 FP8 20-iter e2e at 160 CUs with full opt (multiple reps) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/primus-perf-qwen3-30b-a3b-fp8-turbo.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/primus-perf-qwen3-30b-a3b-fp8-turbo.md index d1c8faf3a..82d728c7d 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/primus-perf-qwen3-30b-a3b-fp8-turbo.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/pr_bodies/primus-perf-qwen3-30b-a3b-fp8-turbo.md @@ -36,10 +36,9 @@ From the FP8 tensorwise config to the last row: −26.8% ms/iter, +36.7% tokens/ | this config alone | 6410.6 | 40,892 | 11.33433 | | this config alone | 6459.6 | 40,582 | 11.33445 | -**Known issue: intermittent NaN.** -- One more config-only run of this branch hit a NaN forward loss on rank 4 at iteration 3. Its parsed arguments match a passing run exactly. -- Across this work, 1 of 13 20-iteration runs at `turbo_deepep_num_cu 160` hit this, as did the only run at 192 (rank 2, iteration 5). -- None of the 5 runs at 80 CUs did. Those runs also predate the fused paths, though, so the cause is not isolated yet. +**NaN mitigation.** +- Intermittent NaN forward loss at `turbo_deepep_num_cu` 160/192 (~1/13). The DeepEP cheap fence (#539) was only row-validated at ≤80 CUs. +- Config now sets `PRIMUS_TURBO_DEEPEP_DISABLE_CHEAP_FENCE=1` (full release/acquire fence). Pair with Turbo `fix/deep_ep/cheap-fence-high-cu`, which defaults to the full fence above 80 CUs. ## Changes From df2bb41a1462bf36e8ffe5bb3c5e1a9f298a9f1f Mon Sep 17 00:00:00 2001 From: xmpeng-dev Date: Sun, 11 Oct 2026 09:07:09 +0000 Subject: [PATCH 14/14] notes(qwen3): mark fence e2e validation as pending GPU availability Co-authored-by: Cursor --- agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/results.md | 3 +++ 1 file changed, 3 insertions(+) diff --git a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/results.md b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/results.md index 19ca8a293..5531b3f7e 100644 --- a/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/results.md +++ b/agent/workspace/qwen3_fp8_tw_gfx950_20261010/e2e/results.md @@ -34,4 +34,7 @@ They are not kept because they contain internal host names. | cfg2 | repeat of cfg1 | 160 | 6410.6 | 20.5 | 930.4 | 40,892 | 223.3 | 11.33433 | | cfg3 | repeat of cfg1 | 160 | 6459.6 | 45.4 | 923.3 | 40,582 | 223.3 | 11.33445 | +| fence1–4 | cfg + `PRIMUS_TURBO_DEEPEP_DISABLE_CHEAP_FENCE=1` | 160 | pending (waiting for GPUs; node shared with vLLM on 0–3) | | | | | | + + The cfg1 and pr1 runs parsed the same 847 arguments, and their wgrad autotune winners match.