Repository navigation
Add sgl-project's NVFP4 MegaMoE as an additive deep_gemm.mega_nvfp4 module - #18
vipulSharma18 wants to merge 1 commit into
Conversation
… module Port the SM100 NVFP4 x NVFP4 MegaMoE decode kernel and the fused `mega_moe_pre_dispatch` kernel from sgl-project/DeepGEMM@ec226bc (deepseek-ai#27, 593e376, deepseek-ai#80, deepseek-ai#81) so Fireworks can drop the separate `sgl_deep_gemm` wheel. Purely additive: the only change to an existing file is the registration hook in csrc/python_api.cpp. No existing header changes, so no existing kernel's JIT cache key changes. - Device code: sgl's headers copied into deep_gemm/include/deep_gemm/mega_nvfp4/ (kernel, pre-dispatch, layout, scheduler, barrier, math/utils/tma_copy, ptx ld_st/tcgen05/tma/utils, sym_buffer, and the `MmaKind` block of types.cuh), moved into namespace `deep_gemm::mega_nvfp4` with includes redirected and nothing else changed (line numbers kept, so even assert-path SASS matches). - Host code: csrc/apis/mega_nvfp4.hpp plus heuristics/nvfp4_mega_moe.hpp and impls/sm100_nvfp4_mega_moe{,_pre_dispatch}.hpp, in `deep_gemm::mega_nvfp4` with `nvfp4`-marked names. Like the separate sgl library, these kernels keep their own SM count, compile for sgl's `sm_100f` target and launch without PDL, so `deep_gemm.set_num_sms`/`set_pdl` and these kernels never affect each other. Not ported: TRT-LLM weight layout, FP8 combine, BF16 and SM90 paths. - Python: deep_gemm/mega_nvfp4 exposes the sgl_deep_gemm subset Fireworks uses (get_symm_buffer_for_mega_moe, transform_weights_for_mega_moe, mega_moe_pre_dispatch, nvfp4_mega_moe, set_num_sms, get_block_m_for_mega_moe, transform_ue4m3_sf_into_required_layout). - tests/test_mega_moe_nvfp4.py ports sgl's pre-dispatch and NVFP4 alphas tests and adds a bytewise check of the NVFP4 pre-dispatch quantization. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
246f92c to
ef210f2
Compare
There was a problem hiding this comment.
Cursor Bugbot has reviewed your changes and found 1 potential issue.
❌ Bugbot Autofix is OFF. To automatically fix reported issues with cloud agents, have a team admin enable autofix in the Cursor dashboard.
Reviewed by Cursor Bugbot for commit ef210f2. Configure here.
|
|
||
| CUTLASS_DEVICE void release_task_info() const { | ||
| task_info_empty_barriers[sched_stage_idx ^ 1].arrive(0u); | ||
| } |
There was a problem hiding this comment.
Scheduler slot reuse missing fence
High Severity
MegaMoEScheduler::release_task_info arrives the empty barrier without a preceding ptx::fence_acq_rel_cta. In-flight task_infos metadata loads can still be pending when the scheduler reuses that pipeline slot, which races and can corrupt the next task. A later grid_sync or nvlink_barrier does not order these slot reads against the arrive.
Triggered by learned rule: MegaMoEScheduler must fence before slot reuse
Reviewed by Cursor Bugbot for commit ef210f2. Configure here.
There was a problem hiding this comment.
Risk: high. Not approved: this adds a large new SM100 NVFP4 MegaMoE kernel and public deep_gemm.mega_nvfp4 API (thousands of lines of device/host code), which is above the low-risk auto-approve threshold. Human review is needed; no reviewers were assigned.
Sent by Cursor Approval Agent: Pull Request Approver




What
This adds sgl-project/DeepGEMM's SM100 NVFP4 × NVFP4 MegaMoE decode kernel and its fused
mega_moe_pre_dispatchkernel, taken fromec226bc(sgl deepseek-ai#27,593e376, deepseek-ai#80, deepseek-ai#81). They are exposed as a new module,deep_gemm.mega_nvfp4, which provides the subset of thesgl_deep_gemmAPI that Fireworks uses:transform_ue4m3_sf_into_required_layouttransform_weights_for_mega_moeget_symm_buffer_for_mega_moemega_moe_pre_dispatchnvfp4_mega_moeget_block_m_for_mega_moeset_num_smsWhy
MiniMax-M3.1 decode (
FIREWORKS_DPEP_MEGAMOE=1) currently ships a second DeepGEMM wheel,sgl_deep_gemm, next to ours. With this module in our fork, Fireworks can drop that dependency and its separate build. That also unblocks landing the M3.1 decode stack (fireworks#55221) onmain.Design: purely additive
csrc/python_api.cppgains an include and aregister_apiscall (+2 lines). Every other file is new: 22 files, +7128/−0.deep_gemm/include/deep_gemm/mega_nvfp4/(kernel, pre-dispatch, layout, scheduler, barrier, math/utils/tma_copy, the ptx helpers, sym_buffer, and theMmaKindblock oftypes.cuh). The only edits are the namespacedeep_gemm::mega_nvfp4and the include paths. Line numbers are kept, so even the assert-path SASS matches sgl._Centry points have separate,nvfp4-marked names in a newcsrc/apis/mega_nvfp4.hpp,heuristics/nvfp4_mega_moe.hppandimpls/sm100_nvfp4_mega_moe{,_pre_dispatch}.hpp.sm_100ftarget and launch without PDL. Sodeep_gemm.set_num_smsanddeep_gemm.set_pdlnever affect them, and they never affect the existing kernels.deepgemm-situ,-mhc-graph-replay,-sigmoid-bf16or-mhc-capabilitychange. All four apply unchanged on top of this branch.Base branch
This PR is stacked on #19, which folds Fireworks' four build patches into
main, and is rebased onto it (ef210f2). Merge #19 first; this PR then retargets tomain. After both merge, Fireworks buildsmainhead with no patches, so any merge style is fine.Verification (B300)
The first round compared
c2e49f5+ 4 patches against this PR + 4 patches. After the rebase onto #19, the agent re-ran three checks: zero new JIT cache keys for existing kernels vs #19, bitwise equality withsgl_deep_gemmat EP2 (92 tensors) and EP8 (356), andtests/test_mega_moe_nvfp4.pypassing..pyfiles are byte-identical.inspect.signatureacrossdeep_gemm,deep_gemm.mega,_C)_C.*nvfp4*names. Nothing removed or changed.test_mega_moe.pyfp8×fp4 and BF16 at 8 ranks,test_mega_moe_compat.py,test_mega_gate.py,test_mega_mhc.py,test_bf16.py, and the dense and m-grouped-contiguous tests fromtest_fp8_fp4.py.sgl_deep_gemm@ec226bcset_num_sms(120). That is 860 tensors across 6 runs. The SASS is identical.tests/test_mega_moe_nvfp4.pypasses. It ports sgl's tests: pre-dispatch bytewise for fp8, mxf4 and nvfp4, and the 2-rank alphas test.Not run:
test_fp8_fp4.pyon both sides. An outside process killed the baseline run after 33 minutes, so only its dense and m-grouped-contiguous tests ran.test_mega_moe_training.py.Notes for reviewers
sgl_deep_gemm, the kernels compile on first use, so callers need an eager forward before CUDA-graph capture. Fireworks already does this.Fireworks follow-up (separate PR)
megamoe_nvfp4.py,import deep_gemm.mega_nvfp4 as deep_gemm.fw-ai/DeepGEMMmainhead (Fireworks PR prepared as a draft) with no build patches.sgl_deep_gemmbuild.🤖 Generated with Claude Code
Note
Medium Risk
New distributed, symmetric-memory MoE inference path with JIT SM100 kernels and strict buffer layout contracts; isolated from existing APIs but operationally critical for decode once adopted.
Overview
Introduces
deep_gemm.mega_nvfp4, a Fireworks port of sgl-project DeepGEMM’s SM100 NVFP4×NVFP4 MegaMoE decode stack so MiniMax-style decode can drop the separatesgl_deep_gemmwheel. The change is additive: onlycsrc/python_api.cppis modified (registermega_nvfp4APIs); everything else is new host glue, JIT launchers, copied device headers underdeep_gemm/include/deep_gemm/mega_nvfp4/, and a Python package (SymmBuffer,get_symm_buffer_for_mega_moe,mega_moe_pre_dispatch,nvfp4_mega_moe, weight/SF transforms).Host entry points live in
csrc/apis/mega_nvfp4.hppwithnvfp4-prefixed_Cnames and namespacedeep_gemm::mega_nvfp4, avoiding shadowing existing MegaMoE APIs. Symmetric-buffer sizing slices views for quantized activations, scale factors, top-k metadata, and ring buffers;nvfp4_mega_moevalidates weights/layout and dispatchessm100_nvfp4_mega_moeon arch 10 only. Heuristics (nvfp4_mega_moe.hpp) keep a separate SM count (set_num_sms_for_nvfp4_mega_moe), compile these kernels withsm_100ffor sgl parity, and launch with PDL disabled so globalset_num_sms/set_pdldo not affect them.Not ported on the host side: TRT-LLM weight layout, FP8 combine (
use_fp8_combineis always false), BF16 MegaMoE, and SM90 kernels. Python rejects legacyDG_USE_FP4_ACTS/DG_USE_MXF4_KINDin favor of explicitmma_type.Reviewed by Cursor Bugbot for commit ef210f2. Bugbot is set up for automated code reviews on this repo. Configure here.