📄 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/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/thop/fp4Quantize.cpp
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp

📄 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/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/thop/fp4Quantize.cpp
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp

📄 CodeRabbit inference engine (CODING_GUIDELINES.md)

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

Files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• tensorrt_llm/_torch/autotuner.py
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• tensorrt_llm/_torch/cute_dsl_kernels/blackwell/grouped_blockscaled_gemm_persistent.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• tests/integration/defs/accuracy/test_llm_api_pytorch.py
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/thop/fp4Quantize.cpp
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• tests/unittest/_torch/thop/parallel/test_cute_dsl_moe.py
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

📄 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/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.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/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• tensorrt_llm/_torch/autotuner.py
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• tensorrt_llm/_torch/cute_dsl_kernels/blackwell/grouped_blockscaled_gemm_persistent.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• tests/integration/defs/accuracy/test_llm_api_pytorch.py
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/thop/fp4Quantize.cpp
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• tests/unittest/_torch/thop/parallel/test_cute_dsl_moe.py
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

📄 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/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.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/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/thop/fp4Quantize.cpp
• cpp/tensorrt_llm/thop/moeUtilOp.cpp

📄 CodeRabbit inference engine (CODING_GUIDELINES.md)

**/*.py: Python code must target Python 3.8+.
Indent Python code with 4 spaces; do not use tabs.
Maintain module namespace when importing; prefer 'from package.subpackage import foo' then 'foo.SomeClass()' instead of importing the class directly.
Python filenames should be snake_case (e.g., some_file.py).
Python classes use PascalCase names.
Functions and methods use snake_case names.
Local variables use snake_case; prefix 'k' for variables that start with a number (e.g., k_99th_percentile).
Global variables use upper SNAKE_CASE prefixed with 'G' (e.g., G_MY_GLOBAL).
Constants use upper SNAKE_CASE (e.g., MY_CONSTANT).
Avoid shadowing variables from an outer scope.
Initialize all externally visible members of a class in the constructor.
Prefer docstrings for interfaces that may be used outside a file; comments for in-function or file-local interfaces.
Use Google-style docstrings for classes and functions (Sphinx-parsable).
Document attributes and variables inline so they render under the class/function docstring.
Avoid reflection when a simpler, explicit approach suffices (e.g., avoid dict(**locals()) patterns).
In try/except, catch the most specific exceptions possible.
For duck-typing try/except, keep the try body minimal and use else for the main logic.

Files:

• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• tensorrt_llm/_torch/autotuner.py
• tensorrt_llm/_torch/cute_dsl_kernels/blackwell/grouped_blockscaled_gemm_persistent.py
• tests/integration/defs/accuracy/test_llm_api_pytorch.py
• tests/unittest/_torch/thop/parallel/test_cute_dsl_moe.py
• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 7104
File: cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/epilogue/fusion/sm90_visitor_scatter.hpp:399-417
Timestamp: 2025-08-21T21:48:35.135Z
Learning: CUTLASS extensions in TensorRT-LLM (located under cpp/tensorrt_llm/cutlass_extensions/) are designed to integrate with and extend functionality in the external CUTLASS repository. When analyzing these extensions, their consumers and functionality wiring may exist in the CUTLASS codebase rather than within TensorRT-LLM itself.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/kernels/CMakeLists.txt
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

Learnt from: nv-lschneider
Repo: NVIDIA/TensorRT-LLM PR: 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/CMakeLists.txt
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/thop/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/cuteDslKernels/CMakeLists.txt
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu

Learnt from: tongyuantongyu
Repo: NVIDIA/TensorRT-LLM PR: 6984
File: cpp/tensorrt_llm/CMakeLists.txt:297-299
Timestamp: 2025-08-18T09:08:07.687Z
Learning: In the TensorRT-LLM project, artifacts are manually copied rather than installed via `cmake --install`, so INSTALL_RPATH properties are not needed - only BUILD_RPATH affects the final artifacts.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt

Learnt from: sklevtsov-nvidia
Repo: NVIDIA/TensorRT-LLM PR: 3294
File: cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/epilogue/fusion/sm90_visitor_scatter.hpp:36-36
Timestamp: 2025-08-08T05:06:31.596Z
Learning: CUTLASS extension files (under cpp/tensorrt_llm/cutlass_extensions/) follow CUTLASS coding style conventions, including using #pragma once instead of TRTLLM_ prefixed header guards, even though they are .hpp files.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/thop/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/kernels/cuteDslKernels/CMakeLists.txt
• cpp/tensorrt_llm/kernels/CMakeLists.txt

Learnt from: fredricz-20070104
Repo: NVIDIA/TensorRT-LLM PR: 7645
File: tests/integration/test_lists/qa/llm_function_core.txt:648-648
Timestamp: 2025-09-09T09:40:45.658Z
Learning: In TensorRT-LLM test lists, it's common and intentional for the same test to appear in multiple test list files when they serve different purposes (e.g., llm_function_core.txt for comprehensive core functionality testing and llm_function_core_sanity.txt for quick sanity checks). This duplication allows tests to be run in different testing contexts.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt
• tests/integration/test_lists/test-db/l0_b200.yml
• tests/integration/test_lists/test-db/l0_dgx_b200.yml
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

Learnt from: nv-lschneider
Repo: NVIDIA/TensorRT-LLM PR: 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/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/thop/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/CMakeLists.txt

Learnt from: tongyuantongyu
Repo: NVIDIA/TensorRT-LLM PR: 7763
File: cpp/tensorrt_llm/CMakeLists.txt:297-301
Timestamp: 2025-09-16T09:30:09.716Z
Learning: In the TensorRT-LLM project, NCCL libraries are loaded earlier by PyTorch libraries or the bindings library, so the main shared library doesn't need NCCL paths in its RPATH - the libraries will already be available in the process address space when needed.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt

Learnt from: galagam
Repo: NVIDIA/TensorRT-LLM PR: 6487
File: tests/unittest/_torch/auto_deploy/unit/singlegpu/test_ad_trtllm_bench.py:1-12
Timestamp: 2025-08-06T13:58:07.506Z
Learning: In TensorRT-LLM, test files (files under tests/ directories) do not require NVIDIA copyright headers, unlike production source code files. Test files typically start directly with imports, docstrings, or code.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt
• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• tests/integration/test_lists/test-db/l0_dgx_b200.yml

Learnt from: jhaotingc
Repo: NVIDIA/TensorRT-LLM PR: 7856
File: cpp/tensorrt_llm/thop/fp8BlockScaleMoe.cpp:159-166
Timestamp: 2025-09-19T21:28:13.751Z
Learning: In TensorRT-LLM blockScaleMoe routing (cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu), the DeepSeek routing method performs reinterpret_cast<float*>(routingLogits) at line 89, which could cause issues if routing_logits are BF16. However, Qwen3-FP8 models use RenormalizeNaive routing method and are not affected by this dtype casting issue.

Applied to files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp

Learnt from: sklevtsov-nvidia
Repo: NVIDIA/TensorRT-LLM PR: 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/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp

Learnt from: ChristinaZ
Repo: NVIDIA/TensorRT-LLM PR: 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/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp

Learnt from: nv-lschneider
Repo: NVIDIA/TensorRT-LLM PR: 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/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

Learnt from: thorjohnsen
Repo: NVIDIA/TensorRT-LLM PR: 6910
File: cpp/tensorrt_llm/batch_manager/kvCacheManager.cpp:0-0
Timestamp: 2025-08-14T21:04:50.248Z
Learning: In KV cache onboarding logic during prefill in cpp/tensorrt_llm/batch_manager/kvCacheManager.cpp, when calculating which blocks fall within the attention window, use getTokensPerBlock() to advance token indices rather than block->getUniqueTokens().size(), because the calculation needs to consider the post-prefill state where blocks will be filled to capacity, not their current token count.

Applied to files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu

Learnt from: ixlmar
Repo: NVIDIA/TensorRT-LLM PR: 8398
File: tensorrt_llm/_torch/pyexecutor/sampling_utils.py:237-272
Timestamp: 2025-10-17T13:21:31.724Z
Learning: The setup.py file in TensorRT-LLM explicitly requires Python 3.10+ via `python_requires=">=3.10, <4"`, making match/case statements and other Python 3.10+ features appropriate throughout the codebase.

Applied to files:

• requirements.txt

Learnt from: tongyuantongyu
Repo: NVIDIA/TensorRT-LLM PR: 7781
File: tests/integration/test_lists/waives.txt:313-313
Timestamp: 2025-09-17T02:48:52.732Z
Learning: In TensorRT-LLM, `tests/integration/test_lists/waives.txt` is specifically for waiving/skipping tests, while other test list files like those in `test-db/` and `qa/` directories are for different test execution contexts (pre-merge, post-merge, QA tests). The same test appearing in both waives.txt and execution list files is intentional - the test is part of test suites but will be skipped due to the waiver.

Applied to files:

• tests/integration/test_lists/test-db/l0_b200.yml
• tests/integration/test_lists/test-db/l0_dgx_b200.yml

Learnt from: moraxu
Repo: NVIDIA/TensorRT-LLM PR: 6303
File: tests/integration/test_lists/qa/examples_test_list.txt:494-494
Timestamp: 2025-07-28T17:06:08.621Z
Learning: In TensorRT-LLM testing, it's common to have both CLI flow tests (test_cli_flow.py) and PyTorch API tests (test_llm_api_pytorch.py) for the same model. These serve different purposes: CLI flow tests validate the traditional command-line workflow, while PyTorch API tests validate the newer LLM API backend. Both are legitimate and should coexist.

Applied to files:

• tests/integration/test_lists/test-db/l0_b200.yml
• tests/integration/test_lists/test-db/l0_dgx_b200.yml
• tests/integration/defs/accuracy/test_llm_api_pytorch.py

Learnt from: pengbowang-nv
Repo: NVIDIA/TensorRT-LLM PR: 7192
File: tests/integration/test_lists/test-db/l0_dgx_b200.yml:56-72
Timestamp: 2025-08-26T09:49:04.956Z
Learning: In TensorRT-LLM test configuration files, the test scheduling system handles wildcard matching with special rules that prevent duplicate test execution even when the same tests appear in multiple yaml files with overlapping GPU wildcards (e.g., "*b200*" and "*gb200*").

Applied to files:

• tests/integration/test_lists/test-db/l0_b200.yml
• tests/integration/test_lists/test-db/l0_dgx_b200.yml

Learnt from: timlee0212
Repo: NVIDIA/TensorRT-LLM PR: 6886
File: tensorrt_llm/_torch/models/modeling_deepseekv3.py:0-0
Timestamp: 2025-08-14T06:36:40.701Z
Learning: In DeepSeek V3 model (tensorrt_llm/_torch/models/modeling_deepseekv3.py), the disagreement between AllReduce.__init__ guard and _compute_mlp_tp_size logic for MNNVL usage is expected by design. The AllReduce component and MLP TP-size computation intentionally use different criteria for MNNVL availability decisions.

Applied to files:

• tests/integration/test_lists/test-db/l0_b200.yml
• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• cpp/tensorrt_llm/thop/moeUtilOp.cpp

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 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/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

Learnt from: nv-lschneider
Repo: NVIDIA/TensorRT-LLM PR: 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/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/thop/fp4Quantize.cpp

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 7104
File: cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h:999-1000
Timestamp: 2025-08-22T01:54:35.850Z
Learning: The `internal_cutlass_kernels` directory in TensorRT-LLM is a mirror of an internal NVIDIA repository and maintains its own implementation and API that may diverge from the public `cutlass_kernels` version. API inconsistencies between these two directories are intentional and by design, not bugs to be fixed.

Applied to files:

• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/CMakeLists.txt
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 7104
File: cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_gemm_kernels.h:141-145
Timestamp: 2025-08-21T02:41:10.565Z
Learning: In TensorRT-LLM MOE GEMM kernels (cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_gemm_kernels.h), the stride_act and stride_weight pointers in TmaWarpSpecializedGroupedGemmInput are intentionally declared as void* rather than typed pointers because the actual stride type is determined at runtime based on factors like the swap_ab flag and layout decisions. This runtime type determination makes compile-time type safety impossible, so void* is the correct approach.

Applied to files:

• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu

Learnt from: sklevtsov-nvidia
Repo: NVIDIA/TensorRT-LLM PR: 3294
File: cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/epilogue/fusion/sm90_visitor_scatter.hpp:0-0
Timestamp: 2025-08-08T05:10:38.906Z
Learning: The ScaledAccPerRowBiasPerColScaleScatter fusion in CUTLASS extensions (cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/epilogue/fusion/sm90_visitor_scatter.hpp) is specifically designed for per-column scaling factors only, so it uses a fixed Stride<_0,_1,int64_t> rather than conditional stride logic.

Applied to files:

• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu

Learnt from: eopXD
Repo: NVIDIA/TensorRT-LLM PR: 6767
File: cpp/tensorrt_llm/batch_manager/kvCacheManager.cpp:0-0
Timestamp: 2025-08-15T06:46:54.897Z
Learning: In cpp/tensorrt_llm/batch_manager/kvCacheManager.cpp addToken function, newly allocated blocks are unshared by design. The beam search path in addToken (when sequence.getNumTokens() > windowSize) is currently broken/non-functional with SWA, so the block allocation doesn't follow a shared-then-unshared pattern.

Applied to files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu

Learnt from: sklevtsov-nvidia
Repo: NVIDIA/TensorRT-LLM PR: 3294
File: cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_tma_warp_specialized_input.cu:118-127
Timestamp: 2025-08-09T20:57:04.084Z
Learning: In the CUTLASS MoE finalize fusion implementation (cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_tma_warp_specialized_input.cu), when setting `fused_finalize_epilogue.stride_final_output` with shape `(hidden_size, num_output_tokens, 1)`, the `num_rows_in_final_output` should be set to `num_output_tokens` (not `hidden_size`) because of a swap+transpose operation that maps rows of the output tensor to `hidden_size` and columns to `num_output_tokens`.

Applied to files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 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/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• tensorrt_llm/_torch/cute_dsl_kernels/blackwell/grouped_blockscaled_gemm_persistent.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp

Learnt from: amitz-nv
Repo: NVIDIA/TensorRT-LLM PR: 7033
File: tensorrt_llm/_torch/pyexecutor/model_engine.py:0-0
Timestamp: 2025-08-19T12:45:11.997Z
Learning: In tensorrt_llm/_torch/pyexecutor/model_engine.py, DoRA (Delta Orthogonal Rank Adaptation) functionality was removed from the PyTorch flow to eliminate issues with inverted DoRA detection logic. The original is_dora condition was checking if scaling_vec_pointer == 0, which was potentially incorrect.

Applied to files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• cpp/tensorrt_llm/thop/moeUtilOp.cpp

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 6915
File: cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu:4010-4012
Timestamp: 2025-08-14T23:23:27.449Z
Learning: For MOE (Mixture of Experts) code reviews in TensorRT-LLM, avoid repeatedly suggesting finalize fusion validation checks and safety assertions. The user djns99 has indicated these suggestions are repetitive and unwanted across multiple MOE-related changes.

Applied to files:

• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• tests/unittest/_torch/thop/parallel/test_cute_dsl_moe.py
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

Learnt from: sklevtsov-nvidia
Repo: NVIDIA/TensorRT-LLM PR: 3294
File: cpp/tensorrt_llm/kernels/internal_cutlass_kernels/include/moe_gemm_kernels.h:152-154
Timestamp: 2025-08-08T22:03:28.403Z
Learning: In CUTLASS, `TagToStrideC_t` template is defined for both pointer types (e.g., `Layout*`) and non-pointer types (e.g., `Layout`). When used with pointer types, it's often wrapped with `std::remove_pointer_t`, while non-pointer usage is direct. Both `cutlass::detail::TagToStrideC_t` and `cutlass::gemm::TagToStrideC_t` support both forms.

Applied to files:

• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 6728
File: cpp/tensorrt_llm/plugins/mixtureOfExperts/mixtureOfExpertsPlugin.cpp:966-966
Timestamp: 2025-08-08T04:10:19.038Z
Learning: TensorRT plugins currently don't support padding functionality, and TensorRT is not getting new features (in maintenance mode). This means that duplicating parameters like mExpertHiddenSize in function calls, even with TODO comments, can be acceptable as pragmatic solutions within these constraints.

Applied to files:

• cpp/tensorrt_llm/thop/moeUtilOp.cpp

Learnt from: nvchenghaoz
Repo: NVIDIA/TensorRT-LLM PR: 8469
File: tensorrt_llm/_torch/auto_deploy/models/patches/nemotron_h.py:98-116
Timestamp: 2025-10-20T17:07:18.745Z
Learning: In NemotronH models (tensorrt_llm/_torch/auto_deploy/models/patches/nemotron_h.py), the gate (self.gate) returns topk_indices and topk_weights that are already in the correct shape to be passed directly to torch_ops.auto_deploy.torch_moe without needing to reshape them when hidden_states is flattened.

Applied to files:

• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py

Learnt from: MatthiasKohl
Repo: NVIDIA/TensorRT-LLM PR: 6904
File: cpp/tensorrt_llm/pybind/thop/bindings.cpp:55-57
Timestamp: 2025-08-14T15:38:01.771Z
Learning: In TensorRT-LLM Python bindings, tensor parameter collections like mla_tensor_params and spec_decoding_tensor_params are kept as required parameters without defaults to maintain API consistency, even when it might affect backward compatibility.

Applied to files:

• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py (1)

• gen_tuning_buckets (369-373)

cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h (1)

• tensorrt_llm (39-147)

cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu (9)

• void (58-113)
• void (182-235)
• void (280-356)
• moePermute (116-160)
• moePermute (116-119)
• moeUnpermute (238-263)
• moeUnpermute (238-240)
• moeActivation (359-417)
• moeActivation (359-362)

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.h (7)

• top_k (240-240)
• n_group (241-241)
• topk_group (243-243)
• local_expert_offset (246-246)
• local_num_experts (247-247)
• getMaxPermutedPaddedCount (105-110)
• getMaxNumCtasInBatchDim (75-103)

tensorrt_llm/_torch/cute_dsl_kernels/blackwell/custom_pipeline.py (5)

• PipelineTmaUmma (73-269)
• PipelineUmmaAsync (273-376)
• producer_acquire (231-255)
• producer_tail (364-376)
• consumer_release (217-229)

tests/integration/defs/conftest.py (2)

• parametrize_with_ids (1818-1844)
• get_sm_version (1890-1893)

cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu (16)

• void (156-275)
• void (309-398)
• void (564-613)
• void (660-707)
• void (710-740)
• void (799-826)
• void (1022-1067)
• void (1070-1082)
• void (1102-1141)
• void (1143-1195)
• void (1198-1236)
• void (1240-1335)
• void (1361-1566)
• void (1683-1757)
• void (1762-1865)
• void (1943-1995)

cpp/include/tensorrt_llm/common/cudaUtils.h (1)

• getMultiProcessorCount (403-450)

tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py (10)

• GroupedGemmInputsHelper (341-415)
• get_max_num_permuted_tokens (359-360)
• infer_num_tokens (362-367)
• gen_tuning_buckets (369-373)
• map_to_tuning_buckets (375-377)
• get_max_num_tiles (351-357)
• _ (325-339)
• _ (665-683)
• _ (879-899)
• cute_dsl_nvfp4_grouped_gemm_blackwell (628-661)

tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py (1)

• cute_dsl_nvfp4_grouped_gemm_ref (93-149)

tensorrt_llm/_torch/utils.py (1)

• unswizzle_sf (144-159)

cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp (10)

• moe_sort (112-125)
• moe_sort (112-114)
• moe_permute (129-200)
• moe_permute (129-132)
• moe_unpermute (204-254)
• moe_unpermute (204-205)
• moe_swiglu (257-303)
• moe_swiglu (257-258)
• moe_swiglu_nvfp4_quantize (305-359)
• moe_swiglu_nvfp4_quantize (305-307)

tensorrt_llm/_torch/custom_ops/cpp_custom_ops.py (1)

• allgather (97-102)

tensorrt_llm/_torch/utils.py (2)

• Fp4QuantizedTensor (110-117)
• shape (116-117)

tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py (1)

• cute_dsl_nvfp4_grouped_gemm_blackwell (628-661)

tensorrt_llm/_torch/autotuner.py (6)

• AutoTuner (514-1181)
• TuningConfig (53-101)
• get_valid_tactics (156-174)
• forward (180-206)
• get (545-548)
• choose_one (623-778)

tensorrt_llm/_torch/utils.py (3)

• fp4_scale_infer_shape (262-267)
• get_last_power_of_2_num_tokens_buckets (252-259)
• last_positive_power_of_2 (229-234)

tensorrt_llm/_torch/cute_dsl_kernels/blackwell/dense_blockscaled_gemm_persistent.py (3)

• Sm100BlockScaledPersistentDenseGemmKernel (65-1877)
• Sm100BlockScaledPersistentDenseGemmKernelWrapper (1901-2015)
• can_implement (1821-1877)

tensorrt_llm/_torch/cute_dsl_kernels/blackwell/grouped_blockscaled_gemm_persistent.py (3)

• Sm100BlockScaledPersistentGroupedGemmKernel (56-2285)
• can_implement (2142-2225)
• wrapper (2228-2285)

tensorrt_llm/_torch/cute_dsl_kernels/blackwell/utils.py (2)

• make_ptr (142-188)
• dtype (118-119)

[error] 18-18: 'tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h' file not found

(clang-diagnostic-error)

[error] 17-17: 'tensorrt_llm/kernels/cuteDslKernels/moeUtils.h' file not found

(clang-diagnostic-error)

389-389: Avoid specifying long messages outside the exception class

(TRY003)

695-695: Unpacked variable bidy is never used

Prefix it with an underscore or any other dummy variable pattern

(RUF059)

695-695: Unpacked variable bidz is never used

Prefix it with an underscore or any other dummy variable pattern

(RUF059)

1060-1060: Loop control variable k_tile not used within loop body

Rename unused k_tile to _k_tile

(B007)

1282-1282: Loop control variable k_tile not used within loop body

Rename unused k_tile to _k_tile

(B007)

1701-1701: Unused method argument: tidx

(ARG002)

1928-1928: Avoid specifying long messages outside the exception class

(TRY003)

306-306: Loop control variable i not used within loop body

Rename unused i to _i

(B007)

117-117: Ambiguous variable name: l

(E741)

136-136: Consider [0, *num_tiles_per_expert.cumsum(dim=0).tolist()] instead of concatenation

Replace with [0, *num_tiles_per_expert.cumsum(dim=0).tolist()]

(RUF005)

139-139: zip() without an explicit strict= parameter

Add explicit value for parameter strict=

(B905)

139-139: Prefer itertools.pairwise() over zip() when iterating over successive pairs

Replace zip() with itertools.pairwise()

(RUF007)

223-223: Unused method argument: output_dtype

(ARG002)

224-224: Unused method argument: all_rank_num_tokens

(ARG002)

225-225: Unused method argument: use_dp_padding

(ARG002)

369-369: Unpacked variable total_num_padded_tokens is never used

Prefix it with an underscore or any other dummy variable pattern

(RUF059)

456-458: Avoid specifying long messages outside the exception class

(TRY003)

418-418: Mutable class attributes should be annotated with typing.ClassVar

(RUF012)

419-419: Mutable class attributes should be annotated with typing.ClassVar

(RUF012)

441-443: Avoid specifying long messages outside the exception class

(TRY003)

448-448: Unused method argument: profile

(ARG002)

449-449: Unused method argument: kwargs

(ARG002)

453-453: Ambiguous variable name: l

(E741)

519-519: Ambiguous variable name: l

(E741)

668-668: Unused function argument: input_scale

(ARG001)

669-669: Unused function argument: weight_scale

(ARG001)

670-670: Unused function argument: alpha

(ARG001)

671-671: Unused function argument: tile_idx_to_group_idx

(ARG001)

672-672: Unused function argument: num_non_exiting_tiles

(ARG001)

673-673: Unused function argument: num_experts

(ARG001)

674-674: Unused function argument: top_k

(ARG001)

675-675: Unused function argument: num_local_experts

(ARG001)

676-676: Unused function argument: local_expert_offset

(ARG001)

677-677: Unused function argument: tile_size

(ARG001)

679-679: Unused function argument: scaling_vector_size

(ARG001)

717-717: Mutable class attributes should be annotated with typing.ClassVar

(RUF012)

738-738: Unused method argument: inputs

(ARG002)

739-739: Unused method argument: profile

(ARG002)

740-740: Unused method argument: kwargs

(ARG002)

772-772: Unpacked variable total_num_padded_tokens is never used

Prefix it with an underscore or any other dummy variable pattern

(RUF059)

881-881: Unused function argument: input_scale

(ARG001)

882-882: Unused function argument: token_selected_experts

(ARG001)

883-883: Unused function argument: token_final_scales

(ARG001)

884-884: Unused function argument: gemm1_weight

(ARG001)

885-885: Unused function argument: gemm1_weight_scale

(ARG001)

886-886: Unused function argument: gemm1_alpha

(ARG001)

887-887: Unused function argument: gemm2_input_global_scale

(ARG001)

888-888: Unused function argument: gemm2_weight

(ARG001)

889-889: Unused function argument: gemm2_weight_scale

(ARG001)

890-890: Unused function argument: gemm2_alpha

(ARG001)

891-891: Unused function argument: num_experts

(ARG001)

892-892: Unused function argument: top_k

(ARG001)

893-893: Unused function argument: num_local_experts

(ARG001)

894-894: Unused function argument: local_expert_offset

(ARG001)

896-896: Unused function argument: scaling_vector_size

(ARG001)

📄 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/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/thop/fp4Quantize.cpp

📄 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/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/thop/fp4Quantize.cpp

📄 CodeRabbit inference engine (CODING_GUIDELINES.md)

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

Files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• tensorrt_llm/_torch/autotuner.py
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• tensorrt_llm/_torch/cute_dsl_kernels/blackwell/grouped_blockscaled_gemm_persistent.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py
• tests/integration/defs/accuracy/test_llm_api_pytorch.py
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• tests/unittest/_torch/thop/parallel/test_cute_dsl_moe.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/thop/fp4Quantize.cpp

📄 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/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• tensorrt_llm/_torch/autotuner.py
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• tensorrt_llm/_torch/cute_dsl_kernels/blackwell/grouped_blockscaled_gemm_persistent.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py
• tests/integration/defs/accuracy/test_llm_api_pytorch.py
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• tests/unittest/_torch/thop/parallel/test_cute_dsl_moe.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/thop/fp4Quantize.cpp

📄 CodeRabbit inference engine (CODING_GUIDELINES.md)

**/*.py: Python code must target Python 3.8+.
Indent Python code with 4 spaces; do not use tabs.
Maintain module namespace when importing; prefer 'from package.subpackage import foo' then 'foo.SomeClass()' instead of importing the class directly.
Python filenames should be snake_case (e.g., some_file.py).
Python classes use PascalCase names.
Functions and methods use snake_case names.
Local variables use snake_case; prefix 'k' for variables that start with a number (e.g., k_99th_percentile).
Global variables use upper SNAKE_CASE prefixed with 'G' (e.g., G_MY_GLOBAL).
Constants use upper SNAKE_CASE (e.g., MY_CONSTANT).
Avoid shadowing variables from an outer scope.
Initialize all externally visible members of a class in the constructor.
Prefer docstrings for interfaces that may be used outside a file; comments for in-function or file-local interfaces.
Use Google-style docstrings for classes and functions (Sphinx-parsable).
Document attributes and variables inline so they render under the class/function docstring.
Avoid reflection when a simpler, explicit approach suffices (e.g., avoid dict(**locals()) patterns).
In try/except, catch the most specific exceptions possible.
For duck-typing try/except, keep the try body minimal and use else for the main logic.

Files:

• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• tensorrt_llm/_torch/autotuner.py
• tensorrt_llm/_torch/cute_dsl_kernels/blackwell/grouped_blockscaled_gemm_persistent.py
• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py
• tests/integration/defs/accuracy/test_llm_api_pytorch.py
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py
• tests/unittest/_torch/thop/parallel/test_cute_dsl_moe.py

📄 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/thop/moeUtilOp.cpp
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp
• cpp/tensorrt_llm/thop/fp4Quantize.cpp

📄 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/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.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/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h

Learnt from: ixlmar
Repo: NVIDIA/TensorRT-LLM PR: 8398
File: tensorrt_llm/_torch/pyexecutor/sampling_utils.py:237-272
Timestamp: 2025-10-17T13:21:31.724Z
Learning: The setup.py file in TensorRT-LLM explicitly requires Python 3.10+ via `python_requires=">=3.10, <4"`, making match/case statements and other Python 3.10+ features appropriate throughout the codebase.

Applied to files:

• requirements.txt

Learnt from: jhaotingc
Repo: NVIDIA/TensorRT-LLM PR: 7856
File: cpp/tensorrt_llm/thop/fp8BlockScaleMoe.cpp:159-166
Timestamp: 2025-09-19T21:28:13.751Z
Learning: In TensorRT-LLM blockScaleMoe routing (cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu), the DeepSeek routing method performs reinterpret_cast<float*>(routingLogits) at line 89, which could cause issues if routing_logits are BF16. However, Qwen3-FP8 models use RenormalizeNaive routing method and are not affected by this dtype casting issue.

Applied to files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp

Learnt from: sklevtsov-nvidia
Repo: NVIDIA/TensorRT-LLM PR: 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/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu

Learnt from: nv-lschneider
Repo: NVIDIA/TensorRT-LLM PR: 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/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu

Learnt from: thorjohnsen
Repo: NVIDIA/TensorRT-LLM PR: 6910
File: cpp/tensorrt_llm/batch_manager/kvCacheManager.cpp:0-0
Timestamp: 2025-08-14T21:04:50.248Z
Learning: In KV cache onboarding logic during prefill in cpp/tensorrt_llm/batch_manager/kvCacheManager.cpp, when calculating which blocks fall within the attention window, use getTokensPerBlock() to advance token indices rather than block->getUniqueTokens().size(), because the calculation needs to consider the post-prefill state where blocks will be filled to capacity, not their current token count.

Applied to files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu

Learnt from: eopXD
Repo: NVIDIA/TensorRT-LLM PR: 6767
File: cpp/tensorrt_llm/batch_manager/kvCacheManager.cpp:0-0
Timestamp: 2025-08-15T06:46:54.897Z
Learning: In cpp/tensorrt_llm/batch_manager/kvCacheManager.cpp addToken function, newly allocated blocks are unshared by design. The beam search path in addToken (when sequence.getNumTokens() > windowSize) is currently broken/non-functional with SWA, so the block allocation doesn't follow a shared-then-unshared pattern.

