📄 CodeRabbit inference engine (CODING_GUIDELINES.md)

**/*.{h,hpp,hh,hxx,cpp,cxx,cc,cu,cuh}: Namespace closing braces must include a trailing comment with the namespace name (e.g., '} // namespace foo').
Prefer const or constexpr variables over #define for constants.
Declare variables that are not modified after initialization as const.
Avoid magic literals in code; except for 0, nullptr, true, false. Use named constants for comparisons and logic.
Use Allman brace style for formatting.
Place the semicolon of an empty for/while loop on a new line.
Bodies of switch/while/do-while/for must be compound statements (brace-delimited), and if/else must always be followed by brace-delimited statements.
Type names (e.g., classes) must be CamelCase starting with an uppercase letter (e.g., FooBar).
Local variables, methods, and namespaces use lowerCamelCase (e.g., localFooBar).
Non-magic-number global variables that are non-static and not in an anonymous namespace must be lowerCamelCase prefixed with 'g' (e.g., gDontUseGlobalFoos).
Non-magic-number globals that are static or in an anonymous namespace use lowerCamelCase prefixed with 's' (e.g., sMutableStaticGlobal).
Locally visible static variables use lowerCamelCase with 's' prefix (e.g., static std::once_flag sFlag).
Private/protected member variables use 'm' prefix with CamelCase (e.g., mNbFooValues). Public members may omit, but 'm' is encouraged for clarity.
Constants (enums, global constants, static constants, and function-scope magic/literal constants) use uppercase SNAKE_CASE with 'k' prefix (e.g., kDIGIT_NUM).
Function-scope constants that are not magic numbers or literals are named like non-constant variables (e.g., bool const pass = a && b).
If macros are necessary, name them in UPPER_SNAKE_CASE (e.g., FOO_VERSION) and prefer constants over #define.
Use LLVM clang-format; wrap lines at a maximum of 120 columns; use '// clang-format off/on' sparingly with justification.
Use smart pointers for heap allocations; prefer unique_ptr for sole ownership, shared_ptr for shared...

Files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.h
• cpp/tensorrt_llm/kernels/moePrepareKernels.h
• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.cu
• cpp/tensorrt_llm/common/envUtils.h
• cpp/tensorrt_llm/kernels/moeLoadBalance/moeLoadBalanceKernels.cu
• cpp/tensorrt_llm/kernels/moePrepareKernels.cu

📄 CodeRabbit inference engine (CODING_GUIDELINES.md)

C++ filenames should be lowerCamelCase (first letter lowercase) and must be case-insensitive unique within a compilation target.

Files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.h
• cpp/tensorrt_llm/kernels/moePrepareKernels.h
• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.cu
• cpp/tensorrt_llm/common/envUtils.h
• cpp/tensorrt_llm/kernels/moeLoadBalance/moeLoadBalanceKernels.cu
• cpp/tensorrt_llm/kernels/moePrepareKernels.cu

📄 CodeRabbit inference engine (CODING_GUIDELINES.md)

Use only spaces, no tabs; indent with 4 spaces.

Files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.h
• cpp/tensorrt_llm/kernels/moePrepareKernels.h
• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.cu
• cpp/tensorrt_llm/common/envUtils.h
• cpp/tensorrt_llm/kernels/moeLoadBalance/moeLoadBalanceKernels.cu
• cpp/tensorrt_llm/kernels/moePrepareKernels.cu

📄 CodeRabbit inference engine (CODING_GUIDELINES.md)

Document new class interfaces and function prototypes with Doxygen; use //! for single-line and //!< for members.

Files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.h
• cpp/tensorrt_llm/kernels/moePrepareKernels.h
• cpp/tensorrt_llm/common/envUtils.h

📄 CodeRabbit inference engine (CODING_GUIDELINES.md)

**/*.{h,hpp,hh,hxx,cpp,cxx,cc}: Prefer anonymous namespaces over 'static' for internal linkage of functions.
All templates (class/function/member/static) must be instantiated at least once; non-POD classes should have private data members.

