Skip to content

Add SM90 MegaMoE by cooperative - #360

Open
mpdfdfl wants to merge 16 commits into
deepseek-ai:mainfrom
mpdfdfl:sm90-mega-moe-pingpong-coop
Open

mpdfdfl wants to merge 16 commits into
deepseek-ai:mainfrom
mpdfdfl:sm90-mega-moe-pingpong-coop

Conversation

@mpdfdfl

@mpdfdfl mpdfdfl commented Jun 13, 2026

Copy link
Copy Markdown

Add SM90 MegaMoE (cooperative)

Co-Authored-By: @OpenDarrenlu

First, thanks to #352 and #323 — this work builds on the ideas from those PRs.

We provide the SM90 fused MegaMoE through a single cooperative kernel
(BLOCK_M=64, BLOCK_N=256). The two math warpgroups cooperatively N-split one
tile — the first warpgroup owns the left 128 columns and the second owns the
right 128 columns (each an m64n128 WGMMA) — and share a single A-tile (weight)
load.

Why cooperative-only, and why the 64×256 tiling

We discovered an accuracy issue during end-to-end testing. We suspect the cause
is that the L1 output of the MoE layer is normally quantized per-128, whereas
our earlier implementation quantized it per-64. To fix this, we settled on a
single cooperative kernel in which the first warpgroup handles the left 128
columns and the other handles the right 128 columns, so the post-SwiGLU L1 output
is quantized per-128 (matching the standard DeepEP runner's
scale_block_size=128).

For the L2 GEMM, cooperative also beats a pingpong variant we tested: the
dequant FMAs inside the k-loop overlap the asynchronous WGMMAs, and the two
warpgroups share a single A-tile load.

The SM90 kernel / scheduler / heuristics are kept separate from the shared SM100
files, so SM100 behavior is unchanged.

Correctness test

tests/test_mega_moe_sm90.py checks the fused kernel against a PyTorch reference
(calc_diff < 0.01) over layered scenarios (smoke → shapes → edge cases →
random stress).

python tests/test_mega_moe_sm90.py --num-processes 8

Benchmark

tests/bench_mega_moe_sm90.py times the fused kernel end-to-end against the
DeepEP + DeepGEMM unfused baselines (V1 contiguous, V1 low-latency, V2
ElasticBuffer), same CUDA-event timing for all.

NVSHMEM_IBGDA_ENABLE=0 NVSHMEM_DISABLE_IBGDA=1 EP_DISABLE_GIN=1 \
python tests/bench_mega_moe_sm90.py --num-processes 8 \
  --hidden 4096 --intermediate-hidden 2048 --num-experts 256 --num-topk 8 \
  --baseline --baseline-version both \
  --batches 8 128 256 512 1024 2048 4096 8192

Here are our latest test results.

1. Single-operator benchmark (vs. deepep)

Shape — hidden=4096, moe_intermediate=2048, experts=256, topk=8 (MiMo-V2.5)

per-rank tokens ours (μs) V1-contig (μs) vs V1-contig V1-ll (μs) vs V1-ll V2 (μs) vs V2
8 262.3 1649.1 6.29x 272.6 1.04x 332.5 1.27x
128 282.5 1533.7 5.43x 327.1 1.16x 380.4 1.35x
256 348.4 1635.3 4.69x 389.3 1.12x 434.8 1.25x
512 497.4 1900.7 3.82x 557.0 1.12x 566.5 1.14x
1024 767.0 2336.9 3.05x 899.6 1.17x 888.0 1.16x
2048 1302.8 3341.1 2.56x 1597.4 1.23x 1543.5 1.18x
4096 2323.0 5320.2 2.29x - - 2878.6 1.24x
8192 4611.9 9408.6 2.04x - - 5516.0 1.20x

2. End-to-end prefill benchmark

Setup: input 16K, output 1, TTFT capped within 1s.

