Skip to content

big M opt: splitT1 + L2 scatter + 2 stream - #5

Merged
AichenF merged 1 commit into
AichenF:megamoe_nvfp4_devfrom
ouershua2299:bigMopt
Sep 3, 2026
Merged

AichenF merged 1 commit into
AichenF:megamoe_nvfp4_devfrom
ouershua2299:bigMopt

Conversation

@ouershua2299

Copy link
Copy Markdown
Collaborator

Summary

Base megamoe_nvfp4_dev @ c277f44. Three large-M split optimizations, plus a unified entry that keeps one weight copy and picks the kernel family per forward. Kernel only, no SGLang integration. 16 files, +1628 / -183.

Dynamic family selection

rho = M * topk / local_experts

rho <= 192  ->  BN256 fused
rho  > 192  ->  BN128 split L1 + L2

Decided in the C++ entry from this call's y.size(0), so decode, small prefill and large prefill can take different families in one process. Model name and hidden size no longer participate. Both families read the same packed-B bytes; block_n only selects the retained scale-metadata view. kernel_family="auto"|"fused"|"split" — production uses "auto".

What changed — large M (split)

Always enabled together, emitted as compile-time defines into the JIT key.

Name Change Nature
splitT1 Emit DG_NVLINK_BARRIER_TRAP_ONLY_TIMEOUT for the split JIT source, replacing the 7-arg device printf in the NVLink barrier timeout branch with a bare trap. dev already did this for fused; split did not. barrier.cuh unchanged No compute logic change
2 stream L1 dispatch-dequant half-stream mode4 + physical-warp remap. New code; dev has no half-stream machinery. split_l1_body.inl 983 -> 1679 lines Dequant publication timing, mbarrier protocol
L2 scatter L2 stages one warp-private 8-row half tile in shared memory, so each destination row leaves as a single 256-byte burst instead of eight scattered 16-byte requests. New code. split_l2_body.inl 659 -> 751 lines Shared-memory layout, cross-rank write path

What changed — small M (fused)

Fused device body is byte-identical to dev. The dequant functions it calls (dequant_braided_quad*, dequant_braided_selector_word, dequant_smem_b_from_packed_braided_lut_window) are unchanged — the +295 lines in mode2_dequant.cuh are split-only functions fused never calls.

Two host lines differ:

Location Change Effect
heuristics/..._small_m.hpp:130 use_mode2_lop3_decoder false -> true, only in the rho <= 6 bucket Needed so one Mode2 braided weight copy serves both families. Both decoders lossless
impls/..._small_m.hpp:287 fused launch passes launch_pdl=false (dev defaults to true) Launch behaviour only

choose_nvfp4_block_n_for_mega_moe_sm90 drops the intermediate_hidden >= 3072 special case (cutoff 190) and always uses 192. Pro's crossover moves M=1520 -> 1536; Flash and MiMo unaffected.

Performance

Microseconds, 8-GPU rank-MAX median. dev = c277f44, W8A8 = deepseek-ai/DeepGEMM PR deepseek-ai#383 @ bc4f33a.

Large M — split (the bigMopt work)

Model M dev bigMopt W8A8 bigMopt vs dev bigMopt vs W8A8
Flash 2048 1234.7 1103.2 957.6 +10.65% -15.20%
Flash 4096 2173.6 1882.8 1698.2 +13.38% -10.87%
Flash 8192 4170.4 3602.9 3239.6 +13.61% -11.21%
Pro 2048 2990.6 2765.3 2392.3 +7.54% -15.59%
Pro 4096 5304.1 4859.4 4140.5 +8.38% -17.36%
Pro 8192 9844.4 9021.8 7841.7 +8.36% -15.05%

bigMopt vs dev positive = bigMopt faster. Geometric mean vs dev: Flash -12.56%, Pro -8.09%, all 6 points -10.35% (1.1155x). Faster than dev at 6/6, still slower than W8A8 at 6/6 (geometric mean +14.19%).

Small M — fused

bigMopt runs dev's fused body byte-for-byte, so this is a precision comparison, not a bigMopt speedup claim. Negative = NVFP4 faster.

Model M NVFP4 fused (dev = bigMopt) W8A8 W8A8 vs NVFP4
Flash 8 205.2 235.9 -15.0%
Flash 16 215.5 262.2 -21.7%
Flash 32 237.8 256.7 -7.9%
Flash 64 245.2 267.9 -9.3%
Flash 128 244.2 276.7 -13.3%
Flash 256 289.7 290.2 -0.2%
Flash 512 453.1 414.7 +8.5%
Flash 1024 747.4 580.2 +22.4%
Pro 8 576.8 699.0 -21.2%
Pro 16 710.9 798.0 -12.3%
Pro 32 724.5 842.3 -16.3%
Pro 64 753.2 856.7 -13.7%
Pro 128 839.5 866.6 -3.2%
Pro 256 822.6 877.8 -6.7%
Pro 512 1234.5 1000.5 +19.0%
Pro 1024 2006.0 1367.8 +31.8%

Geometric mean: M<=256 (12 points) NVFP4 faster by 11.55%; M>=512 (4 points) W8A8 faster by 20.84%. Crossover between M=256 and M=512. PR deepseek-ai#383's FP8 path is always split L1+L2 while NVFP4 here is fused, so these rows compare two kernel structures.

Protocol

Small-M table: bench_kineto, 8 GB L2 flush, num_tests=20 internal mean is one observation, median of 50 observations for M<=128 and 20 for M>=256, same rule for every arm. Large-M table: the protocol behind the bigMopt column was not supplied and raw logs are not archived here; its dev and W8A8 columns agree with the small-M campaign to within 1.2%.

Full detail in README_BIGMOPT.md.

…L2 scatter + 2 stream)

Kernel-only. Base: megamoe_nvfp4_dev @ c277f44.

One weight format, one entry point. Every forward picks the kernel family from
the routed load actually seen by this rank:

    rho = M * topk / local_experts
    rho <= 192  -> BN256 family, small-M fused path
    rho  > 192  -> BN128 family, split L1 + split L2

Selection happens in the C++ entry from this call's y.size(0), so decode, small
prefill and large prefill can take different families inside one serving
process; model name and hidden size no longer participate. Both families read
the same packed-B bytes — block_n selects only the retained scale-metadata view.
kernel_family="auto"|"fused"|"split" is exposed for cross-validation only;
production must use "auto".

Large-M split always enables all three optimizations together, emitted as
compile-time defines into the JIT key so no environment variable can silently
select an unvalidated variant:

  1. splitT1 — compile specialization of the NVLink barrier timeout path:
     DG_NVLINK_BARRIER_TRAP_ONLY_TIMEOUT is now emitted for the split JIT
     source, replacing the 7-argument device printf in the barrier spin-wait
     timeout branch with a bare trap. dev already did this for fused; split
     did not;
  2. 2 stream — L1 dispatch-dequant half-stream mode4 + physical-warp remap
     (DG_NVFP4_L1_DEQUANT_HALF_STREAM=4, DG_NVFP4_L1_DEQUANT_WARP_REMAP=1,
     with a static assert binding remap to mode4);
  3. L2 scatter — split L2 (down) warp-private contiguous 8-row shared-memory
     staging, so each destination row leaves as one 256-byte burst instead of
     eight scattered 16-byte requests billed as 32-byte sectors.

Large-M result versus dev (Flash and Pro, M=2048/4096/8192): faster at all 6
points, +7.54% to +13.61%, equal-weight geometric mean -10.35% latency. Still
slower than the W8A8 calibration reference at all 6 points, geometric mean
+14.19%. The measurement protocol behind those numbers was not supplied and is
not reproducible from this repository; README_BIGMOPT.md records that gap
explicitly. No fused-path measurements are included.

Small M keeps dev's fused device body byte-identical
(sm90_nvfp4_mega_moe_small_m_fused_body.inl is unchanged), but the fused path
is not bit-for-bit dev behaviour — three things differ:

  - heuristics/sm90_nvfp4_mega_moe_small_m.hpp: the rho<=6 bucket flips
    use_mode2_lop3_decoder false -> true, because one common Mode2 braided
    packed-B copy now serves both families;
  - impls/sm90_nvfp4_mega_moe_small_m.hpp: the fused launch passes
    launch_pdl=false;
  - impls/sm90_nvfp4_mega_moe_mode2_dequant.cuh is included by both the split
    and small-M .cuh, so its changes are on the small-M compile path too.

The other five small-M tuning buckets are unchanged. Reverting the first and
third would restore stock dev small-M behaviour but would break the
single-weight-copy premise.

On the weight side, choose_nvfp4_block_n_for_mega_moe_sm90 drops the
intermediate_hidden>=3072 special case (cutoff 190) and always uses
family_threshold=192.

The SGLang integration overlay, built wheels and raw gate evidence are not part
of this branch; README_BIGMOPT.md records where the full sealed delivery is
retained.
@AichenF
AichenF merged commit dbd995f into AichenF:megamoe_nvfp4_dev Sep 3, 2026
@ouershua2299
ouershua2299 deleted the bigMopt branch September 4, 2026 02:25
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