doc: Update doc for Deepseek min latency (NVIDIA#3717)

* Tidy code

Signed-off-by: Zongfei Jing <[email protected]>

* Update doc for min latency deepseek

Signed-off-by: Zongfei Jing <[email protected]>

* Throw exception for RouterKernel when not running on sm90+

Signed-off-by: Zongfei Jing <[email protected]>

---------

Signed-off-by: Zongfei Jing <[email protected]>

Author: zongfeijing

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/trtllmGenSrc/RoutingKernel.cu

@@ -698,6 +698,7 @@ __global__ void __cluster_dims__(NumBlocksPerCluster, 1, 1) __launch_bounds__(Nu
 #else
 __global__ void routingIndicesClusterKernel(KernelParams params)
 {
+    assert(false && "routingIndicesClusterKernel is only supported on SM90+ architectures");
 }
 #endif
 ////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -886,6 +887,8 @@ __global__ void __launch_bounds__(NumThreads) routingIndicesCoopKernel(KernelPar
             params.mPtrPermutedIdxToTokenIdx[permutedIdx] = tokenIdx;
         }
     }
+#else
+    assert(false && "routingIndicesCoopKernel is only supported on SM90+ architectures");
 #endif
 }
 
@@ -973,6 +976,8 @@ __global__ void __launch_bounds__(NumThreads) routingIndicesHistogramKernel(Kern
     // Reduce histograms with atomics.
     int32_t const localExpertCount = smemExpertCount[threadIdx.x];
     atomicAdd(&params.mPtrExpertCounts[threadIdx.x], localExpertCount);
+#else
+    assert(false && "routingIndicesHistogramKernel is only supported on SM90+ architectures");
 #endif
 }
 
@@ -1204,6 +1209,8 @@ __global__ void __launch_bounds__(NumThreads) routingIndicesOffsetsKernel(Kernel
     {
         cudaTriggerProgrammaticLaunchCompletion();
     }
+#else
+    assert(false && "routingIndicesOffsetsKernel is only supported on SM90+ architectures");
 #endif
 }
 

docs/source/blogs/Best_perf_practice_on_DeepSeek-R1_in_TensorRT-LLM.md

@@ -110,6 +110,7 @@ cat >./extra-llm-api-config.yml<<EOF
 pytorch_backend_config:
     enable_overlap_scheduler: true
     use_cuda_graph: true
+    moe_backend: TRTLLM
 speculative_config:
     decoding_type: MTP
     num_nextn_predict_layers: 3
@@ -125,7 +126,7 @@ trtllm-bench --model nvidia/DeepSeek-R1-FP4 \
     --concurrency 1 \
     --max_batch_size 1 \
     --tp 8 \
-    --ep 4 \
+    --ep 2 \
     --extra_llm_api_options ./extra-llm-api-config.yml
 ```
 
@@ -147,12 +148,13 @@ The perf can be different when using different datasets and different machines.
 ===========================================================
 = PERFORMANCE OVERVIEW
 ===========================================================
-Request Throughput (req/sec):                     0.1244
-Total Output Throughput (tokens/sec):             254.5535
-Per User Output Throughput (tokens/sec/user):     254.7634
-Per GPU Output Throughput (tokens/sec/gpu):       31.8192
-Total Latency (ms):                               80368.1616
-Average request latency (ms):                     8036.7546
+Request Throughput (req/sec):                     0.1341
+Total Output Throughput (tokens/sec):             274.4168
+Per User Output Throughput (tokens/sec/user):     274.7188
+Per GPU Output Throughput (tokens/sec/gpu):       34.3021
+Total Token Throughput (tokens/sec):              414.0461
+Total Latency (ms):                               74561.7520
+Average request latency (ms):                     7456.1219
 ```
 
 ### B200 max-throughput

tensorrt_llm/_torch/models/modeling_deepseekv3.py

@@ -624,6 +624,9 @@ def _compute_mlp_tp_size(self, intermediate_size: int,
             )
         return mlp_tp_size
 
+    def _enable_latency_mode(self, num_tokens: int):
+        return num_tokens <= 128 and self.fusion_config.POST_MOE_FUSION and self.is_nvfp4 and self.model_config.moe_backend == 'CUTLASS'
+
     def forward(
         self,
         position_ids: torch.LongTensor,
@@ -650,9 +653,7 @@ def forward(
         using_prev_fusion = self.deepseek_allreduce_disabled or hidden_states.size(
             0) > 128
 
-        min_latency_mode = True if hidden_states.size(
-            0
-        ) <= 128 and self.fusion_config.POST_MOE_FUSION and self.is_nvfp4 and self.model_config.moe_backend == 'CUTLASS' else False
+        min_latency_mode = self._enable_latency_mode(hidden_states.size(0))
 
         if self.fusion_config.PRE_MOE_FUSION:
             # Custom AR Fusion for DeepseekV3