QPS Input token throughput (tok/s) TTFT (ms)
deepep 4.10 69635.40 1005.80
megamoe 5.15 86653.93 1037.03

📈 Compared with deepep, input token throughput improves by 25% (86653 vs 69635 tok/s).

3. Decode benchmark

Since this is a single-node deployment, TPOT is derived from the logs.

bs megamoe (median) deepep (median) Gap
36 18.69 ms 19.43 ms megamoe +4.0%
64 20.14 ms 20.94 ms megamoe +4.0%

TODO / future work

This is just a first version, and there is still a lot of room to optimize
inside the kernel. We would really love your help — let's optimize this kernel
together!
Planned directions:

  • Block-tile scheduling. The L2 cache hit rate is currently low (~40%);
    tuning the block-tile / expert-wave scheduling strategy should reduce HBM
    traffic and help the large-M DRAM-bandwidth-bound cases.
  • Green-context split-kernel. Explore the approach from PR Draft: Add green-context split-kernel MegaMoE features #357 — split the
    MoE into focused kernels (dispatch+L1+SwiGLU / L2+combine / combine-reduce)
    running concurrently on disjoint SM partitions via CUDA green contexts — to
    overlap stages the fused megakernel serializes internally.

mpdfdfl and others added 8 commits June 12, 2026 19:26
Two SM90 (Hopper) FP8 MegaMoE kernels with their own scheduler and
heuristics, kept isolated from the shared SM100 mega_moe scheduler/heuristics
so the SM100 path is untouched:

  * pingpong (BLOCK_M=64): one math warpgroup per tile, the two warpgroups
    overlap one's MMA with the other's epilogue via an OrderedSequenceBarrier.
    Tuned for small/medium M.
  * cooperative (BLOCK_M=128): the two math warpgroups cooperatively M-split a
    single tile and share one B-tile load, halving weight HBM traffic; a
    256-thread cross-warpgroup barrier closes the L2 epilogue. Tuned for large M.

New files: impls/sm90_fp8_mega_moe_{pingpong,cooperative}.cuh, the SM90 host
runtimes, scheduler/sm90_mega_moe.cuh (adds kClusterSize / kL2NMajorSchedule),
and heuristics/sm90_mega_moe.hpp.

Co-Authored-By: darrenlu <2945034270@qq.com>
Wire the two SM90 kernels into the host API and Python:

  * csrc/apis/mega.hpp: SM90-aware get_symm_buffer_size_for_mega_moe (per-arch
    SF dtype/granularity, guarded by is_sm90 so the SM100 path is unchanged);
    fp8_mega_moe routes by token count (DG_SM90_MOE_COOPERATIVE_THRESHOLD,
    default 256) between pingpong (<threshold) and cooperative (>=threshold);
    plus fp8_mega_moe_{pingpong,cooperative} forced entry points for A/B.
  * deep_gemm/mega/__init__.py + deep_gemm/__init__.py: expose fp8_mega_moe,
    fp8_mega_moe_{pingpong,cooperative} and transform_weights_for_mega_moe_sm90.
  * comm/barrier.cuh: guard the NVLink-timeout printf behind DG_NO_DEVICE_PRINTF
    (defined by the SM90 kernels) to avoid the ptxas C7510 WGMMA-pipeline
    serialization a function-call boundary would cause; upstream 300s timeout
    is preserved.

Co-Authored-By: darrenlu <2945034270@qq.com>
  * tests/test_mega_moe_sm90.py: layered correctness suite (L1 smoke .. L5
    stress) against a PyTorch reference via calc_diff. DG_SM90_MOE_KERNEL
    selects auto / pingpong / cooperative.
  * tests/bench_mega_moe_sm90.py: per-config TFLOPS / HBM / NVLink timing with
    an optional DeepEP (V1 contiguous + V2 ElasticBuffer) baseline comparison.
  * tests/_deepep_v1_baseline.py: the DeepEP V1 contiguous baseline used by the
    bench (Triton SwiGLU + FP8 quant + grouped GEMM pipeline).

Co-Authored-By: darrenlu <2945034270@qq.com>
Co-Authored-By: darrenlu <2945034270@qq.com>
Co-Authored-By: darrenlu <2945034270@qq.com>
Co-Authored-By: darrenlu <2945034270@qq.com>
Co-Authored-By: darrenlu <2945034270@qq.com>
@mpdfdfl
mpdfdfl force-pushed the sm90-mega-moe-pingpong-coop branch from 1380676 to ec757bd Compare June 17, 2026 08:54
@Rachmanino

Copy link
Copy Markdown

hi, which machine are you working on, h20, h100 or h200?

@mpdfdfl

mpdfdfl commented Jun 17, 2026

Copy link
Copy Markdown
Author

hi, which machine are you working on, h20, h100 or h200?

we working on H200

mpdfdfl and others added 2 commits June 17, 2026 22:31
Co-Authored-By: darrenlu <2945034270@qq.com>
Co-Authored-By: darrenlu <2945034270@qq.com>
@mpdfdfl mpdfdfl changed the title Add SM90 MegaMoE by pingpong and cooperative Add SM90 MegaMoE by cooperative Jul 9, 2026
@mpdfdfl
mpdfdfl force-pushed the sm90-mega-moe-pingpong-coop branch 2 times, most recently from e2050ed to f983307 Compare July 16, 2026 14:25
mpdfdfl and others added 6 commits July 31, 2026 13:56
…ean comments

- sm90_fp8_mega_moe_cooperative.{cuh,hpp} -> sm90_fp8_mega_moe.{cuh,hpp};
  kernel/runtime/launcher/heuristic symbols renamed accordingly
- remove duplicated fp8_mega_moe_cooperative C++/Python entry and the
  DG_SM90_MOE_KERNEL routing in tests/bench
- compress version-history comment to one-line JIT cache key (v30); comments
  trimmed to SM100 style; fix stale per-64 / M-split wording
- drop unused kSwizzleCDMode; cd_stages default 1 -> 2; order_barriers kept
  (reserved for L2 pingpong)

Co-Authored-By: darrenlu <2945034270@qq.com>
…ns with SM100)

SMEM send staging is now one pull chunk per dispatch warp instead of the full
token (hidden=7168: 14336B -> 7168B), the pull loop streams load/store chunks
through the single buffer with the SF copy overlapping the last chunk's NVLink
load. Pure data movement, results bit-identical. JIT cache key bumped to v31.

Co-Authored-By: darrenlu <2945034270@qq.com>
…g math WGs)

One WG's NVLink scatter now overlaps the other WG's WGMMA, taking the L2 push
tail off the critical path. L1 stays cooperative (per-128 SF needs the cross-WG
amax exchange). Key mechanics:
- scheduler gains an independent L2 tile width (BLOCK_N_L2=128 under pingpong)
- non-owner fast-forwards stage_idx/phase without waiting full/arriving empty;
  owner warps arrive the empty barrier with count 2 (new ptx::mbarrier_arrive
  overload) to keep the init count of 8
- L2 CD becomes per-WG exclusive (halved to 64x128 bf16; 32KB SMEM saved), the
  shared-tile trailing 256-thread barrier is replaced by a per-WG one, and an
  L2->next-wave-L1 rendezvous guards the CD byte aliasing across WGs
- B loads stay 256-wide at 256-aligned n (owner reads its half): loader
  untouched, at the cost of 2x L2 A/B HBM reads
- A/B switch: DG_SM90_MOE_L2_PINGPONG=0 reverts (default on); cache key v32

Co-Authored-By: darrenlu <2945034270@qq.com>
…l/empty

