MXFP4 x BF16 CUTLASS MoE backend perf and profiling improvement on Hopper #8721
Conversation
Signed-off-by: Jiang Shao <[email protected]>
… Hopper MXFP4 x BF16 Grouped GEMM. Signed-off-by: Jiang Shao <[email protected]>
Signed-off-by: Jiang Shao <[email protected]>
StudyingShao
requested review from
Barry-Delaney,
djns99 and
hlu1
and removed request for
hlu1
📝 WalkthroughWalkthroughThe changes add FP4 interleaving support for Hopper mixed GEMM operations, introduce device-level conversion utilities for FP4/BF16 format handling with FP8 scaling overrides, refactor GEMM heuristic configuration selection to enforce COOPERATIVE scheduling and restrict certain tile sizes, and integrate FP4 interleaving into MoE weight preparation pipelines via PyTorch bindings. Changes
Estimated code review effort🎯 3 (Moderate) | ⏱️ ~20 minutes
Pre-merge checks and finishing touches❌ Failed checks (2 warnings)
✅ Passed checks (1 passed)
✨ Finishing touches
🧪 Generate unit tests (beta)
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 5
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
tensorrt_llm/_torch/modules/fused_moe/quantization.py (1)
1393-1402: Fix typo in error messagew3_weight_shard referenced twice. Use the intended variable for clarity.
- f"Invalid shape of w1_weight_shard {w1_weight_shard.shape} and w3_weight_shard {w1_weight_shard.shape}" + f"Invalid shape of w1_weight_shard {w1_weight_shard.shape} and w3_weight_shard {w3_weight_shard.shape}"
🧹 Nitpick comments (10)
cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp (1)
319-330: Consider refactoring repetitive config construction.The four candidate configurations differ only in their cluster shape parameter. This repetition could be reduced by iterating over the cluster shapes.
Apply this diff to refactor:
- CutlassGemmConfig candidate( - tile_config, mainloop_schedule, epilogue_schedule, ClusterShape::ClusterShape_1x1x1); - candidate_configs.push_back(candidate); - candidate = CutlassGemmConfig( - tile_config, mainloop_schedule, epilogue_schedule, ClusterShape::ClusterShape_2x1x1); - candidate_configs.push_back(candidate); - candidate = CutlassGemmConfig( - tile_config, mainloop_schedule, epilogue_schedule, ClusterShape::ClusterShape_1x2x1); - candidate_configs.push_back(candidate); - candidate = CutlassGemmConfig( - tile_config, mainloop_schedule, epilogue_schedule, ClusterShape::ClusterShape_2x2x1); - candidate_configs.push_back(candidate); + ClusterShape const cluster_shapes[] = {ClusterShape::ClusterShape_1x1x1, + ClusterShape::ClusterShape_2x1x1, ClusterShape::ClusterShape_1x2x1, + ClusterShape::ClusterShape_2x2x1}; + for (auto const cluster_shape : cluster_shapes) + { + CutlassGemmConfig candidate(tile_config, mainloop_schedule, epilogue_schedule, cluster_shape); + candidate_configs.push_back(candidate); + }cpp/tensorrt_llm/thop/weightOnlyQuantOp.cpp (2)
402-425: Propagate CUDA stream and add post-kernel error checksThe kernel is launched from the .cu side on the default stream; propagate the current stream and check errors to avoid silent failures and stream hazards.
Apply these coordinated changes:
- In this file (pass stream):
- interleave_fp4_for_Hopper_mixed_gemm(weight_ptr, weight_interleaved_ptr, n, k); + interleave_fp4_for_Hopper_mixed_gemm( + weight_ptr, weight_interleaved_ptr, n, k, at::cuda::getCurrentCUDAStream());
- In moe_gemm_mixed_utils.h (new signature):
-void interleave_fp4_for_Hopper_mixed_gemm(uint8_t* weight, uint8_t* weight_interleaved, int const rows, int const cols); +void interleave_fp4_for_Hopper_mixed_gemm( + uint8_t* weight, uint8_t* weight_interleaved, int const rows, int const cols, cudaStream_t stream);
- In moe_gemm_mixed_utils.cu (use stream + error check):
- interleave_fp4_for_Hopper_mixed_gemm_kernel<<<1024, 1024>>>(weight, weight_interleaved, rows, cols); + interleave_fp4_for_Hopper_mixed_gemm_kernel<<<grid, block, 0, stream>>>(weight, weight_interleaved, rows, cols); + CUDA_CHECK(cudaGetLastError());Also compute grid/block from rows/cols instead of fixed 1024x1024.
Please confirm the kernel expects column-major buffers. If it does, clarify or adjust input layout to avoid interleave mismatches with PyTorch’s row-major tensors. Based on learnings
462-463: Prefer modern Torch operator registrationtorch::RegisterOperators is legacy. Consider switching to TORCH_LIBRARY_FRAGMENT(trtllm, m) for forward‑compatibility.
cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/detail/collective/mixed_input_utils.hpp (1)
102-129: Initialize bf16x8_raw for parity with the non‑interleaved pathNot functionally required, but zero‑init matches the style used above and avoids any tool false‑positives.
- __nv_bf16x8_storage_t bf16x8_raw; + __nv_bf16x8_storage_t bf16x8_raw = {0, 0};cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hpp (1)
903-921: Use CUDA bit-cast helper to avoid strict‑aliasing UBReplace pointer reinterpret/deref with __uint_as_float (or std::bit_cast if available) to keep the conversion well-defined.
- uint32_t temp = 0; - temp = (temp | *reinterpret_cast<uint8_t*>(&scale_ue8m0)) << 23; - return *reinterpret_cast<float*>(&temp); + uint32_t temp = static_cast<uint32_t>(*reinterpret_cast<uint8_t*>(&scale_ue8m0)) << 23; +#if defined(__CUDA_ARCH__) + return __uint_as_float(temp); +#else + float out; + memcpy(&out, &temp, sizeof(out)); + return out; +#endifNote: include if the fallback path is used on host. As per coding guidelines
cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cu (5)
22-24: Add const/restrict to improve safety and optimization.
weightis read‑only; both pointers can be__restrict__to aid codegen. As per coding guidelines.-__global__ void interleave_fp4_for_Hopper_mixed_gemm_kernel( - uint8_t* weight, uint8_t* weight_interleaved, int const rows, int const cols) +__global__ void interleave_fp4_for_Hopper_mixed_gemm_kernel( + uint8_t const* __restrict__ weight, + uint8_t* __restrict__ weight_interleaved, + int const rows, int const cols)
31-33: Use wider index type for large tensors.
row_id * (cols/2) + col_idcan overflow 32‑bit for very large shapes; prefersize_t. As per coding guidelines.- int index_a = row_id * cols / 2 + col_id; - int index_b = (row_id + 8) * cols / 2 + col_id; + size_t const stride = static_cast<size_t>(colBytes); + size_t const index_a = static_cast<size_t>(row_id) * stride + static_cast<size_t>(col_id); + size_t const index_b = static_cast<size_t>(row_id + 8) * stride + static_cast<size_t>(col_id);
37-42: Replace magic numbers with named constexprs.Use named masks/shifts and tile sizes to improve clarity and maintainability. As per coding guidelines.
+// File-scope (anonymous namespace) or before the kernel: +namespace { +constexpr int kHALF_TILE = 8; +constexpr int kTILE_ROWS = 16; +constexpr int kNIBBLE_SHIFT = 4; +constexpr uint8_t kHI_NIBBLE = 0xF0u; +constexpr uint8_t kLO_NIBBLE = 0x0Fu; +} // anonymous @@ - uint8_t fp4_temp_a = (fp4x2_a & 0xF0U) >> 4; - uint8_t fp4_temp_b = (fp4x2_b & 0x0FU) << 4; + uint8_t const hi_a = static_cast<uint8_t>((fp4x2_a & kHI_NIBBLE) >> kNIBBLE_SHIFT); + uint8_t const lo_b = static_cast<uint8_t>((fp4x2_b & kLO_NIBBLE) << kNIBBLE_SHIFT); @@ - fp4x2_a = (fp4x2_a & 0x0FU) | fp4_temp_b; - fp4x2_b = (fp4x2_b & 0xF0U) | fp4_temp_a; + fp4x2_a = static_cast<uint8_t>((fp4x2_a & kLO_NIBBLE) | lo_b); + fp4x2_b = static_cast<uint8_t>((fp4x2_b & kHI_NIBBLE) | hi_a);
49-53: Avoid hardcoded launch dims; add stream and error check.Compute grid/block from problem size, accept a stream (default 0), and check launch status. As per coding guidelines.
-void interleave_fp4_for_Hopper_mixed_gemm(uint8_t* weight, uint8_t* weight_interleaved, int const rows, int const cols) +void interleave_fp4_for_Hopper_mixed_gemm(uint8_t const* weight, + uint8_t* weight_interleaved, int const rows, int const cols, cudaStream_t stream /* = 0 */) { - // column-major input - interleave_fp4_for_Hopper_mixed_gemm_kernel<<<1024, 1024>>>(weight, weight_interleaved, rows, cols); + // Preconditions match kernel’s expectations; consider host-side validation too. + if ((rows & 0xF) != 0 || (cols & 0x1) != 0) + { + // Option: return or assert/log based on project conventions. + return; + } + + int const rowPairs = rows / 2; + int const colBytes = cols / 2; + + // Heuristic: cap block to colBytes; prefer multiples of warpSize. + int const block_x = colBytes >= 256 ? 256 : (colBytes >= 128 ? 128 : 64); + dim3 block(block_x); + // Each block handles one rowPair in grid-stride; cap grid to device limit as needed. + dim3 grid(std::min(rowPairs, 65535)); + + interleave_fp4_for_Hopper_mixed_gemm_kernel<<<grid, block, 0, stream>>>( + weight, weight_interleaved, rows, cols); + // Error check macro per project convention could replace this. + auto const err = cudaPeekAtLastError(); + if (err != cudaSuccess) + { + // Option: propagate via return code/logging per project standards. + } }
17-18: Optional: include if not guaranteed by the header.If
moe_gemm_mixed_utils.hdoesn’t include<cstdint>, add it here to defineuint8_t. As per coding guidelines.#include "moe_gemm_mixed_utils.h" +// #include <cstdint> // if not included transitively
📜 Review details
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Pro
📒 Files selected for processing (7)
cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/detail/collective/mixed_input_utils.hpp(2 hunks)cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hpp(4 hunks)cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp(1 hunks)cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cu(1 hunks)cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.h(1 hunks)cpp/tensorrt_llm/thop/weightOnlyQuantOp.cpp(3 hunks)tensorrt_llm/_torch/modules/fused_moe/quantization.py(3 hunks)
🧰 Additional context used
📓 Path-based instructions (8)
**/*.{h,hpp,hh,hxx,cpp,cxx,cc,cu,cuh}
📄 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/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.hcpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cucpp/tensorrt_llm/thop/weightOnlyQuantOp.cppcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hppcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/detail/collective/mixed_input_utils.hppcpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp
**/*.{cpp,cxx,cc,cu,h,hpp,hh,hxx,cuh}
📄 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/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.hcpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cucpp/tensorrt_llm/thop/weightOnlyQuantOp.cppcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hppcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/detail/collective/mixed_input_utils.hppcpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp
**/*.{h,hpp,hh,hxx,cpp,cxx,cc,cu,cuh,py}
📄 CodeRabbit inference engine (CODING_GUIDELINES.md)
Use only spaces, no tabs; indent with 4 spaces.
Files:
cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.hcpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cucpp/tensorrt_llm/thop/weightOnlyQuantOp.cpptensorrt_llm/_torch/modules/fused_moe/quantization.pycpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hppcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/detail/collective/mixed_input_utils.hppcpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp
**/*.{h,hpp,hh,hxx}
📄 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/moe_gemm/moe_gemm_mixed_utils.hcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hppcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/detail/collective/mixed_input_utils.hpp
**/*.{h,hpp,hh,hxx,cpp,cxx,cc}
📄 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/moe_gemm/moe_gemm_mixed_utils.hcpp/tensorrt_llm/thop/weightOnlyQuantOp.cppcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hppcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/detail/collective/mixed_input_utils.hppcpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp
**/*.{h,hpp,hh,hxx,cuh}
📄 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/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.hcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hppcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/detail/collective/mixed_input_utils.hpp
**/*.{cpp,cxx,cc,h,hpp,hh,hxx,cu,cuh,py}
📄 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/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.hcpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cucpp/tensorrt_llm/thop/weightOnlyQuantOp.cpptensorrt_llm/_torch/modules/fused_moe/quantization.pycpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hppcpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/detail/collective/mixed_input_utils.hppcpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp
**/*.py
📄 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/modules/fused_moe/quantization.py
🧠 Learnings (2)
📚 Learning: 2025-08-21T21:48:35.135Z
Learnt from: djns99
PR: NVIDIA/TensorRT-LLM#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/thop/weightOnlyQuantOp.cpp
📚 Learning: 2025-08-08T05:10:38.906Z
Learnt from: sklevtsov-nvidia
PR: NVIDIA/TensorRT-LLM#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/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hpp
🧬 Code graph analysis (4)
cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.h (1)
cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cu (2)
interleave_fp4_for_Hopper_mixed_gemm(49-53)interleave_fp4_for_Hopper_mixed_gemm(49-49)
cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cu (1)
cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/detail/collective/mixed_input_utils.hpp (9)
void(238-280)void(285-291)void(294-341)void(346-350)void(353-362)void(366-370)void(373-376)void(379-533)void(536-571)
cpp/tensorrt_llm/thop/weightOnlyQuantOp.cpp (1)
cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cu (2)
interleave_fp4_for_Hopper_mixed_gemm(49-53)interleave_fp4_for_Hopper_mixed_gemm(49-49)
tensorrt_llm/_torch/modules/fused_moe/quantization.py (1)
cpp/tensorrt_llm/thop/weightOnlyQuantOp.cpp (2)
fp4_interleave_for_Hopper_mixed_gemm(402-425)fp4_interleave_for_Hopper_mixed_gemm(402-402)
🪛 Clang (14.0.6)
cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.h
[error] 19-19: 'cstdint' file not found
(clang-diagnostic-error)
⏰ Context from checks skipped due to timeout of 90000ms. You can increase the timeout in your CodeRabbit configuration to a maximum of 15 minutes (900000ms). (1)
- GitHub Check: Pre-commit Check
🔇 Additional comments (5)
cpp/tensorrt_llm/thop/weightOnlyQuantOp.cpp (1)
19-19: Include looks correctThe new kernel utility include is appropriate for exposing the interleave launcher.
cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/detail/collective/mixed_input_utils.hpp (1)
361-362: Switch to interleaved converter LGTMUsing the interleaved converter here aligns with the updated data layout.
cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hpp (1)
1107-1115: Call‑site updates LGTMReplacing static_cast(...) with scale_convertor(...) is consistent.
Also applies to: 1210-1212, 1298-1301
cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.h (1)
1-26: Keep #pragma once; update copyright year and add namespace closing commentThe file is in
kernels/cutlass_kernels/, which follows CUTLASS code conventions (using#pragma oncelike all sibling headers). However, update the copyright year to 2025 and add the namespace closing comment per guidelines.-/* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. All rights reserved. +/* + * Copyright (c) 2020-2025, NVIDIA CORPORATION. All rights reserved. */ #pragma once @@ -24,5 +24,5 @@ namespace tensorrt_llm::kernels::cutlass_kernels void interleave_fp4_for_Hopper_mixed_gemm(uint8_t* weight, uint8_t* weight_interleaved, int const rows, int const cols); -} +} // namespace tensorrt_llm::kernels::cutlass_kernelscpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cu (1)
51-53: Verify layout note vs indexing.Comment says “column‑major input” but indices use row-major stride (
row_id * (cols/2) + col_id). Confirm intended layout or fix the comment/indexing.
| if (!has_coop_supported) | ||
| continue; | ||
| // Due to the limitation on the number of registers on SM, | ||
| // cooperative scheduler does not support CtaShape128x128x128B. | ||
| if (tile_config == CutlassTileConfigSM90::CtaShape128x128x128B) | ||
| continue; |
There was a problem hiding this comment.
🛠️ Refactor suggestion | 🟠 Major
Add braces around single-statement if bodies.
The coding guidelines require that if statements always be followed by brace-delimited statements. Both continue statements lack the required braces.
As per coding guidelines.
Apply this diff to add braces:
- if (!has_coop_supported)
- continue;
+ if (!has_coop_supported)
+ {
+ continue;
+ }
// Due to the limitation on the number of registers on SM,
// cooperative scheduler does not support CtaShape128x128x128B.
- if (tile_config == CutlassTileConfigSM90::CtaShape128x128x128B)
- continue;
+ if (tile_config == CutlassTileConfigSM90::CtaShape128x128x128B)
+ {
+ continue;
+ }📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| if (!has_coop_supported) | |
| continue; | |
| // Due to the limitation on the number of registers on SM, | |
| // cooperative scheduler does not support CtaShape128x128x128B. | |
| if (tile_config == CutlassTileConfigSM90::CtaShape128x128x128B) | |
| continue; | |
| if (!has_coop_supported) | |
| { | |
| continue; | |
| } | |
| // Due to the limitation on the number of registers on SM, | |
| // cooperative scheduler does not support CtaShape128x128x128B. | |
| if (tile_config == CutlassTileConfigSM90::CtaShape128x128x128B) | |
| { | |
| continue; | |
| } |
🤖 Prompt for AI Agents
In cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp around lines
310 to 315, two if statements use single-statement bodies with continue and must
be converted to brace-delimited blocks per the coding guidelines; change each
`if (condition) continue;` to `if (condition) { continue; }`, preserving
existing indentation and spacing.
| /* | ||
| * Copyright (c) 2020-2023, NVIDIA CORPORATION. All rights reserved. | ||
| * | ||
| * Licensed under the Apache License, Version 2.0 (the "License"); | ||
| * you may not use this file except in compliance with the License. | ||
| * You may obtain a copy of the License at | ||
| * | ||
| * http://www.apache.org/licenses/LICENSE-2.0 | ||
| * | ||
| * Unless required by applicable law or agreed to in writing, software | ||
| * distributed under the License is distributed on an "AS IS" BASIS, | ||
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| * See the License for the specific language governing permissions and | ||
| * limitations under the License. | ||
| */ |
There was a problem hiding this comment.
Update copyright year to include 2025.
Header stops at 2023; repository guidelines require current year on source files. As per coding guidelines.
- * Copyright (c) 2020-2023, NVIDIA CORPORATION. All rights reserved.
+ * Copyright (c) 2020-2025, NVIDIA CORPORATION. All rights reserved.📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| /* | |
| * Copyright (c) 2020-2023, NVIDIA CORPORATION. All rights reserved. | |
| * | |
| * Licensed under the Apache License, Version 2.0 (the "License"); | |
| * you may not use this file except in compliance with the License. | |
| * You may obtain a copy of the License at | |
| * | |
| * http://www.apache.org/licenses/LICENSE-2.0 | |
| * | |
| * Unless required by applicable law or agreed to in writing, software | |
| * distributed under the License is distributed on an "AS IS" BASIS, | |
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | |
| * See the License for the specific language governing permissions and | |
| * limitations under the License. | |
| */ | |
| /* | |
| * Copyright (c) 2020-2025, NVIDIA CORPORATION. All rights reserved. | |
| * | |
| * Licensed under the Apache License, Version 2.0 (the "License"); | |
| * you may not use this file except in compliance with the License. | |
| * You may obtain a copy of the License at | |
| * | |
| * http://www.apache.org/licenses/LICENSE-2.0 | |
| * | |
| * Unless required by applicable law or agreed to in writing, software | |
| * distributed under the License is distributed on an "AS IS" BASIS, | |
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | |
| * See the License for the specific language governing permissions and | |
| * limitations under the License. | |
| */ |
🤖 Prompt for AI Agents
In cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cu
around lines 1 to 15, the copyright header ends with "2023" but repository
guidelines require the current year; update the copyright range to include 2025
(e.g., "2020-2025" or "2020-2023, 2025" per project convention) and ensure the
license block formatting is preserved exactly as before.
| for (int block_id = blockIdx.x; block_id < rows / 2; block_id += gridDim.x) | ||
| { | ||
| for (int col_id = threadIdx.x; col_id < cols / 2; col_id += blockDim.x) | ||
| { | ||
| int row_id = block_id / 8 * 16 + block_id % 8; | ||
|
|
||
| int index_a = row_id * cols / 2 + col_id; | ||
| int index_b = (row_id + 8) * cols / 2 + col_id; | ||
|
|
There was a problem hiding this comment.
Enforce required shape preconditions and avoid repeated divisions.
Kernel assumes rows % 16 == 0 and cols % 2 == 0. Without checks, odd sizes can index OOB. Also factor rowPairs/colBytes once. As per coding guidelines.
{
- for (int block_id = blockIdx.x; block_id < rows / 2; block_id += gridDim.x)
+ // Preconditions: rows are multiples of 16; cols are even (#fp4 is even -> bytes).
+ if ((rows & 0xF) != 0 || (cols & 0x1) != 0)
+ {
+ return;
+ }
+
+ int const rowPairs = rows / 2;
+ int const colBytes = cols / 2;
+
+ for (int block_id = blockIdx.x; block_id < rowPairs; block_id += gridDim.x)
{
- for (int col_id = threadIdx.x; col_id < cols / 2; col_id += blockDim.x)
+ for (int col_id = threadIdx.x; col_id < colBytes; col_id += blockDim.x)
{
- int row_id = block_id / 8 * 16 + block_id % 8;
+ int const row_id = (block_id / 8) * 16 + (block_id % 8);🤖 Prompt for AI Agents
In cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cu
around lines 25-33, add explicit precondition checks that rows % 16 == 0 and
cols % 2 == 0 (return or assert/error if violated) to prevent OOB indexing, and
refactor the repeated divisions by computing locals once (e.g., halfRows =
rows/2 and halfCols = cols/2 or rowPairs = rows/16 and colBytes = cols/2) and
use those locals in the loops and index calculations instead of recalculating
rows/2 and cols/2 repeatedly.
| if w3_weight_shard.ndim == 2: | ||
| w31_weight_shard = torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm( | ||
| w31_weight_shard) | ||
|
|
||
| dst_w3_w1_weight.copy_(w31_weight_shard.view(dst_w3_w1_weight.dtype), | ||
| non_blocking=True) |
There was a problem hiding this comment.
🧩 Analysis chain
Guard CUDA-only interleave and preserve original device
fp4_interleave_for_Hopper_mixed_gemm requires CUDA tensors (C++ side checks). Today dst_w3_w1_weight can be CPU. Move to CUDA for the op, then back, and assert 16-aligned shapes to avoid C++ TORCH_CHECK.
- if w3_weight_shard.ndim == 2:
- w31_weight_shard = torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm(
- w31_weight_shard)
+ if w3_weight_shard.ndim == 2:
+ n, half_k = w31_weight_shard.shape
+ # C++ op requires n % 16 == 0 and (2*half_k) % 16 == 0
+ assert (n % 16 == 0) and ((2 * half_k) % 16 == 0), "FP4 interleave requires 16-aligned shapes"
+ src_dev = w31_weight_shard.device
+ if src_dev.type != "cuda":
+ w31_weight_shard = w31_weight_shard.cuda(non_blocking=True)
+ w31_weight_shard = torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm(w31_weight_shard)
+ w31_weight_shard = w31_weight_shard.to(src_dev, non_blocking=True)
+ else:
+ w31_weight_shard = torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm(w31_weight_shard)Please confirm the module’s parameters are on CUDA at load time for this method; otherwise this fix is required to prevent runtime TORCH_CHECK failures from the C++ op.
I'll verify this review comment by examining the code context, parameter initialization, and device handling. Let me search the codebase for relevant information.
---
🏁 Script executed:
```shell
#!/bin/bash
# Examine context around the target lines and the broader method
cat -n tensorrt_llm/_torch/modules/fused_moe/quantization.py | sed -n '1380,1430p'
Length of output: 2811
🏁 Script executed:
#!/bin/bash
# Search for dst_w3_w1_weight initialization and device logic
rg -n -B5 -A5 'dst_w3_w1_weight' tensorrt_llm/_torch/modules/fused_moe/quantization.pyLength of output: 31254
🏁 Script executed:
#!/bin/bash
# Find the method definition containing this code
rg -n -B20 'fp4_interleave_for_Hopper_mixed_gemm' tensorrt_llm/_torch/modules/fused_moe/quantization.py | head -50Length of output: 2226
🏁 Script executed:
#!/bin/bash
# Search for how this method is called and what device assumptions exist
rg -n 'def.*quantize' tensorrt_llm/_torch/modules/fused_moe/quantization.pyLength of output: 71
🏁 Script executed:
#!/bin/bash
# Look for CUDA-related assertions or device checks in similar operations
rg -n 'cuda\(\|device\|\.to\(' tensorrt_llm/_torch/modules/fused_moe/quantization.py | head -30Length of output: 0
---
🏁 Script executed:
```shell
#!/bin/bash
# Search for where dst_w3_w1_weights_tensor is allocated/created
rg -n 'dst_w3_w1_weight' tensorrt_llm/_torch/modules/fused_moe/quantization.py | grep -E '(cuda|\.to\(|device|empty|zeros|ones)' | head -20
Length of output: 636
🏁 Script executed:
#!/bin/bash
# Find class definition to see where tensor allocation happens
rg -n 'class.*MoE.*:' tensorrt_llm/_torch/modules/fused_moe/quantization.pyLength of output: 1754
🏁 Script executed:
#!/bin/bash
# Check if there's an __init__ or allocation method that sets up these tensors
rg -n -A10 'def.*__init__' tensorrt_llm/_torch/modules/fused_moe/quantization.py | head -40Length of output: 0
🏁 Script executed:
#!/bin/bash
# Search for other similar ops to see if they have device handling
rg -n 'torch.ops.trtllm' tensorrt_llm/_torch/modules/fused_moe/quantization.py | head -20Length of output: 1498
🏁 Script executed:
#!/bin/bash
# Look for where load_expert_w3_w1_weight is called from
rg -n 'load_expert_w3_w1_weight' tensorrt_llm/_torch/modules/fused_moe/quantization.py | head -10Length of output: 623
🏁 Script executed:
#!/bin/bash
# Examine another class's load_expert_w3_w1_weight that has cuda assert (line 1887)
sed -n '1883,1925p' tensorrt_llm/_torch/modules/fused_moe/quantization.pyLength of output: 2085
---
🏁 Script executed:
```shell
#!/bin/bash
# Find where WFP4A16FusedMoEMethod is instantiated/used
rg -n 'WFP4A16FusedMoEMethod' tensorrt_llm/ -r 'FOUND: $0'
Length of output: 364
🏁 Script executed:
#!/bin/bash
# Look at the load_expert_weights_to_dst method that calls load_expert_w3_w1_weight
sed -n '210,270p' tensorrt_llm/_torch/modules/fused_moe/quantization.pyLength of output: 3091
🏁 Script executed:
#!/bin/bash
# Check how torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm is defined/documented
find tensorrt_llm -name "*.py" -o -name "*.cu" -o -name "*.cpp" | xargs grep -l 'fp4_interleave_for_Hopper_mixed_gemm' 2>/dev/null | head -5Length of output: 53
🏁 Script executed:
#!/bin/bash
# Compare all torch.ops.trtllm usages and check which ones have cuda assertions before them
rg -n -B5 'torch.ops.trtllm.(shuffle_matrix|fp4_interleave|block_scale_interleave)' tensorrt_llm/_torch/modules/fused_moe/quantization.py | grep -E '(assert|cuda|device|torch.ops.trtllm)' | head -40Length of output: 1583
---
🏁 Script executed:
```shell
#!/bin/bash
# Look at fused_moe_cutlass.py to see context of WFP4A16FusedMoEMethod instantiation
sed -n '210,225p' tensorrt_llm/_torch/modules/fused_moe/fused_moe_cutlass.py
Length of output: 758
🏁 Script executed:
#!/bin/bash
# Search for where module.w3_w1_weight tensor is allocated
rg -n 'w3_w1_weight\s*=' tensorrt_llm/_torch/modules/fused_moe/ | head -20Length of output: 577
🏁 Script executed:
#!/bin/bash
# Search for module setup or initialization that might set device
rg -n 'def.*__init__\|w3_w1_weight.*torch\.' tensorrt_llm/_torch/modules/fused_moe/ | head -20Length of output: 0
🏁 Script executed:
#!/bin/bash
# Check if there are any comments or checks about device in the FusedMoE module
rg -n 'cpu\|CUDA\|device.*cuda' tensorrt_llm/_torch/modules/fused_moe/ | head -15Length of output: 0
---
🏁 Script executed:
```shell
#!/bin/bash
# Look at line 178 context where w3_w1_weight is allocated
sed -n '170,195p' tensorrt_llm/_torch/modules/fused_moe/quantization.py
Length of output: 1222
🏁 Script executed:
#!/bin/bash
# Search for the function/method that contains line 178
sed -n '153,185p' tensorrt_llm/_torch/modules/fused_moe/quantization.pyLength of output: 1234
🏁 Script executed:
#!/bin/bash
# Check what arguments are passed to torch.empty
rg -n -A2 'torch.empty.*w3_w1_weight_shape' tensorrt_llm/_torch/modules/fused_moe/quantization.pyLength of output: 205
Guard CUDA-only interleave and preserve original device
The review comment correctly identifies an issue. The fp4_interleave_for_Hopper_mixed_gemm C++ operation requires CUDA tensors, but the code at lines 1409–1411 does not guard this requirement. While torch.empty creates the w3_w1_weight parameter without specifying a device, allowing it to be allocated on CPU, a code comment explicitly notes that tensors don't have to be CUDA, e.g., for "online EPLB".
Other similar methods in the same file (e.g., NVFP4TRTLLMGenFusedMoEMethod at line 1888) include assert device.type == "cuda" before calling torch.ops.trtllm operations, showing the pattern that should be followed here. The proposed fix—moving tensors to CUDA for the operation, executing it, and moving back while asserting 16-aligned shapes—is appropriate.
| if w2_weight_shard.ndim == 2: | ||
| w2_weight_shard = torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm( | ||
| w2_weight_shard) | ||
|
|
There was a problem hiding this comment.
Same CUDA guard for w2 path
Mirror the device/shape guard to avoid failures on CPU tensors.
- if w2_weight_shard.ndim == 2:
- w2_weight_shard = torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm(
- w2_weight_shard)
+ if w2_weight_shard.ndim == 2:
+ n, half_k = w2_weight_shard.shape
+ assert (n % 16 == 0) and ((2 * half_k) % 16 == 0), "FP4 interleave requires 16-aligned shapes"
+ src_dev = w2_weight_shard.device
+ if src_dev.type != "cuda":
+ w2_weight_shard = w2_weight_shard.cuda(non_blocking=True)
+ w2_weight_shard = torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm(w2_weight_shard)
+ w2_weight_shard = w2_weight_shard.to(src_dev, non_blocking=True)
+ else:
+ w2_weight_shard = torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm(w2_weight_shard)📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| if w2_weight_shard.ndim == 2: | |
| w2_weight_shard = torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm( | |
| w2_weight_shard) | |
| if w2_weight_shard.ndim == 2: | |
| n, half_k = w2_weight_shard.shape | |
| assert (n % 16 == 0) and ((2 * half_k) % 16 == 0), "FP4 interleave requires 16-aligned shapes" | |
| src_dev = w2_weight_shard.device | |
| if src_dev.type != "cuda": | |
| w2_weight_shard = w2_weight_shard.cuda(non_blocking=True) | |
| w2_weight_shard = torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm(w2_weight_shard) | |
| w2_weight_shard = w2_weight_shard.to(src_dev, non_blocking=True) | |
| else: | |
| w2_weight_shard = torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm(w2_weight_shard) |
🤖 Prompt for AI Agents
tensorrt_llm/_torch/modules/fused_moe/quantization.py around lines 1443 to 1446:
the interleave call for w2_weight_shard lacks the same device/shape guard as the
w1 path, which can cause failures on CPU tensors; wrap the call so it only runs
when w2_weight_shard is a CUDA tensor and has ndim == 2 (mirror the existing
guard used for w1), i.e., check device/is_cuda and ndim before invoking
torch.ops.trtllm.fp4_interleave_for_Hopper_mixed_gemm so CPU or non-2D tensors
skip the CUDA-specific op.
| } | ||
|
|
||
| // It seems that ping-pong scheduler will never be selected. | ||
| // To shorten the tactic time, remove all alternative options involving ping-pong scheduler. |
There was a problem hiding this comment.
What models have you tested this with? I am hesitant to remove this without a comprehensive sweep of multiple model architectures like Mixtral, DeepSeek, Llama4 and GPT-OSS. Its hard to say what the next DeepSeek moment will look like.
I also dont think tactic selection time is actually a significant concern. There are lots of tactics sure, but weight loading is usually just as long. Maybe we should add a fast profile mode that users can opt into
| // Due to the limitation on the number of registers on SM, | ||
| // cooperative scheduler does not support CtaShape128x128x128B. | ||
| if (tile_config == CutlassTileConfigSM90::CtaShape128x128x128B) | ||
| continue; |
There was a problem hiding this comment.
How much performance are we leaving on the table here? Is there a way to reduce the number of stages or otherwise relieve register pressure
2 participants
Summary by CodeRabbit
Optimization
Refactor
Description
Test Coverage
PR Checklist
Please review the following before submitting your PR:
PR description clearly explains what and why. If using CodeRabbit's summary, please make sure it makes sense.
PR Follows TRT-LLM CODING GUIDELINES to the best of your knowledge.
Test cases are provided for new code paths (see test instructions)
Any new dependencies have been scanned for license and vulnerabilities
CODEOWNERS updated if ownership changes
Documentation updated as needed
The reviewers assigned automatically/manually are appropriate for the PR.
Please check this after reviewing the above items as appropriate for this PR.
GitHub Bot Help
/bot [-h] ['run', 'kill', 'skip', 'reuse-pipeline'] ...Provide a user friendly way for developers to interact with a Jenkins server.
Run
/bot [-h|--help]to print this help message.See details below for each supported subcommand.
run [--reuse-test (optional)pipeline-id --disable-fail-fast --skip-test --stage-list "A10-PyTorch-1, xxx" --gpu-type "A30, H100_PCIe" --test-backend "pytorch, cpp" --add-multi-gpu-test --only-multi-gpu-test --disable-multi-gpu-test --post-merge --extra-stage "H100_PCIe-TensorRT-Post-Merge-1, xxx" --detailed-log --debug(experimental)]Launch build/test pipelines. All previously running jobs will be killed.
--reuse-test (optional)pipeline-id(OPTIONAL) : Allow the new pipeline to reuse build artifacts and skip successful test stages from a specified pipeline or the last pipeline if no pipeline-id is indicated. If the Git commit ID has changed, this option will be always ignored. The DEFAULT behavior of the bot is to reuse build artifacts and successful test results from the last pipeline.--disable-reuse-test(OPTIONAL) : Explicitly prevent the pipeline from reusing build artifacts and skipping successful test stages from a previous pipeline. Ensure that all builds and tests are run regardless of previous successes.--disable-fail-fast(OPTIONAL) : Disable fail fast on build/tests/infra failures.--skip-test(OPTIONAL) : Skip all test stages, but still run build stages, package stages and sanity check stages. Note: Does NOT update GitHub check status.--stage-list "A10-PyTorch-1, xxx"(OPTIONAL) : Only run the specified test stages. Examples: "A10-PyTorch-1, xxx". Note: Does NOT update GitHub check status.--gpu-type "A30, H100_PCIe"(OPTIONAL) : Only run the test stages on the specified GPU types. Examples: "A30, H100_PCIe". Note: Does NOT update GitHub check status.--test-backend "pytorch, cpp"(OPTIONAL) : Skip test stages which don't match the specified backends. Only support [pytorch, cpp, tensorrt, triton]. Examples: "pytorch, cpp" (does not run test stages with tensorrt or triton backend). Note: Does NOT update GitHub pipeline status.--only-multi-gpu-test(OPTIONAL) : Only run the multi-GPU tests. Note: Does NOT update GitHub check status.--disable-multi-gpu-test(OPTIONAL) : Disable the multi-GPU tests. Note: Does NOT update GitHub check status.--add-multi-gpu-test(OPTIONAL) : Force run the multi-GPU tests in addition to running L0 pre-merge pipeline.--post-merge(OPTIONAL) : Run the L0 post-merge pipeline instead of the ordinary L0 pre-merge pipeline.--extra-stage "H100_PCIe-TensorRT-Post-Merge-1, xxx"(OPTIONAL) : Run the ordinary L0 pre-merge pipeline and specified test stages. Examples: --extra-stage "H100_PCIe-TensorRT-Post-Merge-1, xxx".--detailed-log(OPTIONAL) : Enable flushing out all logs to the Jenkins console. This will significantly increase the log volume and may slow down the job.--debug(OPTIONAL) : Experimental feature. Enable access to the CI container for debugging purpose. Note: Specify exactly one stage in thestage-listparameter to access the appropriate container environment. Note: Does NOT update GitHub check status.For guidance on mapping tests to stage names, see
docs/source/reference/ci-overview.mdand the
scripts/test_to_stage_mapping.pyhelper.kill
killKill all running builds associated with pull request.
skip
skip --comment COMMENTSkip testing for latest commit on pull request.
--comment "Reason for skipping build/test"is required. IMPORTANT NOTE: This is dangerous since lack of user care and validation can cause top of tree to break.reuse-pipeline
reuse-pipelineReuse a previous pipeline to validate current commit. This action will also kill all currently running builds associated with the pull request. IMPORTANT NOTE: This is dangerous since lack of user care and validation can cause top of tree to break.