cpp/tensorrt_llm/thop/fp8BlockScaleMoe.cpp (1)

2-2: Update the copyright year to include 2026.

The header currently says 2022-2024, but this file is being modified now. As per coding guidelines: "update year on modified files."

- * Copyright (c) 2022-2024, NVIDIA CORPORATION.  All rights reserved.
+ * Copyright (c) 2022-2026, NVIDIA CORPORATION.  All rights reserved.

tests/unittest/_torch/thop/serial/test_moe.py (2)

1095-1105: Large memory footprint — consider CI impact.

With num_experts=2048, hidden_size=1024, and intermediate_size=1024, gemm1_weights alone is (2048, 2048, 1024) in bf16 (~8 GB), plus gemm2_weights (~4 GB), plus all quantized copies and scales. Combined with the parametrize cross-product (num_tokens × intermediate_size × act_type), this generates multiple heavy test cases.

Consider either:

• Reducing hidden_size/intermediate_size for this specific large-expert case, or
• Limiting the cross-product (e.g., fixing intermediate_size=768 and a single act_type for this param).

---

1189-1199: Same memory concern as test_autotune applies here.

This test_no_autotune variant additionally cross-products with use_topk_as_input=[False, True], further doubling the large-expert test matrix. The use_topk_as_input=True path hits the DeepSeekV3-only skip at line 1357, but the False path still runs all combinations.

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu (1)

29-30: Commented-out code should use #if/#endif instead of comments.

Line 29 uses a comment to disable code (// static constexpr int MaxNumExperts = 128;). If this is dead code, remove it entirely. If it's intentionally kept for reference, use #if 0 / #endif per the coding guidelines.

As per coding guidelines: "Use #if / #endif to disable code, preferably with a mnemonic condition... Do not use comments to disable code."

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.h (1)

37-70: Good refactor to do { ... } while (0) for the LAUNCH_PDL macro.

This is a well-known best practice for multi-statement macros. One minor coding guideline nit: the if on lines 56 and 65 should be followed by brace-delimited statements.

As per coding guidelines: "If and else should always be followed by brace-delimited statements, even if empty or a single statement."

Proposed fix (lines 56-58, 65-67)

             if (smemSize > 48 * 1024)                                                                  \
-                TLLM_CUDA_CHECK(                                                                       \
-                    cudaFuncSetAttribute(kernelTyped, cudaFuncAttributeMaxDynamicSharedMemorySize, smemSize));         \
+            {                                                                                          \
+                TLLM_CUDA_CHECK(                                                                       \
+                    cudaFuncSetAttribute(kernelTyped, cudaFuncAttributeMaxDynamicSharedMemorySize, smemSize));         \
+            }                                                                                          \

Apply similarly to both the true and false UsePdl branches.

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/routingDeepSeek/RoutingDeepSeekCommon.cuh (1)

26-29: std::max with initializer list may need explicit <algorithm> include.

Line 29 uses std::max({...}) with an initializer list, which requires <algorithm> (and implicitly <initializer_list>). These may be provided transitively via RoutingKernel.cuh → DevKernel.h, but relying on transitive includes is fragile. Consider a direct #include <algorithm>, or alternatively, since all three values are compile-time constants, a simpler expression like nested std::max calls or just hardcoding 512 (the actual max) would avoid the dependency entirely.

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/routingDeepSeek/launchInitExpertCounts.cu (1)

27-30: Missing /*coopLaunch=*/ inline comment for consistency.

All other launcher wrappers annotate the coopLaunch parameter with an inline comment (/*coopLaunch=*/false). This one passes bare false.

Suggested fix

-    LAUNCH_ROUTING_DEEPSEEK(data, false, routingInitExpertCounts, (2 * data.mNumExperts - 1) / numThreadsHist + 1,
+    LAUNCH_ROUTING_DEEPSEEK(data, /*coopLaunch=*/false, routingInitExpertCounts, (2 * data.mNumExperts - 1) / numThreadsHist + 1,

As per coding guidelines: "In function calls where parameters are not obvious, use inline C comments to document the parameter."

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/routingDeepSeek/launchCoopKernel.cu (1)

216-222: TODO note: PDL visibility concern is documented but unresolved.

The comment on line 218 explicitly states: "this is not sufficient to ensure visibility in the next kernel!" This suggests the secondary kernel may observe stale data for mPtrCtaIdxXyToBatchIdx, mPtrCtaIdxXyToMnLimit, mPtrNumNonExitingCtas, and mPtrPermutedIdxSize.

Is there a tracking issue for this? If the dependent FC1 kernel relies on these outputs being visible, a missing memory fence before cudaTriggerProgrammaticLaunchCompletion could cause data races on SM90+.

Would you like me to open an issue to track this PDL visibility concern?

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/routingDeepSeek/launchClusterKernel.cu (1)

44-49: Non-SM90 fallback is missing __launch_bounds__ unlike the SM90+ path.

The SM90+ variant (line 27) uses __launch_bounds__(KernelParams::MaxNumExperts) while the fallback (line 45) has no launch bounds annotation. For consistency with the Renormalize variant (which uses __launch_bounds__(NumThreads) on its fallback), consider adding it here. This is minor since the fallback only asserts false.

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/routingRenormalize/launchHistogramScoresKernel.cu (1)

44-44: Unused variable minScore.

minScore is initialized to -INFINITY but never read or passed to any function in this kernel. Likely a copy-paste artifact from the block kernel. Remove it to avoid compiler warnings and dead code.

Suggested fix

-    BaseType minScore = BaseType{-INFINITY};

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/routingRenormalize/launchBlockKernel.cu (1)

100-100: Unused variable minScore.

Same as in launchHistogramScoresKernel.cu — minScore is initialized but never referenced. Remove to avoid dead code.

Suggested fix

-        BaseType minScore = BaseType{-INFINITY};