A local-only fast-forward lets the non-owner WG run 2+ ring wraps ahead of the
physical mbarrier state; parity waits alias every 2 completions, so its next
full-wait returned early, it arrived empty prematurely, and the pipeline hung
from the second L2 tile on an SM (repro: L2.heur.midB.t512; midA.t256 escaped
because each SM only got one L2 tile). The non-owner now waits full and arrives
empty for skipped tiles without computing — the parity invariant (a waiter took
part in the stage's previous use) is restored, the empty barrier keeps its
original 8x count-1 releases, and the count-2 mbarrier_arrive helper is
reverted. Cache key v33.

Verified: test_mega_moe_sm90.py passes all 32 scenarios with pingpong on and off.

Co-Authored-By: darrenlu <2945034270@qq.com>
…e SMEM SFB slots

Math WGs no longer __ldg weight SF from HBM (both L1 and L2); the B loader
stores the 2 covering SF floats per stage next to its B-tile TMA, published
through the existing full/empty mbarrier protocol (arrive release / wait
acquire), so no extra synchronization is needed. Bit-exact vs baseline.

Co-Authored-By: darrenlu <2945034270@qq.com>
Pingpong doubled L2 A/B supply traffic and stayed 15-21% behind cooperative
even with SMEM-staged weight SF, so the experiment is dropped: kL2Pingpong
template param, non-owner participation loop, L2->L1 rendezvous, per-WG L2
CD buffers, order_barriers, and the DG_SM90_MOE_L2_PINGPONG env are all gone.

Co-Authored-By: darrenlu <2945034270@qq.com>
@mpdfdfl
mpdfdfl force-pushed the sm90-mega-moe-pingpong-coop branch from f983307 to e481caf Compare August 3, 2026 06:31
anyunshan added a commit to anyunshan/DeepGEMM that referenced this pull request Aug 13, 2026
Kernel, scheduler, JIT host impl/heuristics and tests taken as-is from
PR deepseek-ai#360 (base 88965b0). No content changes; shared-file reconnection
follows in subsequent commits.
anyunshan added a commit to anyunshan/DeepGEMM that referenced this pull request Aug 13, 2026
From PR deepseek-ai#360: DG_NO_DEVICE_PRINTF suppresses the timeout printf in SM90
TUs (avoids ptxas C7510 WGMMA serialization across the call boundary);
mul2 wraps __fmul2_rn which is SM100+, scalar fallback otherwise.
Kept main wording of the timeout message.
anyunshan added a commit to anyunshan/DeepGEMM that referenced this pull request Aug 13, 2026
The SM90 kernel (PR deepseek-ai#360) tracks L1->L2 readiness per full-pool block via
l1_arrival_count/l2_arrival_mask; upstream 559d79f replaced these with the
ring-buffer full/empty counters. Re-add them as SM90-only fields appended
AFTER all existing fields so every SM100 offset is unchanged, plus:
- num_max_pool_blocks (constructor-derived, mirrors the old base)
- num_ring_tokens defaults to 0 (SM90 constructs Workspace without a ring)
- restore layout::get_num_padded_sf_pool_tokens (SM90 SF buffers span the
  full pool, not the ring)
anyunshan added a commit to anyunshan/DeepGEMM that referenced this pull request Aug 13, 2026
559d79f added a timeout printf to grid_sync (absent on PR deepseek-ai#360 base);
wrap it in DG_NO_DEVICE_PRINTF like the nvlink_barrier one.
anyunshan added a commit to anyunshan/DeepGEMM that referenced this pull request Aug 13, 2026
- csrc/apis/mega.hpp: dedicated get_symm_buffer_size_for_sm90_fp8_mega_moe
  (per-128-K float SF views, full-pool buffers, empty shared slots) and the
  fp8_mega_moe entry, both registered; SM100 sizing path untouched.
- deep_gemm/mega/__init__.py: mma_type=fp8xfp8 routes SymmBuffer to the SM90
  sizing; arch-dependent default (fp8xfp8 on SM90); SM90 weight transform +
  fp8_mega_moe wrappers (PR deepseek-ai#360 semantics, main-style mma_type API).
- exports + .gitignore profiling artifacts.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants