[TRTLLM-14822][feat] deprecate WIDEEP MoE backend - #17119
Conversation
|
PR_Github #63543 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #63551 [ run ] triggered by Bot. Commit: |
|
PR_Github #63551 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #63556 [ run ] triggered by Bot. Commit: |
|
PR_Github #63556 [ run ] completed with state
|
265fc80 to
5dab9bf
Compare
|
/bot run --disable-fail-fast |
|
PR_Github #63587 [ run ] triggered by Bot. Commit: |
|
PR_Github #63587 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #63652 [ run ] triggered by Bot. Commit: |
|
PR_Github #63652 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #63705 [ run ] triggered by Bot. Commit: |
|
PR_Github #63705 [ run ] completed with state |
Raise an explicit error for the WIDEEP MoE backend in both get_moe_cls and create_moe_backend, so the backend cannot be selected through create_moe or by calling create_moe_backend directly. Wide expert parallelism and EPLB remain available through the other backends. TestDeepSeekV4FlashBase FP8 EPLB cases previously required WIDEEP because CutlassFp8BlockScaleGemmRunner takes a Hopper-only wgmma path with no SM100/SM103 implementation. DEEPGEMM routes the same block-scale weights through DeepGEMM, so those cases now use DEEPGEMM. Examples, docs and test lists that referenced WIDEEP are updated to CUTEDSL or DEEPGEMM accordingly. The layer_wise_benchmarks --scaled-from test is skipped because that feature is implemented by monkeypatching WideEPMoE.select_alltoall_method_type. Signed-off-by: xxi <xxi@nvidia.com>
…tion The perf disagg entries named gb200_wideep_* / gb300_wideep_* were removed by mistake. The wideep token is only a filename prefix describing the wide-EP deployment shape (dep32 + eplb288); those configs already set moe_config.backend: CUTEDSL and are unaffected by the WIDEEP deprecation. All ten entries are restored in llm_perf_disagg.yml and llm_perf_multinode.txt. For the DeepSeekV3Lite online EPLB cases the WIDEEP entries were the only coverage of mtp_nextn=0 and fp8kv=False, so deleting them dropped those dimensions entirely. They are now replaced with the equivalent backends still present in each test parametrize list rather than deleted. rtx6k keeps no nvfp4 online EPLB entry: that test parametrize is now TRTLLM-only and TRTLLMGenFusedMoE rejects SM120, so no equivalent backend exists on that platform. The commented-out DeepSeek-V4 EPLB sanity entries now point at the tests that actually exist (4gpus variants) instead of being dropped. Signed-off-by: xxi <xxi@nvidia.com>
WIDEEP dispatched FP8 block-scale MoE on is_sm_100f() internally: Cutlass for SM90/SM120, DeepGEMM for SM100/SM103. Hardcoding DEEPGEMM after the deprecation dropped both Cutlass platforms, because DeepGemmFusedMoE only implements SM100/SM103 and the DEEPGEMM dispatch branch has no SM fallback. Select the backend at runtime with the same is_sm_100f() condition so the test keeps its original per-platform coverage. Signed-off-by: xxi <xxi@nvidia.com>
…uide ep_load_balancer/README.md kept three runnable YAML snippets on the WIDEEP backend. They were doubly invalid: the top-level moe_backend key is not an LlmArgs field (BaseLlmArgs sets extra=forbid) and WideEP does not match the strict Literal on MoeConfig.backend. Drop the stray key and move the snippets to CUTEDSL, which supports FP8 block scales on SM>=90 and is in the EPLB backend allowlist. layer_wise_benchmarks/README.md advertised WIDEEP in six commands and two support notes. TRTLLM_FORCE_ALLTOALL_METHOD is read only by WideEPMoE, so the DeepEP examples now use TRTLLM_FORCE_COMM_METHOD, which CommunicationFactory reads on the ConfigurableMoE path. The supported-backend list is corrected to match runner.py. MOE_DEVELOPER_GUIDE.md described WideEPMoE as Deprecating; it is now rejected outright. Signed-off-by: xxi <xxi@nvidia.com>
…eference Fold in the fix from the closed PR NVIDIA#17210. The checkpoint was missing from references/gsm8k.yaml, so the spec lookup raised Not registered specs and test_fp8_4gpus_static_eplb ended up waived on all four Blackwell platforms. The three measurements recorded there came from the WIDEEP variant this branch deprecates. A DEEPGEMM run on 4x B300 scored 91.13, inside that spread, which is expected because both backends share the DeepGEMM FP8 block-scale path on SM100/SM103. The four waivers go away now that the reference resolves. Signed-off-by: xxi <xxi@nvidia.com>
5dab9bf to
e7e9da5
Compare
|
/bot run --disable-fail-fast |
|
PR_Github #63756 [ run ] triggered by Bot. Commit: |
Summary
Deprecate the
WIDEEPMoE backend by raising an explicit error in bothget_moe_clsandcreate_moe_backend, so it cannot be selected viacreate_moenor by callingcreate_moe_backenddirectly. Wide expert parallelism and EPLB stay available through the remaining backends (DEEPGEMMfor FP8 block-scale checkpoints,TRTLLM/CUTEDSL/CUTLASSotherwise).AUTOnever resolves toWIDEEP, so no implicit user path is affected.Why the FP8 EPLB tests can move off WIDEEP
TestDeepSeekV4FlashBaseFP8 EPLB cases used to requireWIDEEPbecause the default CUTLASS FP8 block-scale path (CutlassFp8BlockScaleGemmRunner::moeGemm) takes a Hopper-onlywgmmaroute with no SM100/SM103 implementation.WIDEEPavoided this via anis_sm_100f()check that switches to DeepGEMM.DEEPGEMMreaches the same kernel directly, so those cases now useDEEPGEMM.Other updates
WIDEEPnow point atCUTEDSL/DEEPGEMM.WIDEEPare remapped to an equivalent case rather than deleted; see Test list handling below.WideEPMoEis dropped from the mirroredswiglu_limitcapability tuple inmodeling_deepseekv4.pyto keep it consistent withcreate_moe.py.--scaled-fromtest is skipped: that feature is implemented by monkeypatchingWideEPMoE.select_alltoall_method_type.Test list handling
Entries are remapped to an equivalent case rather than deleted, so no coverage dimension is lost. In
llm_function_core.txttheWIDEEPentries were the only coverage ofmtp_nextn=0andfp8kv=False, so plain deletion would have dropped those dimensions entirely.WIDEEPentryllm_function_core.txttest_bfloat16_4gpus_online_eplb[mtp_nextn=0-...][mtp_nextn=0-moe_backend=CUTLASS]llm_function_core.txttest_bfloat16_4gpus_online_eplb[mtp_nextn=2-...][mtp_nextn=2-moe_backend=TRTLLM]llm_function_core.txttest_nvfp4_4gpus_online_eplb[fp8kv=False-...][fp8kv=False-moe_backend=TRTLLM]llm_function_core.txttest_nvfp4_4gpus_online_eplb[fp8kv=True-...]TRTLLMvariant was already listedllm_function_rtx6k.txttest_bfloat16_4gpus_online_eplb[mtp_nextn=0/2-...][...-moe_backend=CUTLASS]llm_function_rtx6k.txttest_nvfp4_4gpus_online_eplb[fp8kv=False/True-...]l0_dgx_b200.ymltest_*_8gpus_static_eplb[moe_backend=WIDEEP]test_nvfp4_4gpus_static_eplb[moe_backend=TRTLLM]andtest_fp8_4gpus_static_eplb[moe_backend=DEEPGEMM]; the8gpusvariants do not exist in the sourcewaives.txttest_bfloat16_4gpus_online_eplb[mtp_nextn=2-...] SKIP (nvbugs/6313993)WIDEEPcombination onlyrtx6khas no equivalent for nvfp4 online EPLB. WithWIDEEPgone, that test's parametrize list isTRTLLM-only, andTRTLLMGenFusedMoEexplicitly rejects SM120 (fused_moe_trtllm_gen.py: "does not support SM120 and above").WIDEEPwas the only backend that could run this combination on SM120, so the two entries are dropped rather than remapped. The test is additionally gated by@pytest.mark.skip_device_not_contain(["GB200"]), so those entries never executed on that platform.The
gb200_wideep_*/gb300_wideep_*perf entries are unrelated to this deprecation and are kept. There thewideeptoken is only a filename prefix describing the wide-EP deployment shape (dep32+eplb288); the configs undertests/scripts/perf/disaggregated/already setmoe_config.backend: CUTEDSL. An earlier revision of this PR removed those ten entries fromllm_perf_disagg.ymlandllm_perf_multinode.txtby mistake; they are restored.Test plan
Ran the full
TestDeepSeekV4FlashBaseclass on 4x B300 (SM103), wheel built with-a "100-real;103-real", usingTRTLLM_ACCURACY_NO_REFERENCE=1(this model has no registered MMLU/GSM8K reference entries):test_fp8_4gpus_static_eplb[moe_backend=DEEPGEMM]- PASSED (15m09s)test_auto_dtype[moe_backend=DEEPGEMM]- PASSED (7m46s)test_auto_dtype[moe_backend=TRTLLM]- PASSED (7m58s)test_fp8_chunked_prefill- PASSED (7m18s)test_fp8_4gpus_online_eplb[moe_backend=DEEPGEMM]- fails on a pre-existing incompatibility, see belowNote on the online EPLB case
test_fp8_4gpus_online_eplbfails withAssertionError: t.is_contiguous() must be Trueatmoe_load_balancer.py:151. This is not introduced by this PR:WideEPMoE._get_quant_methodreturns the sameDeepSeekFP8BlockScalesFusedMoEMethodDeepGemmthatDeepGemmFusedMoEreturns, so switching the backend does not change this code path._needs_e8m0_resmooth()is true for all Blackwell architectures, and the resulting scale tensor fromtransform_sf_into_required_layoutis non-contiguous, which the online EPLB host tensor sharer rejects.staticEPLB passeslayer_updates_per_iter=0and short-circuits that requirement.@skip_pre_blackwell, so every platform that can run this case hits the same assertion, including the previousWIDEEPparametrization.Fixing it requires deciding whether forcing contiguity is safe for the DeepGEMM scale layout, which is a separate concern from this deprecation.
Dev Engineer Review
WIDEEPdeprecation is applied across code, configuration, examples, documentation, benchmarks, and tests.get_moe_clsandcreate_moe_backendnow rejectWIDEEPwith an explicit error.QA Engineer Review
test_llm_api_pytorch.py.test_disaggregated_serving.pyandtest_e2e.py.test_moe_lora_validator.py.test_deepseek_r1_gen_scaled_from_16_depbecause--scaled-fromlacks a compatible all-to-all selection hook.llm_function_core.txt,llm_function_rtx6k.txt, andl0_dgx_b200.yml.waives.txt.Verdict: sufficient