Applied to files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu

Learnt from: ChristinaZ
Repo: NVIDIA/TensorRT-LLM PR: 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/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp

Learnt from: sklevtsov-nvidia
Repo: NVIDIA/TensorRT-LLM PR: 3294
File: cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_tma_warp_specialized_input.cu:118-127
Timestamp: 2025-08-09T20:57:04.084Z
Learning: In the CUTLASS MoE finalize fusion implementation (cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_tma_warp_specialized_input.cu), when setting `fused_finalize_epilogue.stride_final_output` with shape `(hidden_size, num_output_tokens, 1)`, the `num_rows_in_final_output` should be set to `num_output_tokens` (not `hidden_size`) because of a swap+transpose operation that maps rows of the output tensor to `hidden_size` and columns to `num_output_tokens`.

Applied to files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py
• tests/integration/defs/accuracy/test_llm_api_pytorch.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 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/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• tests/integration/defs/accuracy/test_llm_api_pytorch.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingDeepSeek.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp

Learnt from: nv-lschneider
Repo: NVIDIA/TensorRT-LLM PR: 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/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/CMakeLists.txt
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingKernel.cuh
• cpp/tensorrt_llm/kernels/cuteDslKernels/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 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/trtllmGenKernels/blockScaleMoe/RoutingLlama4.cu
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• tensorrt_llm/_torch/cute_dsl_kernels/blackwell/grouped_blockscaled_gemm_persistent.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 7104
File: cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/epilogue/fusion/sm90_visitor_scatter.hpp:399-417
Timestamp: 2025-08-21T21:48:35.135Z
Learning: CUTLASS extensions in TensorRT-LLM (located under cpp/tensorrt_llm/cutlass_extensions/) are designed to integrate with and extend functionality in the external CUTLASS repository. When analyzing these extensions, their consumers and functionality wiring may exist in the CUTLASS codebase rather than within TensorRT-LLM itself.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/cuteDslKernels/CMakeLists.txt
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp

