From 1403aeda9efecd08fb83e66998e48ce97b2f1c8a Mon Sep 17 00:00:00 2001 From: xxi Date: Fri, 31 Jul 2026 09:09:14 +0000 Subject: [PATCH 1/6] [TRTLLM-14822][feat] deprecate WIDEEP MoE backend 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 --- ...loyment-guide-for-deepseek-r1-on-trtllm.md | 6 +-- examples/layer_wise_benchmarks/run.py | 4 +- examples/llm-api/llm_sparse_attention.py | 4 +- examples/llm-api/quickstart_advanced.py | 4 +- examples/longbench/eval_longbench_v1.py | 4 +- examples/models/core/exaone/README.md | 2 +- examples/wide_ep/README.md | 4 +- examples/wide_ep/slurm_scripts/config.yaml | 2 +- .../slurm_scripts/kimi-k2-thinking.yaml | 2 +- .../_torch/models/modeling_deepseekv4.py | 4 +- .../_torch/modules/fused_moe/create_moe.py | 32 +++++------- tensorrt_llm/_torch/peft/lora/validation.py | 2 +- .../tools/layer_wise_benchmarks/runner.py | 1 - .../accuracy/test_disaggregated_serving.py | 4 +- .../defs/accuracy/test_llm_api_pytorch.py | 49 ++++++------------- tests/integration/defs/test_e2e.py | 2 +- .../test_lists/qa/llm_function_core.txt | 4 -- .../test_lists/qa/llm_function_rtx6k.txt | 4 -- .../test_lists/qa/llm_perf_disagg.yml | 7 --- .../test_lists/qa/llm_perf_multinode.txt | 10 ---- .../test_lists/test-db/l0_dgx_b200.yml | 4 -- tests/integration/test_lists/waives.txt | 1 - .../_torch/lora/test_moe_lora_validator.py | 3 +- .../tools/test_layer_wise_benchmarks.py | 5 ++ 24 files changed, 52 insertions(+), 112 deletions(-) diff --git a/docs/source/deployment-guide/deployment-guide-for-deepseek-r1-on-trtllm.md b/docs/source/deployment-guide/deployment-guide-for-deepseek-r1-on-trtllm.md index 1a948213246e..8a993f0b08c3 100644 --- a/docs/source/deployment-guide/deployment-guide-for-deepseek-r1-on-trtllm.md +++ b/docs/source/deployment-guide/deployment-guide-for-deepseek-r1-on-trtllm.md @@ -29,8 +29,8 @@ There are multiple MOE backends inside TensorRT LLM, not all of them supporting | H100/H200 | FP8 | CUTLASS | | B200/GB200 EP<=8 | NVFP4 | CUTLASS, TRTLLM | | B200/GB200 EP<=8 | FP8 | DEEPGEMM | -| GB200 NVL72 EP>8 | NVFP4 | WIDEEP | -| GB200 NVL72 EP>8 | FP8 | WIDEEP without EPLB | +| GB200 NVL72 EP>8 | NVFP4 | CUTEDSL | +| GB200 NVL72 EP>8 | FP8 | DEEPGEMM | The default moe backend is `CUTLASS`, so for the combination which is not supported by `CUTLASS`, one must set the `moe_config.backend` explicitly to run the model. @@ -207,7 +207,7 @@ See the [`TorchLlmArgs` class](https://nvidia.github.io/TensorRT-LLM/llm-api/ref Add the following fields to the YAML configuration file `/tmp/config.yml` to enable wide EP: ```yaml moe_config: - backend: WIDEEP + backend: CUTEDSL max_num_tokens: 9216 load_balancer: # configure online EP balancer num_slots: 288 diff --git a/examples/layer_wise_benchmarks/run.py b/examples/layer_wise_benchmarks/run.py index 5a82610ffbde..a23f485ce718 100644 --- a/examples/layer_wise_benchmarks/run.py +++ b/examples/layer_wise_benchmarks/run.py @@ -53,9 +53,7 @@ def comma_separated_floats(s): parser.add_argument("--load-format", type=str, choices=["AUTO", "DUMMY"]) parser.add_argument("--max-num-tokens", type=int) parser.add_argument("--moe-backend", type=str) -parser.add_argument( - "--moe-backend-for-prefill", type=str, choices=["CUTLASS", "DEEPGEMM", "WIDEEP"] -) +parser.add_argument("--moe-backend-for-prefill", type=str, choices=["CUTLASS", "DEEPGEMM"]) parser.add_argument("--moe-max-num-tokens", type=int) group = parser.add_mutually_exclusive_group() group.add_argument( diff --git a/examples/llm-api/llm_sparse_attention.py b/examples/llm-api/llm_sparse_attention.py index 2b1e7901c535..ae001cfb7b06 100644 --- a/examples/llm-api/llm_sparse_attention.py +++ b/examples/llm-api/llm_sparse_attention.py @@ -176,8 +176,8 @@ def parse_arguments(): type=str, default='CUTLASS', choices=[ - 'CUTLASS', 'TRTLLM', 'VANILLA', 'WIDEEP', - 'DEEPGEMM', 'CUTEDSL', 'TRITON' + 'CUTLASS', 'TRTLLM', 'VANILLA', 'DEEPGEMM', + 'CUTEDSL', 'TRITON' ]) parser.add_argument('--tp_size', type=int, default=1) parser.add_argument('--moe_ep_size', type=int, default=-1) diff --git a/examples/llm-api/quickstart_advanced.py b/examples/llm-api/quickstart_advanced.py index 3757d49486d6..9ad3fb50879c 100644 --- a/examples/llm-api/quickstart_advanced.py +++ b/examples/llm-api/quickstart_advanced.py @@ -87,8 +87,8 @@ def add_llm_args(parser): type=str, default='AUTO', choices=[ - 'AUTO', 'CUTLASS', 'TRTLLM', 'VANILLA', 'WIDEEP', 'DEEPGEMM', - 'CUTEDSL', 'TRITON' + 'AUTO', 'CUTLASS', 'TRTLLM', 'VANILLA', 'DEEPGEMM', 'CUTEDSL', + 'TRITON' ], help= 'MoE backend to use. AUTO selects default backend based on model. It currently doesn\'t always give the best choice for all scenarios. The capabilities of auto selection will be improved in future releases.' diff --git a/examples/longbench/eval_longbench_v1.py b/examples/longbench/eval_longbench_v1.py index 829a5df8c13b..43f5568f2e99 100644 --- a/examples/longbench/eval_longbench_v1.py +++ b/examples/longbench/eval_longbench_v1.py @@ -138,8 +138,8 @@ def parse_arguments() -> argparse.Namespace: type=str, default='CUTLASS', choices=[ - 'CUTLASS', 'TRTLLM', 'VANILLA', 'WIDEEP', - 'DEEPGEMM', 'CUTEDSL', 'TRITON' + 'CUTLASS', 'TRTLLM', 'VANILLA', 'DEEPGEMM', + 'CUTEDSL', 'TRITON' ]) parser.add_argument('--tp_size', type=int, default=1) parser.add_argument('--moe_ep_size', type=int, default=-1) diff --git a/examples/models/core/exaone/README.md b/examples/models/core/exaone/README.md index 5050a4be7822..1c6ddad4c067 100644 --- a/examples/models/core/exaone/README.md +++ b/examples/models/core/exaone/README.md @@ -173,7 +173,7 @@ K-EXAONE supports the following MoE backends: |---------|-------------| | `CUTLASS` | Default backend, optimized for general use cases | | `TRTLLM` | TensorRT-LLM backend using TRT-LLM Gen kernels, optimized for low-latency inference | -| `WIDEEP` | Wide expert parallelism backend for cases where EP size exceeds the number of experts | +| `CUTEDSL` | CuTe DSL backend, used for wide expert parallelism where EP size exceeds the number of experts | You can specify the MoE backend using the `--moe_backend` argument: diff --git a/examples/wide_ep/README.md b/examples/wide_ep/README.md index cce3993b3203..4f2382225afd 100644 --- a/examples/wide_ep/README.md +++ b/examples/wide_ep/README.md @@ -49,7 +49,7 @@ When GDRCopy is installed and the kernel module is loaded, you should be able to An example yaml file to enable wide EP: ```yaml moe_config: - backend: WIDEEP + backend: CUTEDSL max_num_tokens: 9216 load_balancer: num_slots: 288 @@ -59,7 +59,7 @@ moe_config: #### `backend` - MoE backend type, defaults to `CUTLASS`. - - Currently, TensorRT LLM has multiple MoE backends that support wide EP, including `WIDEEP`, `CUTLASS`, `TRTLLM` and `CUTEDSL`. There are on-going efforts to refactor the backends so that we don't necessarily need a specific `WIDEEP` backend, and each other backend will support wide EP functionality. + - TensorRT LLM has multiple MoE backends that support wide EP, including `CUTEDSL`, `CUTLASS` and `TRTLLM`. The dedicated `WIDEEP` backend is deprecated and can no longer be selected; use `CUTEDSL` for large-EP NVFP4 deployments, or `DEEPGEMM` for FP8 block-scale checkpoints on Blackwell. #### `max_num_tokens` diff --git a/examples/wide_ep/slurm_scripts/config.yaml b/examples/wide_ep/slurm_scripts/config.yaml index d34a3cc08ca4..7a9f8c9d1f76 100644 --- a/examples/wide_ep/slurm_scripts/config.yaml +++ b/examples/wide_ep/slurm_scripts/config.yaml @@ -81,7 +81,7 @@ worker_config: free_gpu_memory_fraction: 0.6 dtype: fp8 moe_config: - backend: WIDEEP + backend: CUTEDSL use_low_precision_moe_combine: true load_balancer: num_slots: 288 diff --git a/examples/wide_ep/slurm_scripts/kimi-k2-thinking.yaml b/examples/wide_ep/slurm_scripts/kimi-k2-thinking.yaml index 9fd9cb406226..3ad18f971f38 100644 --- a/examples/wide_ep/slurm_scripts/kimi-k2-thinking.yaml +++ b/examples/wide_ep/slurm_scripts/kimi-k2-thinking.yaml @@ -66,7 +66,7 @@ worker_config: free_gpu_memory_fraction: 0.6 dtype: fp8 moe_config: - backend: WIDEEP + backend: CUTEDSL use_low_precision_moe_combine: true load_balancer: num_slots: 416 diff --git a/tensorrt_llm/_torch/models/modeling_deepseekv4.py b/tensorrt_llm/_torch/models/modeling_deepseekv4.py index b988d1dd7e9a..4f2d25355912 100644 --- a/tensorrt_llm/_torch/models/modeling_deepseekv4.py +++ b/tensorrt_llm/_torch/models/modeling_deepseekv4.py @@ -1515,7 +1515,6 @@ def __init__( CutlassFusedMoE, TritonFusedMoE, TRTLLMGenFusedMoE, - WideEPMoE, DeepGemmFusedMoE, ) # NVFP4 routed-expert path: the TRTLLM-Gen fp4-block-scale fused-MoE @@ -1523,8 +1522,7 @@ def __init__( # swiglu_limit is supplied; drop the limit there until the cubin # gains a no-bias clamp variant. MXFP4 variants are unaffected. kernel_requires_bias_for_swiglu_limit = ( - moe_cls in (TRTLLMGenFusedMoE, WideEPMoE) - and experts_quant_config.quant_mode.has_nvfp4() + moe_cls is TRTLLMGenFusedMoE and experts_quant_config.quant_mode.has_nvfp4() ) # DeepSeek-V4 supplies a uniform scalar limit. The TRTLLM-Gen FP8 # path consumes it directly and rejects the redundant tensor. diff --git a/tensorrt_llm/_torch/modules/fused_moe/create_moe.py b/tensorrt_llm/_torch/modules/fused_moe/create_moe.py index 1d22fc1ac47c..25113ac2bcba 100644 --- a/tensorrt_llm/_torch/modules/fused_moe/create_moe.py +++ b/tensorrt_llm/_torch/modules/fused_moe/create_moe.py @@ -26,6 +26,12 @@ from .moe_load_balancer import get_moe_load_balancer from .routing import BaseMoeRoutingMethod +WIDEEP_DEPRECATION_MESSAGE = ( + "The WIDEEP MoE backend is deprecated and can no longer be selected. Wide " + "expert parallelism and EPLB are supported by the other backends: use " + "DEEPGEMM for FP8 block-scale checkpoints, or TRTLLM / CUTEDSL / CUTLASS " + "otherwise.") + def _get_pretrained_megamoe_capability_args( model_config: ModelConfig) -> Dict[str, Optional[object]]: @@ -154,7 +160,7 @@ def get_moe_cls( ) return CutlassFusedMoE elif moe_backend.upper() == "WIDEEP": - return WideEPMoE + raise ValueError(WIDEEP_DEPRECATION_MESSAGE) elif moe_backend.upper() == "TRITON": return TritonFusedMoE elif moe_backend.upper() == "MEGAMOE_DEEPGEMM": @@ -288,6 +294,9 @@ def create_moe_backend( Returns: MoE: MoE backend instance """ + if moe_cls is WideEPMoE: + raise ValueError(WIDEEP_DEPRECATION_MESSAGE) + # Get parameters from pretrained_config if not explicitly provided pretrained_config = model_config.pretrained_config if num_experts is None: @@ -310,7 +319,6 @@ def create_moe_backend( moe_load_balancer = get_moe_load_balancer() if moe_load_balancer is not None: supported_load_balancer_backends = ( - WideEPMoE, CutlassFusedMoE, TRTLLMGenFusedMoE, CuteDslFusedMoE, @@ -336,7 +344,7 @@ def create_moe_backend( if swiglu_limit is not None: assert moe_cls in [ - CutlassFusedMoE, TritonFusedMoE, TRTLLMGenFusedMoE, WideEPMoE, + CutlassFusedMoE, TritonFusedMoE, TRTLLMGenFusedMoE, DeepGemmFusedMoE, MegaMoECuteDsl ], f"swiglu_limit is not supported in {moe_cls.__name__}." @@ -344,7 +352,7 @@ def create_moe_backend( # MegaMoECuteDsl uses the scalar only as a fallback when no per-expert # tensor limit is given (see the MegaMoE branch below). assert moe_cls in [ - CutlassFusedMoE, TRTLLMGenFusedMoE, WideEPMoE, DeepGemmFusedMoE, + CutlassFusedMoE, TRTLLMGenFusedMoE, DeepGemmFusedMoE, MegaMoEDeepGemm, CuteDslFusedMoE, MegaMoECuteDsl ], f"swiglu_limit_scalar is not supported in {moe_cls.__name__}." @@ -394,22 +402,6 @@ def create_moe_backend( init_load_balancer=init_load_balancer, activation_type=activation_type, ) - elif moe_cls == WideEPMoE: - return moe_cls( - routing_method=routing_method, - num_experts=num_experts, - hidden_size=hidden_size, - intermediate_size=intermediate_size, - dtype=dtype, - reduce_results=reduce_results, - model_config=model_config, - aux_stream_dict=aux_stream_dict, - weight_loading_mode=weight_loading_mode, - apply_router_weight_on_input=apply_router_weight_on_input, - layer_idx=layer_idx, - swiglu_limit=swiglu_limit, - swiglu_limit_scalar=swiglu_limit_scalar, - activation_type=activation_type) elif moe_cls == VanillaMoE: assert not apply_router_weight_on_input, "apply_router_weight_on_input is not supported in VanillaMoE." diff --git a/tensorrt_llm/_torch/peft/lora/validation.py b/tensorrt_llm/_torch/peft/lora/validation.py index d773f52e32ec..bbb4a7c3a597 100644 --- a/tensorrt_llm/_torch/peft/lora/validation.py +++ b/tensorrt_llm/_torch/peft/lora/validation.py @@ -93,7 +93,7 @@ def check_moe_lora_supported( Args: moe_backend_name: The resolved `moe_backend` string (e.g. "CUTLASS", - "WIDEEP", "TRTLLM"). Comparison is case-insensitive. + "CUTEDSL", "TRTLLM"). Comparison is case-insensitive. lora_config: The model's `LoraConfig`, or None. quant_config: The model's `QuantConfig`, or None. We only reject when the layer is actually quantized (`quant_mode.has_any_quant`). diff --git a/tensorrt_llm/tools/layer_wise_benchmarks/runner.py b/tensorrt_llm/tools/layer_wise_benchmarks/runner.py index fa025b5fa144..37e1bf239e72 100644 --- a/tensorrt_llm/tools/layer_wise_benchmarks/runner.py +++ b/tensorrt_llm/tools/layer_wise_benchmarks/runner.py @@ -702,7 +702,6 @@ def replace_routing_method_ctx(self, balance_method: BalanceMethod, balance_rati "CUTLASS", "DEEPGEMM", "TRTLLM", - "WIDEEP", ]: raise NotImplementedError( f'Not support replace routing method for moe_backend "{self.model_config.moe_backend}",' diff --git a/tests/integration/defs/accuracy/test_disaggregated_serving.py b/tests/integration/defs/accuracy/test_disaggregated_serving.py index 806d706cd5f1..749ed07f50a0 100644 --- a/tests/integration/defs/accuracy/test_disaggregated_serving.py +++ b/tests/integration/defs/accuracy/test_disaggregated_serving.py @@ -2484,7 +2484,7 @@ class TestDeepSeekV4Flash(LlmapiAccuracyTestHarness): def test_auto_dtype(self): # Disagg smoke test: CTX TP=2 + GEN TP=2 = 4 GPUs. # NVFP4 weights ~71 GB/rank at TP=2, leaving ~107 GB for KV on B200. - # TRTLLM backend required (WIDEEP lacks MXFP4 support for V4-Flash). + # TRTLLM backend required: it is the backend supporting V4-Flash MXFP4. # V4 uses pure-Python KVCacheManagerV2; needs Python transceiver. # NIXL (not DEFAULT) skips the TRTLLM_USE_UCX_KVCACHE=1 fallback. cache_transceiver_config = { @@ -2603,7 +2603,7 @@ def test_auto_dtype(self): # Disagg smoke test: CTX TP=2 + GEN TP=2 = 4 GPUs. # FP8 weights ~71 GB/rank at TP=4 → ~142 GB/rank at TP=2; requires # ≥140 GB per GPU (fits on B300 288 GB, tight on B200 178 GB). - # TRTLLM backend: WIDEEP's FP8 block-scale path is Hopper-only. + # TRTLLM backend: the CUTLASS FP8 block-scale path is Hopper-only. # Compact batching keeps KV cache ~1 GB/rank (default ~100 GB requires fully-clean GPU memory). # V4 uses pure-Python KVCacheManagerV2; needs Python transceiver. # NIXL (not DEFAULT) skips the TRTLLM_USE_UCX_KVCACHE=1 fallback. diff --git a/tests/integration/defs/accuracy/test_llm_api_pytorch.py b/tests/integration/defs/accuracy/test_llm_api_pytorch.py index f0a089120b1d..b8e4ff8d096b 100644 --- a/tests/integration/defs/accuracy/test_llm_api_pytorch.py +++ b/tests/integration/defs/accuracy/test_llm_api_pytorch.py @@ -2258,7 +2258,7 @@ def test_fp8_block_scales_4gpus_static_eplb(self): layer_updates_per_iter=0) pytorch_backend_options = dict(cuda_graph_config=CudaGraphConfig(), moe_config=MoeConfig( - backend="WIDEEP", + backend="DEEPGEMM", load_balancer=eplb_config)) with LLM(f"{llm_models_root()}/DeepSeek-V3-Lite/fp8", tensor_parallel_size=4, @@ -2271,7 +2271,7 @@ def test_fp8_block_scales_4gpus_static_eplb(self): @pytest.mark.skip_less_device(4) @pytest.mark.skip_device_not_contain(["GB200"]) - @parametrize_with_ids("moe_backend", ["WIDEEP", "CUTLASS", "TRTLLM"]) + @parametrize_with_ids("moe_backend", ["CUTLASS", "TRTLLM"]) @parametrize_with_ids("mtp_nextn", [0, 2]) def test_bfloat16_4gpus_online_eplb(self, moe_backend, mtp_nextn): kv_cache_config = KvCacheConfig(free_gpu_memory_fraction=0.7) @@ -2296,7 +2296,7 @@ def test_bfloat16_4gpus_online_eplb(self, moe_backend, mtp_nextn): @pytest.mark.skip_less_device(4) @pytest.mark.skip_device_not_contain(["GB200"]) - @parametrize_with_ids("moe_backend", ["WIDEEP", "TRTLLM"]) + @parametrize_with_ids("moe_backend", ["TRTLLM"]) @parametrize_with_ids("fp8kv", [True, False]) def test_nvfp4_4gpus_online_eplb(self, moe_backend, fp8kv): kv_cache_config = KvCacheConfig(free_gpu_memory_fraction=0.7) @@ -3890,9 +3890,8 @@ class TestDeepSeekV4Flash(LlmapiAccuracyTestHarness): def test_auto_dtype(self): # Aggregate (non-disagg, non-EPLB) coverage. NVFP4 weights are ~71 # GB/rank at TP=2, ~36 GB/rank at TP=4 — TP=4 fits comfortably on - # 4x B200 178GB. TRTLLM backend required because V4-Flash MXFP4 - # routed experts are unsupported by WIDEEP (raises "Unsupported - # quantization mode: [65536]"). + # 4x B200 178GB. TRTLLM backend is pinned because it is the backend + # that supports V4-Flash MXFP4 routed experts. kv_cache_config = KvCacheConfig(free_gpu_memory_fraction=0.5) with LLM(self.MODEL_PATH, tensor_parallel_size=4, @@ -3909,18 +3908,7 @@ def test_auto_dtype(self): task.evaluate(llm) @pytest.mark.skip_less_mpi_world_size(4) - @parametrize_with_ids("moe_backend", [ - pytest.param( - "WIDEEP", - marks=pytest.mark.skip( - reason= - "V4-Flash MXFP4 routed experts: WIDEEP _get_quant_method has " - "no MXFP4 branch (raises 'Unsupported quantization mode: " - "[65536]'). Re-enable once fused_moe_wide_ep.py supports MXFP4." - )), - "TRTLLM", - "MEGAMOE_DEEPGEMM", - ]) + @parametrize_with_ids("moe_backend", ["TRTLLM", "MEGAMOE_DEEPGEMM"]) def test_nvfp4_4gpus_static_eplb(self, moe_backend): eplb_config = _make_deepseekv4_eplb_config(self.MODEL_PATH, layer_updates_per_iter=0, @@ -3929,17 +3917,7 @@ def test_nvfp4_4gpus_static_eplb(self, moe_backend): eplb_config) @pytest.mark.skip_less_mpi_world_size(4) - @parametrize_with_ids("moe_backend", [ - pytest.param( - "WIDEEP", - marks=pytest.mark.skip( - reason= - "V4-Flash MXFP4 routed experts: WIDEEP _get_quant_method has " - "no MXFP4 branch (raises 'Unsupported quantization mode: " - "[65536]'). Re-enable once fused_moe_wide_ep.py supports MXFP4." - )), - "TRTLLM", - ]) + @parametrize_with_ids("moe_backend", ["TRTLLM"]) @parametrize_with_ids("mtp_nextn", [0, 1]) def test_nvfp4_4gpus_online_eplb(self, moe_backend, mtp_nextn): eplb_config = _make_deepseekv4_eplb_config(self.MODEL_PATH, @@ -4042,7 +4020,7 @@ class TestDeepSeekV4FlashBase(LlmapiAccuracyTestHarness): MODEL_PATH = f"{llm_models_root()}/DeepSeek-V4-Flash-Base" @pytest.mark.skip_less_mpi_world_size(4) - @parametrize_with_ids("moe_backend", ["WIDEEP", "TRTLLM"]) + @parametrize_with_ids("moe_backend", ["DEEPGEMM", "TRTLLM"]) def test_auto_dtype(self, moe_backend): # Aggregate (non-disagg, non-EPLB) smoke test. FP8 weights ~71 GB/rank # at TP=4 — fits on 4x B300 (~288 GB/GPU). 1-sample smoke. CUTLASS is @@ -4069,7 +4047,7 @@ def test_fp8_chunked_prefill(self): with LLM(self.MODEL_PATH, tensor_parallel_size=4, moe_expert_parallel_size=4, - moe_config=MoeConfig(backend="WIDEEP"), + moe_config=MoeConfig(backend="DEEPGEMM"), cuda_graph_config=CudaGraphConfig( max_batch_size=DEEPSEEKV4_TEST_MAX_BATCH_SIZE, enable_padding=True), @@ -4083,10 +4061,11 @@ def test_fp8_chunked_prefill(self): task.evaluate(llm, is_integration_test=True) # CUTLASS is omitted: V4 Flash-Base FP8 block-scale weights take a - # Hopper-only kernel path (CutlassFp8BlockScaleGemmRunner::moeGemm) that - # fails on Blackwell. WIDEEP avoids that path and works on B200/B300. + # Hopper-only wgmma kernel path (CutlassFp8BlockScaleGemmRunner::moeGemm) + # that has no SM100/SM103 implementation. DEEPGEMM routes the same weights + # through DeepGEMM instead and works on B200/B300. @pytest.mark.skip_less_mpi_world_size(4) - @parametrize_with_ids("moe_backend", ["WIDEEP"]) + @parametrize_with_ids("moe_backend", ["DEEPGEMM"]) def test_fp8_4gpus_static_eplb(self, moe_backend): eplb_config = _make_deepseekv4_eplb_config(self.MODEL_PATH, layer_updates_per_iter=0, @@ -4095,7 +4074,7 @@ def test_fp8_4gpus_static_eplb(self, moe_backend): eplb_config) @pytest.mark.skip_less_mpi_world_size(4) - @parametrize_with_ids("moe_backend", ["WIDEEP"]) + @parametrize_with_ids("moe_backend", ["DEEPGEMM"]) def test_fp8_4gpus_online_eplb(self, moe_backend): eplb_config = _make_deepseekv4_eplb_config(self.MODEL_PATH, layer_updates_per_iter=2, diff --git a/tests/integration/defs/test_e2e.py b/tests/integration/defs/test_e2e.py index 215731702435..e8cf812b26cc 100644 --- a/tests/integration/defs/test_e2e.py +++ b/tests/integration/defs/test_e2e.py @@ -1354,7 +1354,7 @@ def test_deepseek_r1_mtp_bench(llm_root, llm_venv): }, "enable_attention_dp": True, "moe_config": { - "backend": "WIDEEP", + "backend": "CUTLASS", }, "cuda_graph_config": { "enable_padding": True, diff --git a/tests/integration/test_lists/qa/llm_function_core.txt b/tests/integration/test_lists/qa/llm_function_core.txt index 9fde8bb9d1fa..5977ed874461 100644 --- a/tests/integration/test_lists/qa/llm_function_core.txt +++ b/tests/integration/test_lists/qa/llm_function_core.txt @@ -282,9 +282,7 @@ accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus[tp4-mt accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus[tp4-mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=True] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_kv_cache_aware_routing[mtp_nextn=0] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_kv_cache_aware_routing[mtp_nextn=2] -accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_online_eplb[mtp_nextn=0-moe_backend=WIDEEP] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_online_eplb[mtp_nextn=2-moe_backend=CUTLASS] -accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_online_eplb[mtp_nextn=2-moe_backend=WIDEEP] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_python_scheduler[ep4-mtp_nextn=0] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_python_scheduler[ep4-mtp_nextn=2] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_python_scheduler[tp4-mtp_nextn=0] @@ -438,9 +436,7 @@ accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus[moe_backe accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus[moe_backend=TRTLLM-mtp_nextn=0-tp4-fp8kv=True-attention_dp=True-cuda_graph=True-overlap_scheduler=True-low_precision_combine=True-torch_compile=False] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus[moe_backend=TRTLLM-mtp_nextn=2-ep4-fp8kv=True-attention_dp=True-cuda_graph=True-overlap_scheduler=True-low_precision_combine=False-torch_compile=False] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus[moe_backend=TRTLLM-mtp_nextn=2-tp4-fp8kv=True-attention_dp=True-cuda_graph=True-overlap_scheduler=True-low_precision_combine=False-torch_compile=False] -accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus_online_eplb[fp8kv=False-moe_backend=WIDEEP] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus_online_eplb[fp8kv=True-moe_backend=TRTLLM] -accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus_online_eplb[fp8kv=True-moe_backend=WIDEEP] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_batch_waiting[batch_wait_timeout_iters=10-batch_wait_max_tokens_ratio=1.0-mtp_nextn=0-fp8kv=False-cuda_graph=False-overlap_scheduler=False-torch_compile=False-v2_kv_cache=False] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_batch_waiting[batch_wait_timeout_iters=10-batch_wait_max_tokens_ratio=1.0-mtp_nextn=0-fp8kv=False-cuda_graph=False-overlap_scheduler=False-torch_compile=False-v2_kv_cache=True] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_batch_waiting[batch_wait_timeout_iters=10-batch_wait_max_tokens_ratio=1.0-mtp_nextn=0-fp8kv=False-cuda_graph=False-overlap_scheduler=False-torch_compile=True-v2_kv_cache=False] diff --git a/tests/integration/test_lists/qa/llm_function_rtx6k.txt b/tests/integration/test_lists/qa/llm_function_rtx6k.txt index 4260a5328d77..115bc7cac7b2 100644 --- a/tests/integration/test_lists/qa/llm_function_rtx6k.txt +++ b/tests/integration/test_lists/qa/llm_function_rtx6k.txt @@ -52,8 +52,6 @@ accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mtp_nextn=2- accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=False-enable_chunked_prefill=False-v2_kv_cache=True] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=True-enable_chunked_prefill=False-v2_kv_cache=False] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=True-enable_chunked_prefill=False-v2_kv_cache=True] -accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_online_eplb[mtp_nextn=0-moe_backend=WIDEEP] -accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_online_eplb[mtp_nextn=2-moe_backend=WIDEEP] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_guided_decoding[llguidance-mtp_nextn=0] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_guided_decoding[llguidance-mtp_nextn=2] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_guided_decoding[xgrammar-mtp_nextn=0] @@ -82,8 +80,6 @@ accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4[moe_backend=CUT accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4[moe_backend=CUTLASS-mtp_nextn=2-fp8kv=True-attention_dp=False-cuda_graph=False-overlap_scheduler=False-torch_compile=True] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4[moe_backend=CUTLASS-mtp_nextn=2-fp8kv=True-attention_dp=False-cuda_graph=True-overlap_scheduler=True-torch_compile=False] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4[moe_backend=CUTLASS-mtp_nextn=2-fp8kv=True-attention_dp=False-cuda_graph=True-overlap_scheduler=True-torch_compile=True] -accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus_online_eplb[fp8kv=False-moe_backend=WIDEEP] -accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus_online_eplb[fp8kv=True-moe_backend=WIDEEP] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_batch_waiting[batch_wait_timeout_iters=10-batch_wait_max_tokens_ratio=1.0-mtp_nextn=0-fp8kv=False-cuda_graph=False-overlap_scheduler=False-torch_compile=False-v2_kv_cache=False] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_batch_waiting[batch_wait_timeout_iters=10-batch_wait_max_tokens_ratio=1.0-mtp_nextn=0-fp8kv=False-cuda_graph=False-overlap_scheduler=False-torch_compile=False-v2_kv_cache=True] accuracy/test_llm_api_pytorch.py::TestGPTOSS::test_dflash diff --git a/tests/integration/test_lists/qa/llm_perf_disagg.yml b/tests/integration/test_lists/qa/llm_perf_disagg.yml index 9acefe684999..41fe48ed7ca2 100644 --- a/tests/integration/test_lists/qa/llm_perf_disagg.yml +++ b/tests/integration/test_lists/qa/llm_perf_disagg.yml @@ -42,12 +42,8 @@ llm_perf_disagg: - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_qwen3-235b-fp4_8k1k_con1024_ctx1_tp1_gen1_dep8_eplb0_mtp0_ccb-UCX] TIMEOUT (120) - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_qwen3-235b-fp4_8k1k_con1_ctx1_tp1_gen1_tep4_eplb0_mtp0_ccb-NIXL] TIMEOUT (120) - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_qwen3-235b-fp4_8k1k_con64_ctx1_tp1_gen1_tep4_eplb0_mtp0_ccb-NIXL] TIMEOUT (120) - # GB200 wideep - - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] TIMEOUT (120) - - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_deepseek-v32-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] TIMEOUT (120) # GB200 accuracy cases # GB200 stress cases - - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_stress-deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_ccb-NIXL] TIMEOUT (120) - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_stress-gpt-oss-120b-fp4_8k1k_ctx1_tp1_gen1_tp4_eplb0_eagle3_ccb-NIXL] TIMEOUT (120) # GB200 aggregated ctx_only # GB200 aggregated gen_only @@ -87,8 +83,5 @@ llm_perf_disagg: - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_kimi-k25-thinking-fp4_8k1k_con4096_ctx1_dep4_gen1_dep16_eplb0_mtp0_ccb-UCX] TIMEOUT (120) - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_kimi-k25-thinking-fp4_8k1k_con4_ctx1_dep4_gen1_tep8_eplb0_mtp3_ccb-NIXL] TIMEOUT (120) # GB300 Qwen3-235B - # GB300 wideep - - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_wideep_deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] TIMEOUT (120) - - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_wideep_deepseek-v32-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] TIMEOUT (120) # GB300 aggregated ctx_only # GB300 aggregated gen_only diff --git a/tests/integration/test_lists/qa/llm_perf_multinode.txt b/tests/integration/test_lists/qa/llm_perf_multinode.txt index 9eb5adbcf384..d4e13e353927 100644 --- a/tests/integration/test_lists/qa/llm_perf_multinode.txt +++ b/tests/integration/test_lists/qa/llm_perf_multinode.txt @@ -89,19 +89,9 @@ perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_kimi-k25-thinking-fp4_8k1k_c # GB300 Qwen3-235B -# wideep multi-node -# GB200 wideep -perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] -perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_deepseek-v32-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] - -# GB300 wideep -perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_wideep_deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] -perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_wideep_deepseek-v32-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] - # accuracy cases # stress cases -perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_stress-deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_ccb-NIXL] perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_stress-gpt-oss-120b-fp4_8k1k_ctx1_tp1_gen1_tp4_eplb0_eagle3_ccb-NIXL] # aggregated multi-node (ctx_only and gen_only reuse disagg config yamls) diff --git a/tests/integration/test_lists/test-db/l0_dgx_b200.yml b/tests/integration/test_lists/test-db/l0_dgx_b200.yml index 92267dafe500..bc13bd035e76 100644 --- a/tests/integration/test_lists/test-db/l0_dgx_b200.yml +++ b/tests/integration/test_lists/test-db/l0_dgx_b200.yml @@ -194,10 +194,6 @@ l0_dgx_b200: - accuracy/test_disaggregated_serving.py::TestNemotron3Super120B::test_auto_dtype[mtp_nextn=3-block_reuse=True-use_py_transceiver=False] TIMEOUT (60) - accuracy/test_disaggregated_serving.py::TestNemotron3Super120B::test_auto_dtype[mtp_nextn=0-block_reuse=False-use_py_transceiver=False] TIMEOUT (60) - accuracy/test_disaggregated_serving.py::TestGLM52NVFP4::test_nvfp4_nixl[cache_mgr_v1] TIMEOUT (60) - # DeepSeek-V4 EPLB pre-merge sanity (uncomment once DeepSeek-V4-Flash/Flash-Base - # checkpoints are staged under llm_models_root()). - # - accuracy/test_llm_api_pytorch.py::TestDeepSeekV4Flash::test_nvfp4_8gpus_static_eplb[moe_backend=WIDEEP] TIMEOUT (120) - # - accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_fp8_8gpus_static_eplb[moe_backend=WIDEEP] TIMEOUT (120) - accuracy/test_disaggregated_serving.py::TestNemotron3Super120B::test_ctx_dp2_gen_tp4 TIMEOUT (60) - accuracy/test_disaggregated_serving.py::TestQwen3NextInstruct::test_auto_dtype[use_py_transceiver=True] TIMEOUT (60) - accuracy/test_disaggregated_serving.py::TestQwen3NextInstruct::test_auto_dtype[use_py_transceiver=False] TIMEOUT (60) diff --git a/tests/integration/test_lists/waives.txt b/tests/integration/test_lists/waives.txt index 2fc51b9b3f87..5a4df968870a 100644 --- a/tests/integration/test_lists/waives.txt +++ b/tests/integration/test_lists/waives.txt @@ -36,7 +36,6 @@ accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus[pp4-mt accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus[tp2pp2-mtp_nextn=0-attention_dp=False-cuda_graph=False-overlap_scheduler=False-torch_compile=True] SKIP (https://nvbugs/6428096) accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus[tp2pp2-mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=True] SKIP (https://nvbugs/6198774) accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus[tp4-mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=True] SKIP (https://nvbugs/6198774) -accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_online_eplb[mtp_nextn=2-moe_backend=WIDEEP] SKIP (https://nvbugs/6313993) accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_python_scheduler[mtp_nextn=0-attention_dp=False-cuda_graph=False-overlap_scheduler=False-enable_chunked_prefill=True] SKIP (https://nvbugs/6388139) accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_python_scheduler[mtp_nextn=2-attention_dp=False-cuda_graph=False-overlap_scheduler=False-enable_chunked_prefill=True] SKIP (https://nvbugs/6507095) accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_fp8_block_scales_4gpus[pp4-mtp_nextn=0-fp8kv=False-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=False-sampler_async_worker=False] SKIP (https://nvbugs/6427411) diff --git a/tests/unittest/_torch/lora/test_moe_lora_validator.py b/tests/unittest/_torch/lora/test_moe_lora_validator.py index 993f65a67aa8..f2d4c0175eea 100644 --- a/tests/unittest/_torch/lora/test_moe_lora_validator.py +++ b/tests/unittest/_torch/lora/test_moe_lora_validator.py @@ -56,7 +56,7 @@ def test_has_moe_lora_targets_each_module(name): def test_check_no_lora_is_noop(): # No LoRA at all; validator must not raise regardless of backend/quant. check_moe_lora_supported( - moe_backend_name="WIDEEP", + moe_backend_name="CUTEDSL", lora_config=None, quant_config=_FakeQuantConfig(_FP8_BLOCK_SCALE), ) @@ -117,7 +117,6 @@ def test_check_moe_lora_rejects_fp8_block_scale(): @pytest.mark.parametrize( "backend", [ - "WIDEEP", "TRITON", "DEEPGEMM", "VANILLA", diff --git a/tests/unittest/tools/test_layer_wise_benchmarks.py b/tests/unittest/tools/test_layer_wise_benchmarks.py index 2cc05a5bf0f7..c3e4eee83030 100644 --- a/tests/unittest/tools/test_layer_wise_benchmarks.py +++ b/tests/unittest/tools/test_layer_wise_benchmarks.py @@ -243,6 +243,11 @@ def test_deepseek_v32_ctx_dep(llm_root, world_size): # The pinned DeepSeek FP4 checkpoint requires SM100+. +@pytest.mark.skip( + reason="--scaled-from rewrites WideEPMoE.select_alltoall_method_type, which is " + "the only alltoall-selection hook it patches. The WIDEEP backend is deprecated, " + "so weak scaling has no equivalent backend until the hook is generalized." +) @skip_pre_blackwell @pytest.mark.parametrize("world_size", [4]) def test_deepseek_r1_gen_scaled_from_16_dep(llm_root, world_size): From 3b41cdaf35eb39c5730573227874d2258aa3e89c Mon Sep 17 00:00:00 2001 From: xxi Date: Mon, 3 Aug 2026 01:36:01 +0000 Subject: [PATCH 2/6] [TRTLLM-14822][fix] restore test list coverage lost in WIDEEP deprecation 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 --- tests/integration/test_lists/qa/llm_function_core.txt | 7 +++++-- tests/integration/test_lists/qa/llm_function_rtx6k.txt | 2 ++ tests/integration/test_lists/qa/llm_perf_disagg.yml | 7 +++++++ tests/integration/test_lists/qa/llm_perf_multinode.txt | 10 ++++++++++ tests/integration/test_lists/test-db/l0_dgx_b200.yml | 4 ++++ 5 files changed, 28 insertions(+), 2 deletions(-) diff --git a/tests/integration/test_lists/qa/llm_function_core.txt b/tests/integration/test_lists/qa/llm_function_core.txt index 5977ed874461..e49063e103ca 100644 --- a/tests/integration/test_lists/qa/llm_function_core.txt +++ b/tests/integration/test_lists/qa/llm_function_core.txt @@ -282,7 +282,9 @@ accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus[tp4-mt accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus[tp4-mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=True] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_kv_cache_aware_routing[mtp_nextn=0] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_kv_cache_aware_routing[mtp_nextn=2] +accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_online_eplb[mtp_nextn=0-moe_backend=CUTLASS] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_online_eplb[mtp_nextn=2-moe_backend=CUTLASS] +accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_online_eplb[mtp_nextn=2-moe_backend=TRTLLM] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_python_scheduler[ep4-mtp_nextn=0] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_python_scheduler[ep4-mtp_nextn=2] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_python_scheduler[tp4-mtp_nextn=0] @@ -436,6 +438,7 @@ accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus[moe_backe accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus[moe_backend=TRTLLM-mtp_nextn=0-tp4-fp8kv=True-attention_dp=True-cuda_graph=True-overlap_scheduler=True-low_precision_combine=True-torch_compile=False] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus[moe_backend=TRTLLM-mtp_nextn=2-ep4-fp8kv=True-attention_dp=True-cuda_graph=True-overlap_scheduler=True-low_precision_combine=False-torch_compile=False] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus[moe_backend=TRTLLM-mtp_nextn=2-tp4-fp8kv=True-attention_dp=True-cuda_graph=True-overlap_scheduler=True-low_precision_combine=False-torch_compile=False] +accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus_online_eplb[fp8kv=False-moe_backend=TRTLLM] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus_online_eplb[fp8kv=True-moe_backend=TRTLLM] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_batch_waiting[batch_wait_timeout_iters=10-batch_wait_max_tokens_ratio=1.0-mtp_nextn=0-fp8kv=False-cuda_graph=False-overlap_scheduler=False-torch_compile=False-v2_kv_cache=False] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_batch_waiting[batch_wait_timeout_iters=10-batch_wait_max_tokens_ratio=1.0-mtp_nextn=0-fp8kv=False-cuda_graph=False-overlap_scheduler=False-torch_compile=False-v2_kv_cache=True] @@ -986,8 +989,8 @@ llmapi/test_llm_api_pytorch_moe_lora.py::test_mixtral_moe_routed_expert_fp8_mult accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_cute_dsl_nvfp4[fp8kv=False-attention_dp=False-cuda_graph=False-overlap_scheduler=False-torch_compile=False] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_cute_dsl_nvfp4[fp8kv=False-attention_dp=False-cuda_graph=False-overlap_scheduler=False-torch_compile=True] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_cute_dsl_nvfp4_4gpus[tp4-fp8kv=False-attention_dp=False-cuda_graph=False-overlap_scheduler=False-torch_compile=False] -accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_auto_dtype[moe_backend=WIDEEP] +accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_auto_dtype[moe_backend=DEEPGEMM] accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_auto_dtype[moe_backend=TRTLLM] accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_fp8_chunked_prefill -accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_fp8_4gpus_static_eplb[moe_backend=WIDEEP] +accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_fp8_4gpus_static_eplb[moe_backend=DEEPGEMM] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_cute_dsl_nvfp4_4gpus[tp4-fp8kv=False-attention_dp=False-cuda_graph=False-overlap_scheduler=False-torch_compile=True] diff --git a/tests/integration/test_lists/qa/llm_function_rtx6k.txt b/tests/integration/test_lists/qa/llm_function_rtx6k.txt index 115bc7cac7b2..016ed64513bf 100644 --- a/tests/integration/test_lists/qa/llm_function_rtx6k.txt +++ b/tests/integration/test_lists/qa/llm_function_rtx6k.txt @@ -52,6 +52,8 @@ accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mtp_nextn=2- accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=False-enable_chunked_prefill=False-v2_kv_cache=True] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=True-enable_chunked_prefill=False-v2_kv_cache=False] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=True-enable_chunked_prefill=False-v2_kv_cache=True] +accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_online_eplb[mtp_nextn=0-moe_backend=CUTLASS] +accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus_online_eplb[mtp_nextn=2-moe_backend=CUTLASS] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_guided_decoding[llguidance-mtp_nextn=0] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_guided_decoding[llguidance-mtp_nextn=2] accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_guided_decoding[xgrammar-mtp_nextn=0] diff --git a/tests/integration/test_lists/qa/llm_perf_disagg.yml b/tests/integration/test_lists/qa/llm_perf_disagg.yml index 41fe48ed7ca2..9acefe684999 100644 --- a/tests/integration/test_lists/qa/llm_perf_disagg.yml +++ b/tests/integration/test_lists/qa/llm_perf_disagg.yml @@ -42,8 +42,12 @@ llm_perf_disagg: - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_qwen3-235b-fp4_8k1k_con1024_ctx1_tp1_gen1_dep8_eplb0_mtp0_ccb-UCX] TIMEOUT (120) - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_qwen3-235b-fp4_8k1k_con1_ctx1_tp1_gen1_tep4_eplb0_mtp0_ccb-NIXL] TIMEOUT (120) - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_qwen3-235b-fp4_8k1k_con64_ctx1_tp1_gen1_tep4_eplb0_mtp0_ccb-NIXL] TIMEOUT (120) + # GB200 wideep + - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] TIMEOUT (120) + - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_deepseek-v32-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] TIMEOUT (120) # GB200 accuracy cases # GB200 stress cases + - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_stress-deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_ccb-NIXL] TIMEOUT (120) - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_stress-gpt-oss-120b-fp4_8k1k_ctx1_tp1_gen1_tp4_eplb0_eagle3_ccb-NIXL] TIMEOUT (120) # GB200 aggregated ctx_only # GB200 aggregated gen_only @@ -83,5 +87,8 @@ llm_perf_disagg: - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_kimi-k25-thinking-fp4_8k1k_con4096_ctx1_dep4_gen1_dep16_eplb0_mtp0_ccb-UCX] TIMEOUT (120) - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_kimi-k25-thinking-fp4_8k1k_con4_ctx1_dep4_gen1_tep8_eplb0_mtp3_ccb-NIXL] TIMEOUT (120) # GB300 Qwen3-235B + # GB300 wideep + - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_wideep_deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] TIMEOUT (120) + - perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_wideep_deepseek-v32-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] TIMEOUT (120) # GB300 aggregated ctx_only # GB300 aggregated gen_only diff --git a/tests/integration/test_lists/qa/llm_perf_multinode.txt b/tests/integration/test_lists/qa/llm_perf_multinode.txt index d4e13e353927..9eb5adbcf384 100644 --- a/tests/integration/test_lists/qa/llm_perf_multinode.txt +++ b/tests/integration/test_lists/qa/llm_perf_multinode.txt @@ -89,9 +89,19 @@ perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_kimi-k25-thinking-fp4_8k1k_c # GB300 Qwen3-235B +# wideep multi-node +# GB200 wideep +perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] +perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_deepseek-v32-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] + +# GB300 wideep +perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_wideep_deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] +perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb300_wideep_deepseek-v32-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_con1024_ccb-NIXL] + # accuracy cases # stress cases +perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_wideep_stress-deepseek-r1-fp4_8k1k_ctx2_gen1_dep32_bs128_eplb288_mtp3_ccb-NIXL] perf/test_perf_sanity.py::test_e2e[disagg-e2e-gb200_stress-gpt-oss-120b-fp4_8k1k_ctx1_tp1_gen1_tp4_eplb0_eagle3_ccb-NIXL] # aggregated multi-node (ctx_only and gen_only reuse disagg config yamls) diff --git a/tests/integration/test_lists/test-db/l0_dgx_b200.yml b/tests/integration/test_lists/test-db/l0_dgx_b200.yml index bc13bd035e76..438cedc309a2 100644 --- a/tests/integration/test_lists/test-db/l0_dgx_b200.yml +++ b/tests/integration/test_lists/test-db/l0_dgx_b200.yml @@ -194,6 +194,10 @@ l0_dgx_b200: - accuracy/test_disaggregated_serving.py::TestNemotron3Super120B::test_auto_dtype[mtp_nextn=3-block_reuse=True-use_py_transceiver=False] TIMEOUT (60) - accuracy/test_disaggregated_serving.py::TestNemotron3Super120B::test_auto_dtype[mtp_nextn=0-block_reuse=False-use_py_transceiver=False] TIMEOUT (60) - accuracy/test_disaggregated_serving.py::TestGLM52NVFP4::test_nvfp4_nixl[cache_mgr_v1] TIMEOUT (60) + # DeepSeek-V4 EPLB pre-merge sanity (uncomment once DeepSeek-V4-Flash/Flash-Base + # checkpoints are staged under llm_models_root()). + # - accuracy/test_llm_api_pytorch.py::TestDeepSeekV4Flash::test_nvfp4_4gpus_static_eplb[moe_backend=TRTLLM] TIMEOUT (120) + # - accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_fp8_4gpus_static_eplb[moe_backend=DEEPGEMM] TIMEOUT (120) - accuracy/test_disaggregated_serving.py::TestNemotron3Super120B::test_ctx_dp2_gen_tp4 TIMEOUT (60) - accuracy/test_disaggregated_serving.py::TestQwen3NextInstruct::test_auto_dtype[use_py_transceiver=True] TIMEOUT (60) - accuracy/test_disaggregated_serving.py::TestQwen3NextInstruct::test_auto_dtype[use_py_transceiver=False] TIMEOUT (60) From 07fff065605f21523ea9f46ac634078432899e06 Mon Sep 17 00:00:00 2001 From: xxi Date: Mon, 3 Aug 2026 06:12:02 +0000 Subject: [PATCH 3/6] [TRTLLM-14822][fix] keep Hopper coverage for FP8 block-scale static EPLB 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 --- .../defs/accuracy/test_llm_api_pytorch.py | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/tests/integration/defs/accuracy/test_llm_api_pytorch.py b/tests/integration/defs/accuracy/test_llm_api_pytorch.py index b8e4ff8d096b..ad77dec63cb7 100644 --- a/tests/integration/defs/accuracy/test_llm_api_pytorch.py +++ b/tests/integration/defs/accuracy/test_llm_api_pytorch.py @@ -2256,10 +2256,14 @@ def test_fp8_block_scales_4gpus_static_eplb(self): num_slots=num_slots, initial_global_assignments=initial_global_assignments, layer_updates_per_iter=0) - pytorch_backend_options = dict(cuda_graph_config=CudaGraphConfig(), - moe_config=MoeConfig( - backend="DEEPGEMM", - load_balancer=eplb_config)) + # Replaces the deprecated WIDEEP backend, which dispatched on + # is_sm_100f() internally: DeepGEMM covers SM100/SM103 FP8 block + # scales, CUTLASS covers SM90/SM120. + pytorch_backend_options = dict( + cuda_graph_config=CudaGraphConfig(), + moe_config=MoeConfig( + backend="DEEPGEMM" if is_sm_100f() else "CUTLASS", + load_balancer=eplb_config)) with LLM(f"{llm_models_root()}/DeepSeek-V3-Lite/fp8", tensor_parallel_size=4, moe_expert_parallel_size=4, From d6d38e27daf2c8fa70a37d4f5a64aa02d2717f18 Mon Sep 17 00:00:00 2001 From: xxi Date: Mon, 3 Aug 2026 06:40:32 +0000 Subject: [PATCH 4/6] [TRTLLM-14822][doc] complete WIDEEP deprecation in examples and MoE guide 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 --- examples/layer_wise_benchmarks/README.md | 20 +++++++++---------- examples/wide_ep/ep_load_balancer/README.md | 7 +++---- .../modules/fused_moe/MOE_DEVELOPER_GUIDE.md | 7 ++++--- 3 files changed, 17 insertions(+), 17 deletions(-) diff --git a/examples/layer_wise_benchmarks/README.md b/examples/layer_wise_benchmarks/README.md index 2b8070f955c1..77ff7d15b6f6 100644 --- a/examples/layer_wise_benchmarks/README.md +++ b/examples/layer_wise_benchmarks/README.md @@ -86,8 +86,8 @@ NP=4 ./mpi_launch.sh ./run.sh config_gen.yaml --batch-size 32 --seq-len-q 4 NP=4 ./mpi_launch.sh ./run.sh config_ctx.yaml --layer-indices 5,6,7,8 NP=4 ./mpi_launch.sh ./run.sh config_gen.yaml --layer-indices 5,6,7,8 -# Scale DEP=16 to 4 GPUs: reduces the number of experts; uses MNNVL A2A if applicable -NP=4 ./mpi_launch.sh ./run.sh config_gen.yaml --scaled-from 16 --moe-backend WIDEEP +# Scale DEP=16 to 4 GPUs: reduces the number of experts +NP=4 ./mpi_launch.sh ./run.sh config_gen.yaml --scaled-from 16 --moe-backend CUTEDSL # Scale TEP=16 to 4 GPUs: reduces the number of attention heads and experts NP=4 ./mpi_launch.sh ./run.sh config_gen.yaml --scaled-from 16 --no-enable-attention-dp @@ -101,8 +101,8 @@ NP=2 ./mpi_launch.sh ./run.sh config_ctx.yaml --model Qwen/Qwen3-Next-80B-A3B-In NP=2 ./mpi_launch.sh ./run.sh config_gen.yaml --model Qwen/Qwen3-Next-80B-A3B-Instruct --layer-indices 6,7 --no-enable-attention-dp --mamba-ssm-cache-dtype float16 --batch-size 512 # Run with DeepEP A2A -NP=4 ./mpi_launch.sh -x TRTLLM_FORCE_ALLTOALL_METHOD=DeepEP ./run.sh config_ctx.yaml --moe-backend WIDEEP -NP=4 ./mpi_launch.sh -x TRTLLM_FORCE_ALLTOALL_METHOD=DeepEP ./run.sh config_gen.yaml --moe-backend WIDEEP +NP=4 ./mpi_launch.sh -x TRTLLM_FORCE_COMM_METHOD=DEEPEP ./run.sh config_ctx.yaml --moe-backend CUTEDSL +NP=4 ./mpi_launch.sh -x TRTLLM_FORCE_COMM_METHOD=DEEPEP ./run.sh config_gen.yaml --moe-backend CUTEDSL # Run with imbalanced ranks: in addition to activating all experts, the specified ratio of tokens is sent to rank 0 # Note: if balance ratio is 0, the "activate all experts" behavior is not applied @@ -157,14 +157,14 @@ python3 scripts/build_wheel.py --cuda_architectures native --no-venv --skip_buil **Step 3:** Run benchmarks to generate profiles. Run the following command on the controller node, where `NODES` ≤ the number of allocated nodes: ```bash -# Run DeepSeek-R1 NVFP4 with wide EP; uses MNNVL A2A if applicable -NODES=4 NP=16 ./slurm_launch.sh ./run.sh config_gen.yaml --moe-backend WIDEEP +# Run DeepSeek-R1 NVFP4 with wide EP +NODES=4 NP=16 ./slurm_launch.sh ./run.sh config_gen.yaml --moe-backend CUTEDSL # Run with TRTLLMGen NODES=4 NP=16 ./slurm_launch.sh ./run.sh config_gen.yaml --moe-backend TRTLLM # Run with DeepEPLowLatency -NODES=4 NP=16 TRTLLM_FORCE_ALLTOALL_METHOD=DeepEPLowLatency ./slurm_launch.sh ./run.sh config_gen.yaml --moe-backend WIDEEP +NODES=4 NP=16 TRTLLM_FORCE_COMM_METHOD=DEEPEPLOWLATENCY ./slurm_launch.sh ./run.sh config_gen.yaml --moe-backend CUTEDSL # You can run 4-GPU and 8-GPU tasks without reallocating the Slurm job NODES=1 NP=4 ./slurm_launch.sh ./run.sh config_ctx.yaml @@ -187,7 +187,7 @@ Run with OpenMPI: ```bash NP=4 ./mpi_launch.sh ./run.sh config_ctx.yaml --batch-size 1,2,4 --seq-len-q 1024,8192 -NP=4 ./mpi_launch.sh ./run.sh config_gen.yaml --scaled-from 16 --moe-backend WIDEEP --batch-size 32,64,128,256,512 --seq-len-q 1,2,3,4 +NP=4 ./mpi_launch.sh ./run.sh config_gen.yaml --scaled-from 16 --moe-backend CUTEDSL --batch-size 32,64,128,256,512 --seq-len-q 1,2,3,4 ``` ## Parse profiles @@ -354,7 +354,7 @@ Two E2E traces are required because the two pieces of information cannot be capt Limitations: 1. Pipeline parallelism is not supported. -2. Only the CUTLASS and WIDEEP MoE backends are supported. +2. Only the CUTEDSL, CUTLASS, DEEPGEMM and TRTLLM MoE backends are supported. 3. Only tested with the GEN phase and attention DP. ## Developer utilities @@ -372,7 +372,7 @@ Limitations: 1. Error `fp8 blockscale gemm only support Hopper` on Blackwell. - The default MoE backend "CUTLASS" does not support FP8 weights. Please choose the same MoE backend as your end-to-end config. A typical solution is to add the `--moe-backend DEEPGEMM` (or `TRTLLM`, `WIDEEP`) and `--moe-backend-for-prefill DEEPGEMM` (or `WIDEEP`) options. + The default MoE backend "CUTLASS" does not support FP8 weights. Please choose the same MoE backend as your end-to-end config. A typical solution is to add the `--moe-backend DEEPGEMM` (or `TRTLLM`, `CUTEDSL`) and `--moe-backend-for-prefill DEEPGEMM` options. 2. Error `huggingface_hub.errors.HfHubHTTPError: 429 Client Error: Too Many Requests for url: https://huggingface.co/nvidia/DeepSeek-R1-0528-FP4-v2/resolve/main/config.json`. diff --git a/examples/wide_ep/ep_load_balancer/README.md b/examples/wide_ep/ep_load_balancer/README.md index 7c4c77b432be..c250e2565280 100644 --- a/examples/wide_ep/ep_load_balancer/README.md +++ b/examples/wide_ep/ep_load_balancer/README.md @@ -27,9 +27,8 @@ Run 32-way expert parallelism inference on the prepared dataset. Please refer to cat > ./config.yaml < ./config_eplb.yaml < ./config_eplb.yaml < Date: Mon, 3 Aug 2026 23:48:28 +0000 Subject: [PATCH 5/6] [https://nvbugs/6546609][fix] register DeepSeek-V4-Flash-Base GSM8K reference Fold in the fix from the closed PR #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 --- .../integration/defs/accuracy/references/gsm8k.yaml | 12 ++++++++++++ tests/integration/test_lists/waives.txt | 4 ---- 2 files changed, 12 insertions(+), 4 deletions(-) diff --git a/tests/integration/defs/accuracy/references/gsm8k.yaml b/tests/integration/defs/accuracy/references/gsm8k.yaml index ef32f74f1ba0..8a020f2e6ebc 100644 --- a/tests/integration/defs/accuracy/references/gsm8k.yaml +++ b/tests/integration/defs/accuracy/references/gsm8k.yaml @@ -145,6 +145,18 @@ deepseek-ai/DeepSeek-V4-Flash: # 95.11 reference still holds for the hypothesis test. - quant_algo: FP8_BLOCK_SCALES accuracy: 95.11 +deepseek-ai/DeepSeek-V4-Flash-Base: + # Base (pretrained, non-instruct) checkpoint, so GSM8K lands well below the + # instruct DeepSeek-V4-Flash above. GSM8K measurements from + # test_fp8_4gpus_static_eplb at TP=4/EP=4 over the full 1319 samples: + # * 90.90 / 91.02 / 91.43 on 4x B200 183GB. + # * 91.13 on 4x B300 (SM103). + # SM100 and SM103 route FP8 block scales through the same DeepGEMM path, so + # the spread above is run-to-run noise rather than a per-GPU difference. The + # reference records the low end; the hypothesis test's ~3.2-point margin + # absorbs the drift. + - quant_algo: FP8_BLOCK_SCALES + accuracy: 90.90 deepseek-ai/DeepSeek-V4-Pro: # Full GSM8K aggregate gate for the Pro deployment path: TP=8, EP=8, # attention DP, TRTLLM MoE, FP8 KV cache, MTP max_draft_len=1, padded CUDA diff --git a/tests/integration/test_lists/waives.txt b/tests/integration/test_lists/waives.txt index 5a4df968870a..f058b87e6a08 100644 --- a/tests/integration/test_lists/waives.txt +++ b/tests/integration/test_lists/waives.txt @@ -148,7 +148,6 @@ full:B200/accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mt full:B200/accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus[tp4-mtp_nextn=0-attention_dp=False-cuda_graph=False-overlap_scheduler=False-torch_compile=False] SKIP (https://nvbugs/6525007) full:B200/accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_cute_dsl_nvfp4_4gpus[tp4-fp8kv=False-attention_dp=False-cuda_graph=False-overlap_scheduler=False-torch_compile=False] SKIP (https://nvbugs/6526186) full:B200/accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_cute_dsl_nvfp4_4gpus[tp4-fp8kv=False-attention_dp=False-cuda_graph=False-overlap_scheduler=False-torch_compile=True] SKIP (https://nvbugs/6474888) -full:B200/accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_fp8_4gpus_static_eplb[moe_backend=WIDEEP] SKIP (https://nvbugs/6546609) full:B200/accuracy/test_llm_api_pytorch.py::TestLagunaXS::test_fp8 SKIP (https://nvbugs/6525011) full:B200/accuracy/test_llm_api_pytorch.py::TestLlama3_1_8BInstruct::test_fp8[fp8kv=True-attn_backend=TRTLLM-torch_compile=True] SKIP (https://nvbugs/6473161) full:B200/accuracy/test_llm_api_pytorch.py::TestLlama3_1_8BInstruct::test_fp8_4gpus[tp4-fp8kv=True-attn_backend=TRTLLM-torch_compile=True] SKIP (https://nvbugs/6473161) @@ -181,7 +180,6 @@ full:B300/accuracy/test_llm_api_pytorch.py::TestDeepSeekV32::test_nvfp4_multi_gp full:B300/accuracy/test_llm_api_pytorch.py::TestDeepSeekV32::test_nvfp4_multi_gpus_chunked_prefill[latency_qsplit] SKIP (https://nvbugs/6423866) full:B300/accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=True-enable_chunked_prefill=True-v2_kv_cache=True] SKIP (https://nvbugs/6422343) full:B300/accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_nvfp4_4gpus[moe_backend=CUTLASS-mtp_nextn=0-ep4-fp8kv=True-attention_dp=True-cuda_graph=True-overlap_scheduler=True-low_precision_combine=False-torch_compile=True] SKIP (https://nvbugs/6474888) -full:B300/accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_fp8_4gpus_static_eplb[moe_backend=WIDEEP] SKIP (https://nvbugs/6546609) full:B300/accuracy/test_llm_api_pytorch.py::TestGemma3_1BInstruct::test_fp8_prequantized[torch_compile=True] SKIP (https://nvbugs/6475346) full:B300/accuracy/test_llm_api_pytorch.py::TestLagunaXS::test_fp8 SKIP (https://nvbugs/6525011) full:B300/accuracy/test_llm_api_pytorch.py::TestMiniMaxM3::test_auto_dtype[tp_size=8-ep_size=8] SKIP (https://nvbugs/6445375) @@ -206,7 +204,6 @@ full:GB200/accuracy/test_llm_api_autodeploy.py::TestNemotronNanoV3::test_accurac full:GB200/accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mtp_nextn=0-attention_dp=False-cuda_graph=False-overlap_scheduler=False-torch_compile=False-enable_chunked_prefill=False-v2_kv_cache=False] SKIP (https://nvbugs/6525896) full:GB200/accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_cute_dsl_bf16_gemm[cuda_graph=True] SKIP (https://nvbugs/6525897) full:GB200/accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_fp8_block_scales_4gpus[ep4-mtp_nextn=0-fp8kv=False-attention_dp=False-cuda_graph=True-overlap_scheduler=False-torch_compile=False-sampler_async_worker=False] SKIP (https://nvbugs/6547150) -full:GB200/accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_fp8_4gpus_static_eplb[moe_backend=WIDEEP] SKIP (https://nvbugs/6546609) full:GB200/accuracy/test_llm_api_pytorch.py::TestMiniMaxM3::test_mxfp8[use_msa=False] SKIP (https://nvbugs/6479471) full:GB200/accuracy/test_llm_api_pytorch.py::TestMiniMaxM3::test_nvfp4[use_msa=False] SKIP (https://nvbugs/6479471) full:GB200/accuracy/test_llm_api_pytorch.py::TestMinistral8BInstruct::test_auto_dtype SKIP (https://nvbugs/6547151) @@ -231,7 +228,6 @@ full:GB300/accuracy/test_llm_api_autodeploy.py::TestNemotronNanoV3::test_accurac full:GB300/accuracy/test_llm_api_autodeploy.py::TestNemotronNanoV3::test_accuracy[nvfp4-1-attn_dp_off-trtllm] SKIP (https://nvbugs/6329165) full:GB300/accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16[mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=True-enable_chunked_prefill=True-v2_kv_cache=True] SKIP (https://nvbugs/6422343) full:GB300/accuracy/test_llm_api_pytorch.py::TestDeepSeekV3Lite::test_bfloat16_4gpus[pp4-mtp_nextn=2-attention_dp=True-cuda_graph=True-overlap_scheduler=True-torch_compile=False] SKIP (https://nvbugs/6525057) -full:GB300/accuracy/test_llm_api_pytorch.py::TestDeepSeekV4FlashBase::test_fp8_4gpus_static_eplb[moe_backend=WIDEEP] SKIP (https://nvbugs/6546609) full:GB300/accuracy/test_llm_api_pytorch.py::TestGPTOSS::test_w4_4gpus[v1_kv_cache-dp4-trtllm-fp8] SKIP (https://nvbugs/6474894) full:GB300/accuracy/test_llm_api_pytorch.py::TestGemma3_1BInstruct::test_fp8_prequantized[torch_compile=True] SKIP (https://nvbugs/6475346) full:GB300/accuracy/test_llm_api_pytorch.py::TestLlama3_1_8BInstruct::test_bfloat16_4gpus[pp4-attn_backend=FLASHINFER-torch_compile=False] SKIP (https://nvbugs/6385771) From 2224fdf56bab969fca32075422541a2775877463 Mon Sep 17 00:00:00 2001 From: xxi Date: Thu, 6 Aug 2026 10:50:42 +0000 Subject: [PATCH 6/6] [TRTLLM-14822][fix] drop WIDEEP from the Qwen3-Next MoE quant probe test_qwen3_next_moe_quant.py landed on main (#17120) after this branch was cut, and parametrizes its two backend-resolution probes over WIDEEP. test_unexcluded_layer_keeps_configured_backend_and_layer_quant_config reaches resolve_moe_cls() -> get_moe_cls() with moe_backend=WIDEEP, which this branch turns into a ValueError, so that parametrization now fails. test_excluded_layer_builds_bf16_on_cutlass only passes by accident: Qwen3NextSparseMoeBlock rewrites moe_backend to CUTLASS before the call. Remove the WIDEEP entries from both parametrize lists, matching how the rest of this branch handles WIDEEP in unit tests. The other backends keep full coverage of both probes. Signed-off-by: xxi --- tests/unittest/_torch/models/test_qwen3_next_moe_quant.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tests/unittest/_torch/models/test_qwen3_next_moe_quant.py b/tests/unittest/_torch/models/test_qwen3_next_moe_quant.py index d28b6403291c..7b2370e60503 100644 --- a/tests/unittest/_torch/models/test_qwen3_next_moe_quant.py +++ b/tests/unittest/_torch/models/test_qwen3_next_moe_quant.py @@ -221,7 +221,7 @@ def _capture(*args, **kwargs): return captured -@pytest.mark.parametrize("backend", ["CUTLASS", "TRTLLM", "DEEPGEMM", "WIDEEP", "CUTEDSL"]) +@pytest.mark.parametrize("backend", ["CUTLASS", "TRTLLM", "DEEPGEMM", "CUTEDSL"]) @pytest.mark.parametrize("layer_idx", [5, MTP_LAYER_IDX]) def test_excluded_layer_builds_bf16_on_cutlass(backend, layer_idx): per_layer_quant_config = QuantConfig(quant_algo=QuantAlgo.FP8_BLOCK_SCALES) @@ -245,7 +245,6 @@ def test_excluded_layer_builds_bf16_on_cutlass(backend, layer_idx): "CUTLASS": "CutlassFusedMoE", "TRTLLM": "TRTLLMGenFusedMoE", "DEEPGEMM": "DeepGemmFusedMoE", - "WIDEEP": "WideEPMoE", "CUTEDSL": "CuteDslFusedMoE", }