Files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.h
• cpp/tensorrt_llm/kernels/moePrepareKernels.h
• cpp/tensorrt_llm/common/envUtils.h

📄 CodeRabbit inference engine (CODING_GUIDELINES.md)

Use include guards named 'TRTLLM_<FILE_NAME_IN_CAPS_WITH_UNDERSCORES>_H' (no leading or trailing underscore; directory names excluded).

Files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.h
• cpp/tensorrt_llm/kernels/moePrepareKernels.h
• cpp/tensorrt_llm/common/envUtils.h

📄 CodeRabbit inference engine (CODING_GUIDELINES.md)

Prepend the NVIDIA Apache-2.0 copyright header with current year to the top of all source files (e.g., .cpp, .h, .cu, .py).

Files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.h
• cpp/tensorrt_llm/kernels/moePrepareKernels.h
• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.cu
• cpp/tensorrt_llm/common/envUtils.h
• cpp/tensorrt_llm/kernels/moeLoadBalance/moeLoadBalanceKernels.cu
• cpp/tensorrt_llm/kernels/moePrepareKernels.cu

Learnt from: nv-lschneider
PR: NVIDIA/TensorRT-LLM#7910
File: cpp/tensorrt_llm/kernels/nccl_device/config.cu:15-17
Timestamp: 2025-09-23T15:01:00.070Z
Learning: In TensorRT-LLM NCCL device kernels, the <sstream> header is not needed as an explicit include in config.cu because it's provided transitively through other headers. Local compilation testing confirms this works without the explicit include.

Applied to files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.h
• cpp/tensorrt_llm/kernels/moePrepareKernels.h
• cpp/tensorrt_llm/kernels/moeLoadBalance/moeLoadBalanceKernels.cu

Learnt from: nv-lschneider
PR: NVIDIA/TensorRT-LLM#7910
File: cpp/tensorrt_llm/kernels/nccl_device/config.cu:15-17
Timestamp: 2025-09-23T15:01:00.070Z
Learning: In TensorRT-LLM NCCL device kernels (cpp/tensorrt_llm/kernels/nccl_device/config.cu), std::ostringstream is used but <sstream> doesn't need to be explicitly included because it's provided transitively through other headers like tensorrt_llm/common/cudaUtils.h or config.h. Local compilation testing confirms this works without the explicit include.

Applied to files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.h
• cpp/tensorrt_llm/kernels/moePrepareKernels.h
• cpp/tensorrt_llm/kernels/moeLoadBalance/moeLoadBalanceKernels.cu

Learnt from: nv-lschneider
PR: NVIDIA/TensorRT-LLM#7910
File: cpp/tensorrt_llm/kernels/nccl_device/multimem.h:20-30
Timestamp: 2025-09-23T15:13:48.819Z
Learning: TRT-LLM targets modern CUDA toolkits that support FP8 datatypes, so cuda_fp8.h can be included unconditionally without version guards in TRT-LLM code.

Applied to files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.h
• cpp/tensorrt_llm/kernels/moePrepareKernels.h
• cpp/tensorrt_llm/kernels/moeLoadBalance/moeLoadBalanceKernels.cu

Learnt from: djns99
PR: NVIDIA/TensorRT-LLM#7104
File: cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu:1475-1480
Timestamp: 2025-08-21T02:39:12.009Z
Learning: The min latency mode functionality in TensorRT-LLM MOE kernels (cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu) is deprecated and no longer being maintained/updated, as confirmed by djns99. Bug reports and optimization suggestions for the computeStridesTmaWarpSpecializedLowLatencyKernel and related min latency code paths should be deprioritized.

Applied to files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.cu

Learnt from: djns99
PR: NVIDIA/TensorRT-LLM#6915
File: cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu:4616-4626
Timestamp: 2025-08-19T03:35:20.866Z
Learning: In the MOE profiler TMA workspace preparation (cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu), the overlapping of TMA WS regions for NONE and FINALIZE variants is deliberate design to save memory space, as confirmed by djns99. The comment "reuse the same pointers to save space" reflects this intentional behavior.

Applied to files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.cu

Learnt from: sklevtsov-nvidia
PR: NVIDIA/TensorRT-LLM#3294
File: cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu:1198-1209
Timestamp: 2025-08-08T22:03:40.707Z
Learning: In the CUTLASS MoE kernels (cpp/tensorrt_llm/cutlass_extensions), when `layout_info.fusion` is set to `TmaWarpSpecializedGroupedGemmInput::EpilogueFusion::FINALIZE`, the `router_scales` parameter must be non-null by design. The fused finalize kernel epilogue does not perform nullptr checks and requires valid router scales to function correctly. This is an implicit contract that callers must satisfy when enabling the FINALIZE fusion mode.

Applied to files:

• cpp/tensorrt_llm/kernels/fusedMoeCommKernels.cu

Learnt from: ChristinaZ
PR: NVIDIA/TensorRT-LLM#7068
File: cpp/tensorrt_llm/kernels/moeTopKFuncs.cuh:169-172
Timestamp: 2025-08-20T07:43:36.447Z
Learning: In TensorRT-LLM MOE kernels, when processing up to 128 experts across 32 threads, each thread handles at most 4 experts (N < 5 constraint), where N represents candidates per thread rather than total system capacity.

Applied to files:

• cpp/tensorrt_llm/kernels/moeLoadBalance/moeLoadBalanceKernels.cu
• cpp/tensorrt_llm/kernels/moePrepareKernels.cu

Learnt from: djns99
PR: NVIDIA/TensorRT-LLM#7104
File: cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu:1185-1189
Timestamp: 2025-08-25T00:03:39.294Z
Learning: TLLM_CHECK_WITH_INFO is a host-side utility function and cannot be called from CUDA device functions (those marked with __device__ or __global__). In device code, assert() is the primary mechanism for handling "should never happen" conditions, and like standard C++ assert, CUDA's assert only works in debug builds and is compiled out in release builds.

Applied to files:

• cpp/tensorrt_llm/kernels/moeLoadBalance/moeLoadBalanceKernels.cu

Learnt from: nv-lschneider
PR: NVIDIA/TensorRT-LLM#7910
File: cpp/tensorrt_llm/kernels/nccl_device/config.cu:42-49
Timestamp: 2025-09-23T14:58:05.372Z
Learning: In TensorRT-LLM NCCL device kernels (cpp/tensorrt_llm/kernels/nccl_device/), the token partitioning intentionally uses ceil-like distribution (same token_per_rank for all ranks) to ensure all ranks launch the same number of blocks. This is required for optimal NCCL device API barrier performance, even though it may launch extra blocks for non-existent tokens on later ranks. Runtime bounds checking in the kernel (blockID validation) handles the overshoot cases.

Applied to files:

• cpp/tensorrt_llm/kernels/moeLoadBalance/moeLoadBalanceKernels.cu
• cpp/tensorrt_llm/kernels/moePrepareKernels.cu

cpp/tensorrt_llm/kernels/moePrepareKernels.cu (5)

• grid (302-302)
• grid (360-360)
• block (301-301)
• block (335-335)
• block (359-359)

cpp/tensorrt_llm/kernels/decoderMaskedMultiheadAttention/decoderMaskedMultiheadAttentionLaunch.h (1)

• grid (259-345)

cpp/tensorrt_llm/kernels/fusedMoeCommKernels.cu (4)

• tokenCount (1066-1128)
• tokenCount (1066-1066)
• tokenCount (1140-1213)
• tokenCount (1140-1140)

cpp/tensorrt_llm/kernels/fusedMoeCommKernels.cu (4)

• tokenCount (1066-1128)
• tokenCount (1066-1066)
• tokenCount (1140-1213)
• tokenCount (1140-1140)