Learnt from: tongyuantongyu
Repo: NVIDIA/TensorRT-LLM PR: 6984
File: cpp/tensorrt_llm/CMakeLists.txt:297-299
Timestamp: 2025-08-18T09:08:07.687Z
Learning: In the TensorRT-LLM project, artifacts are manually copied rather than installed via `cmake --install`, so INSTALL_RPATH properties are not needed - only BUILD_RPATH affects the final artifacts.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt

Learnt from: sklevtsov-nvidia
Repo: NVIDIA/TensorRT-LLM PR: 3294
File: cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/epilogue/fusion/sm90_visitor_scatter.hpp:36-36
Timestamp: 2025-08-08T05:06:31.596Z
Learning: CUTLASS extension files (under cpp/tensorrt_llm/cutlass_extensions/) follow CUTLASS coding style conventions, including using #pragma once instead of TRTLLM_ prefixed header guards, even though they are .hpp files.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/CMakeLists.txt
• cpp/tensorrt_llm/thop/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cuteDslKernels/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp

Learnt from: nv-lschneider
Repo: NVIDIA/TensorRT-LLM PR: 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/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h

Learnt from: galagam
Repo: NVIDIA/TensorRT-LLM PR: 6487
File: tests/unittest/_torch/auto_deploy/unit/singlegpu/test_ad_trtllm_bench.py:1-12
Timestamp: 2025-08-06T13:58:07.506Z
Learning: In TensorRT-LLM, test files (files under tests/ directories) do not require NVIDIA copyright headers, unlike production source code files. Test files typically start directly with imports, docstrings, or code.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt
• tests/integration/test_lists/test-db/l0_dgx_b200.yml

Learnt from: fredricz-20070104
Repo: NVIDIA/TensorRT-LLM PR: 7645
File: tests/integration/test_lists/qa/llm_function_core.txt:648-648
Timestamp: 2025-09-09T09:40:45.658Z
Learning: In TensorRT-LLM test lists, it's common and intentional for the same test to appear in multiple test list files when they serve different purposes (e.g., llm_function_core.txt for comprehensive core functionality testing and llm_function_core_sanity.txt for quick sanity checks). This duplication allows tests to be run in different testing contexts.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt
• tests/integration/test_lists/test-db/l0_b200.yml
• tests/integration/test_lists/test-db/l0_dgx_b200.yml
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

Learnt from: tongyuantongyu
Repo: NVIDIA/TensorRT-LLM PR: 7763
File: cpp/tensorrt_llm/CMakeLists.txt:297-301
Timestamp: 2025-09-16T09:30:09.716Z
Learning: In the TensorRT-LLM project, NCCL libraries are loaded earlier by PyTorch libraries or the bindings library, so the main shared library doesn't need NCCL paths in its RPATH - the libraries will already be available in the process address space when needed.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt

Learnt from: nv-lschneider
Repo: NVIDIA/TensorRT-LLM PR: 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/CMakeLists.txt
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/thop/fp4Quantize.cpp

Learnt from: xinhe-nv
Repo: NVIDIA/TensorRT-LLM PR: 8534
File: scripts/format_test_list.py:1-6
Timestamp: 2025-10-22T06:53:47.017Z
Learning: The file `scripts/format_test_list.py` in the TensorRT-LLM repository does not require the NVIDIA Apache-2.0 copyright header.

Applied to files:

• cpp/tensorrt_llm/CMakeLists.txt

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 6728
File: cpp/tensorrt_llm/plugins/mixtureOfExperts/mixtureOfExpertsPlugin.cpp:966-966
Timestamp: 2025-08-08T04:10:19.038Z
Learning: TensorRT plugins currently don't support padding functionality, and TensorRT is not getting new features (in maintenance mode). This means that duplicating parameters like mExpertHiddenSize in function calls, even with TODO comments, can be acceptable as pragmatic solutions within these constraints.

Applied to files:

• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/RoutingRenormalize.cu
• cpp/tensorrt_llm/thop/moeUtilOp.cpp

Learnt from: tongyuantongyu
Repo: NVIDIA/TensorRT-LLM PR: 7781
File: tests/integration/test_lists/waives.txt:313-313
Timestamp: 2025-09-17T02:48:52.732Z
Learning: In TensorRT-LLM, `tests/integration/test_lists/waives.txt` is specifically for waiving/skipping tests, while other test list files like those in `test-db/` and `qa/` directories are for different test execution contexts (pre-merge, post-merge, QA tests). The same test appearing in both waives.txt and execution list files is intentional - the test is part of test suites but will be skipped due to the waiver.

Applied to files:

• tests/integration/test_lists/test-db/l0_b200.yml
• tests/integration/test_lists/test-db/l0_dgx_b200.yml

Learnt from: pengbowang-nv
Repo: NVIDIA/TensorRT-LLM PR: 7192
File: tests/integration/test_lists/test-db/l0_dgx_b200.yml:56-72
Timestamp: 2025-08-26T09:49:04.956Z
Learning: In TensorRT-LLM test configuration files, the test scheduling system handles wildcard matching with special rules that prevent duplicate test execution even when the same tests appear in multiple yaml files with overlapping GPU wildcards (e.g., "*b200*" and "*gb200*").

Applied to files:

• tests/integration/test_lists/test-db/l0_b200.yml
• tests/integration/test_lists/test-db/l0_dgx_b200.yml

Learnt from: moraxu
Repo: NVIDIA/TensorRT-LLM PR: 6303
File: tests/integration/test_lists/qa/examples_test_list.txt:494-494
Timestamp: 2025-07-28T17:06:08.621Z
Learning: In TensorRT-LLM testing, it's common to have both CLI flow tests (test_cli_flow.py) and PyTorch API tests (test_llm_api_pytorch.py) for the same model. These serve different purposes: CLI flow tests validate the traditional command-line workflow, while PyTorch API tests validate the newer LLM API backend. Both are legitimate and should coexist.

Applied to files:

• tests/integration/test_lists/test-db/l0_b200.yml
• tests/integration/test_lists/test-db/l0_dgx_b200.yml
• tests/integration/defs/accuracy/test_llm_api_pytorch.py

Learnt from: timlee0212
Repo: NVIDIA/TensorRT-LLM PR: 6886
File: tensorrt_llm/_torch/models/modeling_deepseekv3.py:0-0
Timestamp: 2025-08-14T06:36:40.701Z
Learning: In DeepSeek V3 model (tensorrt_llm/_torch/models/modeling_deepseekv3.py), the disagreement between AllReduce.__init__ guard and _compute_mlp_tp_size logic for MNNVL usage is expected by design. The AllReduce component and MLP TP-size computation intentionally use different criteria for MNNVL availability decisions.

Applied to files:

• tests/integration/test_lists/test-db/l0_b200.yml
• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• cpp/tensorrt_llm/thop/moeUtilOp.cpp

Learnt from: amitz-nv
Repo: NVIDIA/TensorRT-LLM PR: 7033
File: tensorrt_llm/_torch/pyexecutor/model_engine.py:0-0
Timestamp: 2025-08-19T12:45:11.997Z
Learning: In tensorrt_llm/_torch/pyexecutor/model_engine.py, DoRA (Delta Orthogonal Rank Adaptation) functionality was removed from the PyTorch flow to eliminate issues with inverted DoRA detection logic. The original is_dora condition was checking if scaling_vec_pointer == 0, which was potentially incorrect.

Applied to files:

• tensorrt_llm/_torch/models/modeling_deepseekv3.py
• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• tensorrt_llm/_torch/autotuner.py
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 6915
File: cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu:4010-4012
Timestamp: 2025-08-14T23:23:27.449Z
Learning: For MOE (Mixture of Experts) code reviews in TensorRT-LLM, avoid repeatedly suggesting finalize fusion validation checks and safety assertions. The user djns99 has indicated these suggestions are repetitive and unwanted across multiple MOE-related changes.

Applied to files:

• cpp/tensorrt_llm/thop/moeUtilOp.cpp
• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py
• tests/integration/defs/accuracy/test_llm_api_pytorch.py
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu
• cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp

Learnt from: nvchenghaoz
Repo: NVIDIA/TensorRT-LLM PR: 8469
File: tensorrt_llm/_torch/auto_deploy/transform/library/rms_norm.py:180-182
Timestamp: 2025-10-20T17:09:21.560Z
Learning: In tensorrt_llm/_torch/auto_deploy/transform/library/rms_norm.py, the _gated_rmsnorm_replacement function does not need to cast the output of torch.ops.auto_deploy.torch_rmsnorm_gated back to the input dtype, even though the custom op returns fp32. The dtype handling is managed elsewhere or the fp32 output is acceptable for downstream consumers.

Applied to files:

• tensorrt_llm/_torch/autotuner.py

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 7104
File: cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_gemm_kernels.h:141-145
Timestamp: 2025-08-21T02:41:10.565Z
Learning: In TensorRT-LLM MOE GEMM kernels (cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_gemm_kernels.h), the stride_act and stride_weight pointers in TmaWarpSpecializedGroupedGemmInput are intentionally declared as void* rather than typed pointers because the actual stride type is determined at runtime based on factors like the swap_ab flag and layout decisions. This runtime type determination makes compile-time type safety impossible, so void* is the correct approach.

Applied to files:

• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.cu
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h

Learnt from: djns99
Repo: NVIDIA/TensorRT-LLM PR: 7104
File: cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h:999-1000
Timestamp: 2025-08-22T01:54:35.850Z
Learning: The `internal_cutlass_kernels` directory in TensorRT-LLM is a mirror of an internal NVIDIA repository and maintains its own implementation and API that may diverge from the public `cutlass_kernels` version. API inconsistencies between these two directories are intentional and by design, not bugs to be fixed.

Applied to files:

• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.h
• cpp/tensorrt_llm/kernels/CMakeLists.txt
• tensorrt_llm/_torch/cute_dsl_kernels/blackwell/grouped_blockscaled_gemm_persistent.py
• cpp/tensorrt_llm/kernels/cuteDslKernels/CMakeLists.txt
• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h

Learnt from: sklevtsov-nvidia
Repo: NVIDIA/TensorRT-LLM PR: 3294
File: cpp/tensorrt_llm/kernels/internal_cutlass_kernels/include/moe_gemm_kernels.h:152-154
Timestamp: 2025-08-08T22:03:28.403Z
Learning: In CUTLASS, `TagToStrideC_t` template is defined for both pointer types (e.g., `Layout*`) and non-pointer types (e.g., `Layout`). When used with pointer types, it's often wrapped with `std::remove_pointer_t`, while non-pointer usage is direct. Both `cutlass::detail::TagToStrideC_t` and `cutlass::gemm::TagToStrideC_t` support both forms.

Applied to files:

• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu

Learnt from: sklevtsov-nvidia
Repo: NVIDIA/TensorRT-LLM PR: 3294
File: cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/epilogue/fusion/sm90_visitor_scatter.hpp:0-0
Timestamp: 2025-08-08T05:10:38.906Z
Learning: The ScaledAccPerRowBiasPerColScaleScatter fusion in CUTLASS extensions (cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/epilogue/fusion/sm90_visitor_scatter.hpp) is specifically designed for per-column scaling factors only, so it uses a fixed Stride<_0,_1,int64_t> rather than conditional stride logic.

Applied to files:

• cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu
• cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h

Learnt from: nvchenghaoz
Repo: NVIDIA/TensorRT-LLM PR: 8469
File: tensorrt_llm/_torch/auto_deploy/models/patches/nemotron_h.py:98-116
Timestamp: 2025-10-20T17:07:18.745Z
Learning: In NemotronH models (tensorrt_llm/_torch/auto_deploy/models/patches/nemotron_h.py), the gate (self.gate) returns topk_indices and topk_weights that are already in the correct shape to be passed directly to torch_ops.auto_deploy.torch_moe without needing to reshape them when hidden_states is flattened.

Applied to files:

• tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py

Learnt from: MatthiasKohl
Repo: NVIDIA/TensorRT-LLM PR: 6904
File: cpp/tensorrt_llm/pybind/thop/bindings.cpp:55-57
Timestamp: 2025-08-14T15:38:01.771Z
Learning: In TensorRT-LLM Python bindings, tensor parameter collections like mla_tensor_params and spec_decoding_tensor_params are kept as required parameters without defaults to maintain API consistency, even when it might affect backward compatibility.

Applied to files:

• tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py (1)

• gen_tuning_buckets (369-373)

cpp/tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h (1)

• tensorrt_llm (39-147)

cpp/tensorrt_llm/kernels/cuteDslKernels/moeUtils.cu (9)

• void (58-113)
• void (182-235)
• void (280-356)
• moePermute (116-160)
• moePermute (116-119)
• moeUnpermute (238-263)
• moeUnpermute (238-240)
• moeActivation (359-417)
• moeActivation (359-362)

tensorrt_llm/_torch/cute_dsl_kernels/blackwell/custom_pipeline.py (5)

• PipelineTmaUmma (73-269)
• PipelineUmmaAsync (273-376)
• producer_acquire (231-255)
• producer_tail (364-376)
• consumer_release (217-229)

tensorrt_llm/_torch/custom_ops/cpp_custom_ops.py (1)

• allgather (97-102)

tensorrt_llm/_torch/utils.py (2)

• Fp4QuantizedTensor (110-117)
• shape (116-117)

cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp (8)

• moe_sort (112-125)
• moe_sort (112-114)
• moe_permute (129-200)
• moe_permute (129-132)
• moe_swiglu_nvfp4_quantize (305-359)
• moe_swiglu_nvfp4_quantize (305-307)
• moe_unpermute (204-254)
• moe_unpermute (204-205)

tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py (1)

• cute_dsl_nvfp4_grouped_gemm_blackwell (628-661)

tests/integration/defs/conftest.py (2)

• parametrize_with_ids (1818-1844)
• get_sm_version (1890-1893)

tensorrt_llm/_utils.py (1)

• get_sm_version (733-735)

tensorrt_llm/_torch/autotuner.py (4)

• ConstraintSpec (39-49)
• DynamicTensorSpec (23-35)
• TuningConfig (53-101)
• get (545-548)

tensorrt_llm/_torch/utils.py (3)

• fp4_scale_infer_shape (262-267)
• get_last_power_of_2_num_tokens_buckets (252-259)
• last_positive_power_of_2 (229-234)

tensorrt_llm/_torch/cute_dsl_kernels/blackwell/dense_blockscaled_gemm_persistent.py (3)

• Sm100BlockScaledPersistentDenseGemmKernel (65-1877)
• Sm100BlockScaledPersistentDenseGemmKernelWrapper (1901-2015)
• can_implement (1821-1877)

tensorrt_llm/_torch/cute_dsl_kernels/blackwell/grouped_blockscaled_gemm_persistent.py (3)

• Sm100BlockScaledPersistentGroupedGemmKernel (56-2285)
• can_implement (2142-2225)
• wrapper (2228-2285)

tensorrt_llm/_torch/cute_dsl_kernels/blackwell/utils.py (2)

• make_ptr (142-188)
• dtype (118-119)

tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py (10)

• GroupedGemmInputsHelper (341-415)
• get_max_num_permuted_tokens (359-360)
• infer_num_tokens (362-367)
• gen_tuning_buckets (369-373)
• map_to_tuning_buckets (375-377)
• get_max_num_tiles (351-357)
• _ (325-339)
• _ (665-683)
• _ (879-899)
• cute_dsl_nvfp4_grouped_gemm_blackwell (628-661)

tensorrt_llm/_torch/modules/fused_moe/fused_moe_cute_dsl.py (1)

• cute_dsl_nvfp4_grouped_gemm_ref (93-149)

tensorrt_llm/_torch/utils.py (1)

• unswizzle_sf (144-159)

cpp/tensorrt_llm/thop/cuteDslMoeUtilsOp.cpp (10)

• moe_sort (112-125)
• moe_sort (112-114)
• moe_permute (129-200)
• moe_permute (129-132)
• moe_unpermute (204-254)
• moe_unpermute (204-205)
• moe_swiglu (257-303)
• moe_swiglu (257-258)
• moe_swiglu_nvfp4_quantize (305-359)
• moe_swiglu_nvfp4_quantize (305-307)

cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu (16)

• void (156-275)
• void (309-398)
• void (564-613)
• void (660-707)
• void (710-740)
• void (799-826)
• void (1022-1067)
• void (1070-1082)
• void (1102-1141)
• void (1143-1195)
• void (1198-1236)
• void (1240-1335)
• void (1361-1566)
• void (1683-1757)
• void (1762-1865)
• void (1943-1995)

cpp/include/tensorrt_llm/common/cudaUtils.h (1)

• getMultiProcessorCount (403-450)

cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/runner.h (7)

• top_k (240-240)
• n_group (241-241)
• topk_group (243-243)
• local_expert_offset (246-246)
• local_num_experts (247-247)
• getMaxPermutedPaddedCount (105-110)
• getMaxNumCtasInBatchDim (75-103)

[error] 18-18: 'tensorrt_llm/kernels/cutlass_kernels/include/moe_kernels.h' file not found

(clang-diagnostic-error)

[error] 17-17: 'tensorrt_llm/kernels/cuteDslKernels/moeUtils.h' file not found

(clang-diagnostic-error)

389-389: Avoid specifying long messages outside the exception class

(TRY003)

695-695: Unpacked variable bidy is never used

Prefix it with an underscore or any other dummy variable pattern

(RUF059)

695-695: Unpacked variable bidz is never used

Prefix it with an underscore or any other dummy variable pattern

(RUF059)

1060-1060: Loop control variable k_tile not used within loop body

Rename unused k_tile to _k_tile

(B007)

1282-1282: Loop control variable k_tile not used within loop body

Rename unused k_tile to _k_tile

(B007)

1701-1701: Unused method argument: tidx

(ARG002)

1928-1928: Avoid specifying long messages outside the exception class

(TRY003)

117-117: Ambiguous variable name: l

(E741)

136-136: Consider [0, *num_tiles_per_expert.cumsum(dim=0).tolist()] instead of concatenation

Replace with [0, *num_tiles_per_expert.cumsum(dim=0).tolist()]

(RUF005)

139-139: zip() without an explicit strict= parameter

Add explicit value for parameter strict=

(B905)

139-139: Prefer itertools.pairwise() over zip() when iterating over successive pairs

Replace zip() with itertools.pairwise()

(RUF007)

223-223: Unused method argument: output_dtype

(ARG002)

224-224: Unused method argument: all_rank_num_tokens

(ARG002)

225-225: Unused method argument: use_dp_padding

(ARG002)

369-369: Unpacked variable total_num_padded_tokens is never used

Prefix it with an underscore or any other dummy variable pattern

(RUF059)

456-458: Avoid specifying long messages outside the exception class

(TRY003)

418-418: Mutable class attributes should be annotated with typing.ClassVar

(RUF012)

419-419: Mutable class attributes should be annotated with typing.ClassVar

(RUF012)

441-443: Avoid specifying long messages outside the exception class

(TRY003)

448-448: Unused method argument: profile

(ARG002)

449-449: Unused method argument: kwargs

(ARG002)

453-453: Ambiguous variable name: l

(E741)

519-519: Ambiguous variable name: l

(E741)

668-668: Unused function argument: input_scale

(ARG001)

669-669: Unused function argument: weight_scale

(ARG001)

670-670: Unused function argument: alpha

(ARG001)

671-671: Unused function argument: tile_idx_to_group_idx

(ARG001)

672-672: Unused function argument: num_non_exiting_tiles

(ARG001)

673-673: Unused function argument: num_experts

(ARG001)

674-674: Unused function argument: top_k

(ARG001)

675-675: Unused function argument: num_local_experts

(ARG001)

676-676: Unused function argument: local_expert_offset

(ARG001)

677-677: Unused function argument: tile_size

(ARG001)

679-679: Unused function argument: scaling_vector_size

(ARG001)

717-717: Mutable class attributes should be annotated with typing.ClassVar

(RUF012)

738-738: Unused method argument: inputs

(ARG002)

739-739: Unused method argument: profile

(ARG002)

740-740: Unused method argument: kwargs

(ARG002)

772-772: Unpacked variable total_num_padded_tokens is never used

Prefix it with an underscore or any other dummy variable pattern

(RUF059)

881-881: Unused function argument: input_scale

(ARG001)

882-882: Unused function argument: token_selected_experts

(ARG001)

883-883: Unused function argument: token_final_scales

(ARG001)

884-884: Unused function argument: gemm1_weight

(ARG001)

885-885: Unused function argument: gemm1_weight_scale

(ARG001)

886-886: Unused function argument: gemm1_alpha

(ARG001)

887-887: Unused function argument: gemm2_input_global_scale

(ARG001)

888-888: Unused function argument: gemm2_weight

(ARG001)

889-889: Unused function argument: gemm2_weight_scale

(ARG001)

890-890: Unused function argument: gemm2_alpha

(ARG001)

891-891: Unused function argument: num_experts

(ARG001)

892-892: Unused function argument: top_k

(ARG001)

893-893: Unused function argument: num_local_experts

(ARG001)

894-894: Unused function argument: local_expert_offset

(ARG001)

896-896: Unused function argument: scaling_vector_size

(ARG001)

306-306: Loop control variable i not used within loop body

Rename unused i to _i

(B007)