Skip to content

[WS1][CUDA][Qwen3-Next] MoE route/combine contract with a batch-invariant grouped expert GEMM (RFC #428 stack 5/11) - #503

Open
fusheng-ji wants to merge 17 commits into
RL-Align:test-qwennextfrom
fusheng-ji:feat/ws1-qwen3-next-moe-route-combine
Open

fusheng-ji wants to merge 17 commits into
RL-Align:test-qwennextfrom
fusheng-ji:feat/ws1-qwen3-next-moe-route-combine

Conversation

@fusheng-ji

@fusheng-ji fusheng-ji commented Oct 9, 2026 •

Copy link
Copy Markdown
Contributor

Stack: #466 (merged) → #467 → #468 → #469 → this PR → #507 → #543 → #544 → #549 → #557 → benchmarks (#551), base test-qwennext; merges after #469. This PR's own diff. RFC #428 row moe_route_combine_contract, CUDA.

What

The routed-MoE contract for Qwen3-Next-80B-A3B (512 experts, top-10), per stage (full table in docs/operators/qwen3-next-moe-route-combine.md):

Stage Order and dtype Atomics
router BF16 x, W upcast; one pinned vLLM batch-invariant GEMM, IEEE FP32 none
selection argsort(logits, descending, stable)[:, :10]: score descending, ties to the lower expert id none
weights fixed-order softmax and renormalisation over a fixed pairwise tree, FP32 none
experts per (token, route) row: gate_up GEMM → SwiGLU in FP32 → BF16 → down GEMM; no routed weight in the kernel, no sum over routes none, one writer per row
combine e0·w0, then + e_k·w_k for k = 1..9 in route order, FP32; one BF16 cast none
backward per expert in ascending id; dW slices written once; dx summed per token in ascending expert order; route-weight gradient: the sum over H goes through fixed_order_row_sum (custom autograd function in combine_routes) none
  • rl_engine/models/qwen3_next/qwen3_next_forward.py: shared_linear, shared_router, stable_top10_routes, combine_routes, grouped_route_linear, shared_moe. Fails closed on CPU or ROCm tensors, on any vLLM other than 0.30.0, on VLLM_TRITON_USE_TD and on TF32 Triton dots.
  • rl_engine/models/qwen3_next/qwen3_next_tp_blocks.py: TP4MoE with the HF load/export of all 512 experts, a replicated router and shared-expert gate, and MCore TP attributes.
  • rl_engine/models/qwen3_next/qwen3_next_tp.py: Megatron-style copy/reduce autograd boundaries through RL-Kernel's fixed-order collective.
  • rl_engine/validation/common/tensor_identity.py: raw-bit identity (signed zero distinct, NaN/Inf rejected, dtype must match).
  • tools/validation/models/qwen3_next/qwen3_next_tp_moe_check.py: the TP4 gate on real layer-0 weights (torchrun, 4 GPUs).
  • combine_routes has a custom autograd function: the stock autograd backward summed grad * expert_output over H with torch.sum, whose reduction order depends on the batch shape, so the route-weight gradient (and dx through the router) stopped being batch-invariant from 16 rows up on H100 and B200. The function sums H with fixed_order_row_sum. New test: first 8 rows alone vs the same rows inside larger batches.
  • Layout as reviewed on [WS1][CUDA][Qwen3-Next] TP distributed logprob parity: fixed reduction order, selected-token logprob exact across NCCL settings (RFC #428 stack 9/11) #549: the prior-art runner and plot live in the benchmark PR ([WS1][CUDA][Qwen3-Next] Comparison benchmarks and plots for the Qwen3-Next stack: MoE, attention, logprob, EP (RFC #428 stack 11/11) #551), raw evidence is a release ZIP with a summary in reports/experiments/qwen3_next/.

Prior art & reuse decision

Decision: reuse vLLM's batch-invariant GEMM (router, backward) and its fused_moe Triton kernel (expert GEMMs); implement the FP32 routing, tie rule, combine, backward and TP4 boundaries. No surveyed implementation is batch-invariant and has a batch-invariant backward. The batch-invariant ones (vLLM BI=1, SGLang deterministic) are inference-only and route in BF16, which picks a different expert set than FP64 routing for 11 of 256 tokens; Megatron-core + TE, VIME's training path, routes in FP32 but changes routing, output, dx and dW with batch size.

Library Implementation (version / file) Forward BI Backward BI Tokens routed unlike FP64 Verdict
CUTLASS via FlashInfer cutlass_fused_moe (grouped GEMM) no (≥ 256 tokens) no backward 11 / 256 not BI; inference only
FlashInfer 0.6.18.post1 fused_moe.cutlass_fused_moe no (≥ 256) no backward 11 / 256 not BI
FlashAttention — — — — N/A: attention library, no MoE
vLLM 0.30.0 fused_topk + fused_experts, VLLM_BATCH_INVARIANT=0 no (≥ 256) no backward 11 / 256 not BI
vLLM 0.30.0 same, VLLM_BATCH_INVARIANT=1 yes no backward 11 / 256 BI, inference only → expert GEMM kernel reused
SGLang 0.5.21 Triton fused_moe, default / deterministic inference no / yes no backward 11 / 256 BI only in deterministic mode, inference only
Miles (radixark/miles@23d3d38) miles/kernels/moe/fused_experts.py (SGLang forward) + fused_experts_bwd.py, imported unmodified output equal over 20 launches, dx of probe rows 0 / 10 batch cells no, measured: over 20 identical launches dx, dW gate/up and dW down equal the first in 1 / 20 (BF16 tl.atomic_add); route-weight gradient 20 / 20 not measured backward not reproducible; 10 ms vs our 158 ms fwd+bwd (summary)
verl (volcengine/verl@e5a627c) no MoE kernel (searched verl/); trains through Megatron-core, HF/FSDP or VeOmni, rolls out through vLLM/SGLang — — — covered by those rows
Vime 0.3.2 Megatron-core MoELayer (no MoE kernel of its own) see Megatron row covered by Megatron row
Megatron-core 0.16 + TE 2.16 MoELayer, TE grouped GEMM, FP32 router, all-to-all (VIME's Qwen3-Next config) no (routes ≥ 256, output 1024) no (1024) 0 / 256 not BI
transformers 5.17.0 Qwen3NextExperts eager (verl FSDP path) yes no (1024) 11 / 256 backward not BI
RL-Kernel (this PR) shared_moe / TP4MoE yes yes 0 / 256

Batch invariance: 8 probe tokens alone vs first and last in 16-1024-token batches, forward (routes, output) and backward (dx, dW), bitwise. Miles was measured by importing its own kernels at the pinned commit (20 identical launches on the real layer-0 weights, bit for bit against the first; its forward uses SGLang's kernel with SGLang's deterministic tile).

Results (B200, H = 2048, 512 experts, top-10, width 512, BF16)

Qwen3-Next routed MoE prior art

RL-Kernel Megatron + TE vLLM BI=1 SGLang det. vLLM BI=0 SGLang default FlashInfer HF eager
max abs vs FP64 2.1e-3 2.2e-3 1.1e-1 1.1e-1 1.1e-1 1.1e-1 1.1e-1 1.1e-1
rel. L2 vs FP64 3.9e-3 4.6e-3 7.0e-2 7.0e-2 7.0e-2 7.0e-2 7.0e-2 7.0e-2
forward, 1 / 64 / 1024 / 4096 tokens (ms) 1.39 / 1.69 / 2.04 / 3.05 3.8 / 9.7 / 11.7 / 12.1 0.39 / 0.66 / 0.87 / 1.05 0.35 / 0.66 / 0.84 / 1.37 0.38 / 0.66 / 0.87 / 1.06 0.34 / 0.61 / 0.83 / 1.35 0.37 / 0.72 / 0.90 / 1.12 2.1 / 62 / 89 / 87
forward + backward, 64 / 1024 tokens (ms) 114 / 169 45 / 51 – – – – – 877 / 1284
  • The accuracy gap is routing: with BF16 router logits, 11 of 256 tokens select a different expert set than FP64 routing; with an FP32 router (this PR, Megatron-core) none do.
  • TP4, real layer-0 weights: all 12 cases pass on 4 ranks (HF round trip of all 512 experts, identical routes on every rank, full == chunked and reversed-order batches at 8/64/256/1024 tokens, training forward == no-grad forward, identical replicated router and shared-gate gradients).

Summary, tested commits and commands: reports/experiments/qwen3_next/qwen3-next-moe-route-b200.md; raw reports ZIP (SHA-256 9fcc7042…58cb0). Prior art measured at c58816b, TP4 gate at 2e950f3 (identical MoE code).

Status

  • Head e514d42f, on top of [WS1][CUDA][Qwen3-Next] C6: GDN decode-step goldens and recurrent/conv state contract (RFC #428 stack 4/4) #469 and d5253a5c. The last commit is the route-weight gradient fix above; it changes backward arithmetic only, forward is untouched. CI is green on this head: lint, unit-tests, docs, DCO and the bot review pass; provider and ws1-gtest skip because no self-hosted runner is attached.
  • Ready for review. No known acceptance failure.
  • Known: slower than vLLM's non-strict MoE forward (2.04 vs 0.87 ms at 1024 tokens); about 0.5 ms is the FP32 IEEE router GEMM the contract requires. The backward is a per-expert loop (169 ms forward + backward at 1024 tokens vs 51 ms for Megatron + TE, which is not batch-invariant); it is the next thing to speed up.
  • Rerun after aligning the stack with the reviewed layout (clean clone of the stack top 90c09164, 4 × B200): the standalone TP4 MoE gate passed 12/12 on all four ranks, and the layers stages of the EP1 and EP4 chains passed with it. (before the route-weight fix; B200 rerun on e514d42f below, see History)

History

  • 2026-10-26 (review): two OpenCodeReview rounds on the stack (the bot re-reviews every PR's whole diff after each push); every finding of these two rounds is answered in its own thread or comment, most as repeats of earlier answers or as not applicable with the code reference. Fixes in this PR: d53ac95e triggers the provider job on rl_engine/validation/common/**. The fixes of the PRs below it are merged in; new head 9800e1a4. Kernel arithmetic is unchanged, so the measured results stand.
  • 2026-10-26: merged test-qwennext after [WS1][CUDA][Qwen3-Next] C1: zero-centred RMSNorm references and CUDA kernel (RFC #428 stack 2/4) #467 landed (c382dd4b), through [WS1][CUDA][Qwen3-Next] C6: GDN decode-step goldens and recurrent/conv state contract (RFC #428 stack 4/4) #469's new head 27212b65, in 68240ec3. One conflict, in the provider workflow's pytest step: [WS1][CUDA][Qwen3-Next] C6: GDN decode-step goldens and recurrent/conv state contract (RFC #428 stack 4/4) #469's fail-closed skip check now also runs this PR's check modules.
  • 2026-11-14 (finding and fix, H100 and B200): on H100 the stock combine_routes backward gave a route-weight gradient that was not batch-invariant at 16, 32, 64, 128, 256 and 1024 rows (equal at 8 and 9). e514d42f fixes it with a fixed-order row sum. H100 A/B on 4 × H100: previous head d5253a5c 322 passed, probe unequal at T ≥ 16; tree ce3f499d (= e514d42f) 323 passed (the extra one is the new regression test), probe equal at every T, TP4 MoE gate 4/4 ranks on both. Output and expert gradients were equal on both. Same on B200 (4 × B200, same A/B and probe): d5253a5c 322 passed, probe unequal at T ≥ 16; e514d42f 323 passed, new test passed, probe equal at every T; TP4 MoE gate 12/12 cases on 4/4 ranks on both. So the break is not Hopper-specific, and the fix holds on both architectures. Raw data: hopper-qwen3-next/pr503/b200-e514d42f (TP4 gate on the index plus the single layer-0 shard of revision 9c7f2fbe, which is all the gate reads).

Tests

python setup.py build_ext --inplace
python -m pytest -q tests/validation/common/test_tensor_identity.py tests/models/qwen3_next/test_qwen3_next_forward_contract.py tests/models/qwen3_next/test_qwen3_next_tp_blocks.py
TRITON_F32_DEFAULT=ieee python -m pytest -q tests/models/qwen3_next/check_qwen3_next_forward.py      # needs vLLM 0.30.0
TRITON_F32_DEFAULT=ieee torchrun --nproc-per-node 4 tools/validation/models/qwen3_next/qwen3_next_tp_moe_check.py \
    --checkpoint <Qwen3-Next-80B-A3B-Instruct> --output <dir>
Check Result
CPU: tensor identity, forward contract, TP blocks 39 passed
layout check, pre-commit on all files, mkdocs build --strict clean
check_qwen3_next_forward.py (GEMM batch/chunk/reorder, route ties, combine order, MoE VJP vs FP64, route and output BI, grouped == per-expert bitwise) 23 passed
TP4 gate, 4 × B200 12 / 12 cases, 4 ranks
check_qwen3_next_forward.py -k combine_gradients (new: route-weight gradient, 8 rows alone vs inside larger batches), H100 and B200 1 passed on each
tests/models/qwen3_next/test_*.py + check_*.py, 4 × H100, tree ce3f499d (= e514d42f) 323 passed, 0 failed, 0 skipped
tests/models/qwen3_next/test_*.py + check_*.py, 4 × B200, e514d42f 323 passed, 0 failed, 0 skipped
TP4 gate, 4 × B200, e514d42f 12 / 12 cases, 4 ranks

CI: the CPU tests are in ci.yml; check_qwen3_next_forward.py imports vLLM and runs in the self-hosted Qwen3-Next-provider-GPU workflow with TRITON_F32_DEFAULT=ieee, which now also triggers on rl_engine/validation/models/qwen3_next/** and tools/validation/models/qwen3_next/**.

Scope

CUDA, TP4 × EP1 (all 512 experts on every rank). EP4 is in #557; ROCm and other TP sizes are separate claims. Model-level parity uses this block but belongs to full_model_chain (#543).

Results (H100 / Hopper, H = 2048, 512 experts, top-10, width 512, BF16)

Run at tree ce3f499d (= head e514d42f). 4 × NVIDIA H100, driver 580.173.02, CUDA 13.0, torch 2.13.0+cu130, triton 3.7.1, vLLM 0.30.0, transformers 5.17.0, VLLM_BATCH_INVARIANT=0. Fresh worktrees, fresh builds, clean tree before and after. Baseline run at the previous head d5253a5c. Raw data and report: hopper-qwen3-next/pr503/b960f8ce.

Check d5253a5c (before) e514d42f tree ce3f499d
tests/models/qwen3_next/test_*.py + check_*.py 322 passed 323 passed
TP4 MoE gate (4 ranks, real layer-0 weights) 4/4 passed 4/4 passed
route-weight gradient equal across batch size, T = 8 / 9 yes yes
route-weight gradient equal across batch size, T = 16, 32, 64, 128, 256, 1024 no yes
output and expert gradients equal across batch size yes yes

Qwen3-Next routed MoE prior art on H100

RL-Kernel Megatron + TE vLLM BI=1 vLLM BI=0 HF eager FlashInfer
forward batch invariant (routes, output) yes no yes yes yes yes
backward dW batch invariant yes no no backward no backward yes –
backward dx batch invariant yes no no backward no backward yes –
tokens routed unlike FP64 0 / 256 0 / 256 11 / 256 11 / 256 11 / 256 11 / 256
max abs vs FP64 1.7e-3 2.1e-3 9.6e-2 9.6e-2 9.6e-2 9.6e-2
forward 1 / 64 / 1024 / 4096 tokens (ms) 1.77 / 2.70 / 3.79 / 5.58 5.8 / 14.9 / 18.3 / 18.4 0.50 / 1.80 / 2.61 / 3.06 0.50 / 1.51 / 2.14 / 2.43 3.7 / 125 / 177 / 178 0.50 / 4.0 / 5.1 / 4.7
forward + backward 64 / 1024 tokens (ms) 231 / 353 70 / 82 – – 2188 / 3183 –

The prior-art figure and table above were measured before e514d42f (the RL-Kernel row records f9caee24) and were not rerun at e514d42f. FlashInfer sm90a JIT needed libnvrtc.so on LD_LIBRARY_PATH (symlink to libnvrtc.so.13). Non-RL-Kernel rows use the #551 runner at 77d59a59. Diagnostics for the original torch.sum behaviour: hopper-qwen3-next/pr503/diagnostics.

…rouped expert GEMM

RFC RL-Align#428 moe_route_combine_contract (CUDA).

- Routing: FP32 router through the pinned vLLM batch-invariant GEMM, a
  fixed-order softmax, stable descending argsort (ties go to the lower
  expert id) and FP32 renormalisation of the ten selected weights.
- Experts: per-route gate_up and down projections through vLLM 0.30.0's
  fused_moe Triton kernel with the batch-invariant tile passed explicitly
  (BLOCK 64/64/32, SPLIT_K=1); no routed-weight multiply and no sum inside
  the kernel, so every (token, route) row has one writer. Bitwise equal to
  the previous per-expert GEMM loop at TP4 and TP1 widths; fails closed if
  VLLM_TRITON_USE_TD is enabled.
- Combine: ten ordered FP32 multiply-adds and one BF16 cast; no atomics.
- Backward: per-expert pinned GEMMs into one dense gradient buffer; dx is
  summed in ascending expert order.
- TP4 block: HF load/export for all 512 experts, replicated router and
  shared-expert gate, Megatron-style copy/reduce autograd boundaries moved
  into qwen3_next_tp.py.
- Prior-art runner and plot comparing HF transformers, vLLM fused_moe
  (BI=0/1) and FlashInfer cutlass_fused_moe; TP4 real-weight gate script.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…nce (B200, 2e950f3)

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
@coderabbitai

coderabbitai Bot commented Oct 9, 2026 •

Copy link
Copy Markdown

Important

Review skipped

Auto reviews are disabled on base/target branches other than the default branch.

Please check the settings in the CodeRabbit UI or the .coderabbit.yaml file in this repository. To trigger a single review, invoke the @coderabbitai review command.

Run configuration

  • Source mode: Auto
  • Repository: Unavailable
  • Branch: Unavailable
  • File: No configuration file from this repository applied
  • Commit: Unavailable
  • Layers: defaults
  • Review profile: CHILL
  • Plan: Advanced
  • Run ID: 911f6787-4a78-4e55-b73a-1c7a19d28fdd

You can disable this status message by setting the reviews.review_status to false in the CodeRabbit configuration file.

Use the checkbox below for a quick retry:

  • Trigger review
  • Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

…r-art runner

- Megatron-core 0.16 MoELayer with Transformer Engine grouped GEMM and fused
  permute, configured as VIME runs Qwen3-Next (FP32 softmax router, top-10,
  all-to-all dispatcher), including its backward.
- SGLang 0.5.21 Triton fused_moe in default and deterministic-inference mode.
  sgl_kernel cannot load next to torch 2.13, so its two kernels on this path
  are replaced by SGLang's own Triton moe_sum_reduce and JIT
  moe_align_block_size; any other sgl_kernel call fails loudly.
- --only and --merge combine reports from different Python environments.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…ured (B200, c58816b)

Every library in the reuse table is now measured: Megatron-core 0.16 + TE 2.16 as VIME configures Qwen3-Next, and SGLang 0.5.21 in default and deterministic-inference mode. The TP4 gate evidence is unchanged (2e950f3, identical MoE code).

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
…p as E402-exempt

The gate scripts insert the repo root into sys.path before importing
rl_engine and tools, which flake8 reports as E402. Mark those imports
with noqa, as the other tools/ scripts do.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Annotate the latency samples and accept the default-argument lambdas used as
timed calls. No behaviour change.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
The tree is unchanged: RL-Align#469's refactor merge carries nothing this
branch lacks; this only makes the current RL-Align#469 tip an ancestor.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
Per the layout review on RL-Align#549:
- the TP4 MoE gate moves to tools/validation/models/qwen3_next/;
- the prior-art runner and its plot move to the Qwen3-Next benchmark PR
  (benchmarks/models/qwen3_next/);
- the committed evidence directory becomes a ZIP release attachment with
  a short summary in reports/experiments/qwen3_next/ (download link,
  SHA-256, tested commits, commands); the bytes are unchanged.
The GPU workflow now also triggers on files under the qwen3_next/
validation directories.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
@fusheng-ji fusheng-ji changed the title [WS1][CUDA][Qwen3-Next] MoE route/combine contract with a batch-invariant grouped expert GEMM (RFC #428 stack 5/8) [WS1][CUDA][Qwen3-Next] MoE route/combine contract with a batch-invariant grouped expert GEMM (RFC #428 stack 5/11) Oct 15, 2026
The raw results (JSON, figures, audit files) now live on the fork's pr_assets branch;
the documents link to them at a fixed commit.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
@fusheng-ji
fusheng-ji force-pushed the feat/ws1-qwen3-next-moe-route-combine branch from f9caee2 to 8943c14 Compare October 23, 2026 21:44
@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · rl_engine/_C.pyi (L281-L282): pybind11 converts std::vector<torch::Tensor> to a Python list, so list[torch.Tensor] is correct here.

Not changed: the stub carries types only, like every other entry in _C.pyi. activation is validated at runtime (TORCH_CHECK(activation == 0 || activation == 1) in the binding), and list[torch.Tensor] matches pybind11's conversion, as the finding notes.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · rl_engine/models/qwen3_next/qwen3_next_forward.py (L207-L224): n/a

💡 Suggested Change Before: ``` def stable_top10_routes(router_logits): """Select ten distinct experts and normalize their FP32 probabilities.

No change: this item has no description, and the suggestion replaces the body of stable_top10_routes with x, which would drop the shape, dtype and finiteness checks, the fixed-order softmax and the stable tie rule that the route contract requires.

Brings in RL-Align#468's a203719: the norm manifest's top-level keys, chain semantics and duplicate-op checks, the gate dtype message and the new norm tests.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
fusheng-ji added a commit to fusheng-ji/RL-Kernel that referenced this pull request Oct 26, 2026
Brings in RL-Align#468's a203719: the norm manifest's top-level keys, chain semantics and duplicate-op checks, the gate dtype message and the new norm tests.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · csrc/cuda/norm/rmsnorm.cu (L254-L256): Verified: block_reduce_sum (line 47) ends its loop with a __syncthreads() executed by all threads, so the store of rstd[row] followed by __syncthreads() is correctly ordered and cannot deadlock.

Agreed, no change: block_reduce_sum ends with a barrier that every thread executes, so the store of rstd[row] and the reads after it are ordered.

Comment on lines +85 to +90
- name: Checkout the reviewed commit
uses: actions/checkout@v4
with:
repository: ${{ github.event_name == 'workflow_dispatch' && inputs.source_repository || github.repository }}
ref: ${{ github.event_name == 'workflow_dispatch' && inputs.source_sha || github.event.pull_request.head.sha || github.sha }}
persist-credentials: false

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

security · high
actions/checkout is told to fetch github.event.pull_request.head.sha for pull_request events, so the PR-controlled head commit is what gets built and executed on the self-hosted runner. Combined with the PR trigger above, build/execute provenance is the unreviewed fork/branch tree, not a reviewed commit. Only use the dispatch source_sha (or the trusted base ref) here.

Suggestion:

Suggested change
- name: Checkout the reviewed commit
uses: actions/checkout@v4
with:
repository: ${{ github.event_name == 'workflow_dispatch' && inputs.source_repository || github.repository }}
ref: ${{ github.event_name == 'workflow_dispatch' && inputs.source_sha || github.event.pull_request.head.sha || github.sha }}
persist-credentials: false
ref: ${{ github.event_name == 'workflow_dispatch' && inputs.source_sha || github.sha }}

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Same answer as 549-I4381852969, 507-I4383081716 and r4372884187; not changed. Pushing a branch here already requires write access, the job runs with contents: read and no pull_request_target, and fork PRs never reach the runner; this is the trust model of ws1-chain-npu.yml. Checking out github.sha would test the merge ref instead of the PR head.

Comment on lines +123 to +137
- name: Check pinned providers in their own process
# The FP32 router needs IEEE Triton dots from process start.
env:
TRITON_F32_DEFAULT: ieee
run: |
set -o pipefail
log="$RUNNER_TEMP/provider-pytest.log"
python3 -m pytest -q -rs tests/models/qwen3_next/check_qwen3_next_norm_providers.py tests/models/qwen3_next/check_gdn_recurrent_golden.py tests/models/qwen3_next/check_qwen3_next_forward.py | tee "$log"
# Fail closed: these modules skip as a whole when CUDA or vLLM is missing, and pytest
# exits 0 for an all-skipped run. A skipped provider check is not provider evidence.
if grep -Eq '^SKIPPED|[0-9]+ skipped' "$log"; then
echo "::error::a provider check was skipped"
exit 1
fi
grep -Eq '[0-9]+ passed' "$log"

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

bug · high
set -o pipefail is on but set -e is not, and the step's final status is the exit code of the last command. Two concrete holes: (1) the grep -Eq '[0-9]+ passed' at the end also matches pytest's non-pass lines such as 12 passed, 3 failed / 10 passed, 2 skipped, so a run with skips or failures can still exit 0; (2) ^SKIPPED never matches pytest's summary (it prints SKIPPED [1] ... only with -rs, i.e. after the reason prefix), so a genuinely skipped provider check is only caught by the loosely-matching [0-9]+ skipped pattern. Anchor the checks on the pytest exit code and on the exact summary line.

Suggestion:

Suggested change
- name: Check pinned providers in their own process
# The FP32 router needs IEEE Triton dots from process start.
env:
TRITON_F32_DEFAULT: ieee
run: |
set -o pipefail
log="$RUNNER_TEMP/provider-pytest.log"
python3 -m pytest -q -rs tests/models/qwen3_next/check_qwen3_next_norm_providers.py tests/models/qwen3_next/check_gdn_recurrent_golden.py tests/models/qwen3_next/check_qwen3_next_forward.py | tee "$log"
# Fail closed: these modules skip as a whole when CUDA or vLLM is missing, and pytest
# exits 0 for an all-skipped run. A skipped provider check is not provider evidence.
if grep -Eq '^SKIPPED|[0-9]+ skipped' "$log"; then
echo "::error::a provider check was skipped"
exit 1
fi
grep -Eq '[0-9]+ passed' "$log"
run: |
set -eo pipefail
log="$RUNNER_TEMP/provider-pytest.log"
python3 -m pytest -q -rs tests/models/qwen3_next/check_qwen3_next_norm_providers.py tests/models/qwen3_next/check_gdn_recurrent_golden.py tests/models/qwen3_next/check_qwen3_next_forward.py | tee "$log"
# Fail closed: these modules skip as a whole when CUDA or vLLM is missing, and pytest
# exits 0 for an all-skipped run. A skipped provider check is not provider evidence.
if grep -Eq '^SKIPPED \[|[0-9]+ skipped' "$log"; then
echo "::error::a provider check was skipped"
exit 1
fi
grep -Eq '^[0-9]+ passed(,[^,]*)? in [0-9.]+s' "$log"
grep -qv 'failed\|error' "$log"

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Same point as 507-I4383081721, 557-I4381857006; not changed. The step runs under bash -e, and set -o pipefail is set in the same shell before the pytest | tee pipeline, so a failing pytest (exit 1, 2, 4 or 5) stops the step before either grep.

Comment on lines +78 to +79
dispatch(
rows.contiguous(),

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

bug · high
The top_k argument is inferred as top_k if rows.shape[0] == tokens else 1. When rows is the per-route [T*top_k, K] activation, rows.shape[0] == tokens is false so top_k=1, which is the intended per-route mode. But when rows is [T, K] shared across routes, the kernel is asked for top_k=10 with only T source rows — the fused MoE kernel indexes source rows per (token, slot). This dual-mode inference silently changes kernel semantics based on an accidental shape comparison (rows.shape[0] == tokens). If a caller passes [T, K] rows while tokens == T*10, or vice versa, the branches flip without any error. Make the mode an explicit parameter instead of inferring it from a shape comparison.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Same point as 549-I4383088158; not changed. rows.shape[0] is T for the token layout and 10 * T for the per-route layout; they are equal only for T == 0, which returns before this line.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · .github/workflows/qwen3-next-provider-gpu.yml: The pull_request trigger routes same-repository PR branches straight onto the long-lived self-hosted runner (`if: ...

Same answer as 549-I4381852969, 507-I4383081716 and r4372884187; not changed. Pushing a branch here already requires write access, the job runs with contents: read and no pull_request_target, and fork PRs never reach the runner; this is the trust model of ws1-chain-npu.yml.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · .github/workflows/qwen3-next-provider-gpu.yml (L123-L126): TRITON_F32_DEFAULT: ieee is only set on the pytest step, but check_qwen3_next_forward.py exercises shared_linear in FP32 (logits/weights created as float32), and rl_engine/models/qwen3_next/qwen3_next_forward.py fails when triton.knobs.language.fp32_default != "ieee".

Not changed. Each run: step is its own process, and the variable is set from process start on the pytest step, the only step that runs FP32 Triton dots; the build and smoke steps import only torch and _C. shared_linear raises unless the setting is ieee, so a misconfigured step fails instead of passing.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · csrc/bindings/ops.cpp (L421-L423): rmsnorm_check_weight only bounds x.numel() <= INT_MAX, but the kernels compute the row base as int row * H in 32-bit (x + row * H).

Agreed, no change: the binding bounds x.numel() by INT_MAX and requires a positive width, so T, H and every row * H fit in int.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · rl_engine/backends/cuda/norm/rmsnorm.py (L404-L409): parameter_vjp_contributions_fp32 ignores the weight_offset.

Not changed (see 468-S1). dweight does not depend on the offset, because d/dw (w + c) = 1; test_cuda_gated_backward_is_offset_aware asserts dweight is equal with and without it, and that dx and dgate see 1 + w.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · rl_engine/backends/cuda/norm/rmsnorm.py (L312-L314): backward assumes the saved x is 2-D when doing rstd.reshape(x.shape[:-1]) and dy.float() * normed * scale.

Agreed, no change: a float32 weight with a bf16 x is a supported, separately dispatched combination.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · rl_engine/_C.pyi (L275-L282): Stub uses list[torch.Tensor], while the rest of this file uses the list[...] form too — consistent, no change needed.

Agreed, no change: the stub matches the binding's return type, as the finding notes.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · rl_engine/models/qwen3_next/qwen3_next_forward.py (L133-L133): shared_linear wraps x as x.reshape(-1, x.shape[-1]) before calling _SharedLinear.apply, so x saved in ctx is always 2-D.

Agreed, no change: linear(grad.t(), x.t()) is grad.T @ x, of shape [N, K], which matches weight.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · rl_engine/models/qwen3_next/qwen3_next_forward.py (L274-L276): route_indices.flatten().tolist() forces a GPU→host sync inside backward (and the .flatten() on a possibly non-contiguous tensor).

Not changed. flatten() follows logical index order whatever the strides (it copies when it must), so slot t * top_k + k maps back to token slot // top_k. The host-side grouping per expert is the deliberate fixed-order, per-expert backward.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · rl_engine/models/qwen3_next/qwen3_next_forward.py (L161-L162): shared_route/shared_linear FP32 path: the router does shared_linear(x.float(), weight.float()) on every forward.

Not changed. Autograd's backward for .float() casts the gradient back to the BF16 parameter's dtype, so the weight does get its gradient. The fp32 copy is 4 MiB per call; caching it would go stale after each optimizer step.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · rl_engine/models/qwen3_next/qwen3_next_tp_blocks.py (L125-L130): self.gate.weight is registered as a plain nn.Linear parameter and is NOT marked with _parallel_parameter(..., duplicate=True) even though the router weight is replicated across TP ranks (each rank computes the full 512-way router).

Not changed. With no tensor_model_parallel attribute, MCore's param_is_not_tensor_parallel_duplicate counts the replicated router weight only on TP rank 0, so the grad norm does not double-count it; marking it duplicate=True would set shared and drop it from the norm. test_qwen3_next_tp_blocks.py pins that replicated weights carry no TP attributes.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · rl_engine/models/qwen3_next/qwen3_next_tp_blocks.py (L27-L28): moe_hf_shapes() rebuilds a 1538-entry nested dict on every call, and it is called once per shard_moe_weight/assemble_moe_weight per tensor name plus once per loop iteration in load_hf_weights/export_local_hf_weights (i.e. ~1538 calls of a 1538-entry construction during a single load = quadratic startup cost).

Not changed. Rebuilding the shape table takes about 0.1 ms per call (about 0.2 s per full load), which is negligible next to reading the checkpoint's 1538 expert tensors, and only loaders and tools call it.

@Flink-ddd Flink-ddd left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

LGTM, Thank you.

fusheng-ji added a commit to fusheng-ji/RL-Kernel that referenced this pull request Nov 14, 2026
…th the combine_routes patch

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
The autograd backward of combine_routes summed grad * expert_output over H
with torch.sum, whose reduction order depends on the batch shape, so the
route-weight gradient stopped being batch-invariant from 16 rows up. Add a
custom autograd Function that sums H with fixed_order_row_sum, and a test that
compares the first 8 rows alone against the same rows inside larger batches.

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
fusheng-ji added a commit to fusheng-ji/RL-Kernel that referenced this pull request Nov 14, 2026
… vs public head d5253a5

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
fusheng-ji added a commit to fusheng-ji/RL-Kernel that referenced this pull request Nov 14, 2026
d5253a5

Signed-off-by: Wenbo Ji <36562829+fusheng-ji@users.noreply.github.com>
@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · .github/workflows/qwen3-next-provider-gpu.yml (L28): the rl_engine/integrations/engines/train/vllm/qwen3_next* trigger path matches no tracked file, so the filter is dead.

Not changed. This workflow file is shared by the whole Qwen3-Next stack, and the next PR, #507, adds the files this entry covers: rl_engine/integrations/engines/train/vllm/qwen3_next_conv.py, qwen3_next_gdn.py and qwen3_next_provider.py. In this PR alone the entry matches nothing; once #507 lands it triggers the GPU job for those adapters. Dropping it here would only mean adding it back in #507.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · tests/models/qwen3_next/check_qwen3_next_forward.py (L70): the empty-token sub-check relies on weight.grad/bias.grad being allocated; a provider that short-circuits zero-row inputs would leave them None, and torch.count_nonzero(None) raises TypeError instead of asserting the zero-gradient contract.

Not changed. For an empty input the backward returns torch.zeros_like(weight) and a zero bias gradient for every operand that requires grad (_SharedLinear.backward, empty branch), so both gradients exist and the check asserts they are zero. If a provider ever left them None, torch.count_nonzero(None) would raise and the test would fail; it cannot pass silently, so the check does tell a missing gradient apart from a correct one.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · tests/models/qwen3_next/test_qwen3_next_forward_contract.py (L43): comparing fixed_order_row_sum with an FP64 sum at rtol=1e-5, atol=1e-4 looks flaky for a plain FP32 sequential reduction at width 512.

Not changed. fixed_order_row_sum is a pairwise tree of elementwise adds, not a sequential chain, so its rounding error grows with log2(width), and the inputs come from a seeded CPU generator, so every run sums the same values in the same order and gets the same bits. On this test's own inputs, the largest error is 1.5e-5 at width 512 against a tolerance of at least 3.4e-4. The worst error-to-tolerance ratio over all five widths is 0.018, and 20 repeated calls are bitwise identical. The reference is the double-precision sum rounded to FP32.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · tests/models/qwen3_next/check_qwen3_next_forward.py (L143): the inactive-expert gradient assertion does not validate sparse dispatch, since a dense implementation that multiplies by a zero routing weight would pass too.

Not changed. This assertion checks the gradient contract, not compute sparsity: an expert that no token selects must get an exactly zero dW, so nothing leaks into unused expert slices (a stray write, or a 0 * inf from a dense path). This PR does not claim sparse compute as a contract. The dispatch itself is checked bit for bit against per-expert GEMMs by test_grouped_experts_are_bitwise_the_per_expert_gemms and test_grouped_projection_is_bitwise_the_per_expert_gemm.

@fusheng-ji

Copy link
Copy Markdown
Contributor Author

OpenCodeReview · tests/models/qwen3_next/check_qwen3_next_forward.py (L99): the reference upstream gradient comes from the BF16 actual output (2 * actual.float()); building it from the double expected would make the reference VJP exact.

Not changed. The check compares vector-Jacobian products, so both sides must use the same upstream vector. combine_routes' backward receives g = 2 * actual (the gradient of actual.float().square().sum()), and the double reference is evaluated at that same g. The comparison therefore measures only the backward's error. With 2 * expected, the reference would use a different g from the one the provider received, and the difference would then include the forward rounding.

@KJLdefeated KJLdefeated left a comment •

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Overall looks good, just some questions before approve:

  1. Is TP invariance supported? Is Tp4 == Tp1? If this RFC does not require TP invariance, that's ok ignore my questions.
  2. Non-blocking: to make best performance, consider using MegaMoE to overlap communication and computation? And I believe MegaMoE can also reach BI by some configuration. (Maybe it is not this PR's scope, you can ignore me.)
  3. This query on RFC (Non-blocking): Is Qwen3-Next really needs TP4? Normally MoE only use EP, TP4 would cause more overhead. I just want to check if TP4 is really worth to do.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

new-model platform: cuda Specific optimizations or bugs in NVIDIA graphics cards (such as FlashInfer, TMA optimizations) qwen WS1

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants