[None][feat] TRT-LLM Gen MoE optimize DeepSeek Fp8 activation kernel #9175
Conversation
Signed-off-by: Nikita Korobov <14355239+nekorobov@users.noreply.github.com>
📝 WalkthroughWalkthroughThe PR introduces vectorized processing utilities (Float4Max, Float2Max functors, type-packing functions) and a KernelTraits template for adaptive operations based on token packing. The activation kernel is extended to process multiple tokens per CTA using packed types with per-token shared arrays, local work arrays, and block-level reductions. Host-side scheduling logic is updated with heuristics to determine optimal grid and thread configurations for the DeepSeek FP8 path. Changes
Estimated code review effort🎯 4 (Complex) | ⏱️ ~45–60 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: 0
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu (1)
223-331: Restructure the synchronization pattern or enforce strong innerDim divisibility constraints before kernel launch.The deadlock risk is real and verified. The
activationDeepSeekKernelhas a critical synchronization bug:Loop configuration:
NUM_THREADS_PER_CTA = 128(line 369, DevKernel.cu)gridSizeX = ceil((params.innerDim / 2) / 128)(lines 377–379)- Thread stride:
128 * gridSizeXDeadlock trigger:
Whenparams.innerDim / 2is not an exact multiple of128, the last CTA (whereblockIdx.x = gridSizeX - 1) has threads that skip the stridedhiddenIdxloop entirely. Specifically, any threadtwheret + 128 * (gridSizeX - 1) >= params.innerDim / 2never enters the loop and thus never reachesBlockReduce::Reduce()or__syncthreads()at line 329. Meanwhile, other threads in the same block do reach these calls, violating the requirement that all threads participate uniformly. This causes warp divergence and block deadlock.Current state:
- No validation of
innerDimdivisibility exists before launch (runner.cu, lines 402–428)innerDimis set toargs.intermediate_size * 2with no constraint checksRequired fix (choose one):
- Enforce at host level (line 403–428 in runner.cu): Add a
TLLM_CHECKthatdata.innerDim / 2is divisible by128before callingmoe::dev::activation::run(activationData, stream)at line 537.- Restructure kernel (lines 260–348 in DevKernel.cu): Accumulate per-thread maxima over the strided loop in local arrays, then perform a single
BlockReduceper(tokenCtaIdx, k)outside the loop, followed by__syncthreads()once.
🧹 Nitpick comments (5)
cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.h (1)
85-122: Token-per-CTA launch macros look consistent; consider tightening error handling.The new
LAUNCH_NUM_TOKENS_PER_CTA/LAUNCH_ACTIVATIONmacros correctly threadNumTokensPerCtaintoKernelParamsviaLAUNCH_ESC(type, 4/2/1), and the dtype branching mirrors the existingLAUNCHpattern.One thing to consider is that the
elsebranch for unsupportednumTokensPerCtaonly logs and silently skips the launch. If this path is ever hit, a hard failure (TLLM_CHECK_WITH_INFOor similar) would make debugging much easier than a quiet no-op.cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu (4)
103-216: Vector max and pack/unpack helpers look good; consider guarding unsupported sizes at compile time.The
Float4Max/Float2Maxfunctors and thepackedTypeFromArray/arrayFromPackedTypespecializations for(float4,4),(float2,2), and(float,1)are straightforward and match theKernelTraits<4/2/1>usage.The generic templates currently return zero-initialized values, which silently “works” if an unsupported
(VecType, size)combination is ever instantiated, but hides the bug. Consider adding astatic_assertin the generic versions to fail compilation for unsupported combinations so future refactors can’t accidentally pick an invalidNumTokensPerCta/PackedTypepair without also updating these helpers.
246-305: Avoid reading uninitialized per-token arrays whenpermutedIdx == -1or token index is out of range.The per-token arrays
float scale1Arr[NumTokensPerCta]; float scale2Arr[NumTokensPerCta]; float dataX1Arr[NumTokensPerCta]; float dataX2Arr[NumTokensPerCta]; float outArr[NumTokensPerCta]; float absOutArr[NumTokensPerCta]; int permutedIdxArr[NumTokensPerCta];are written in the first
tokenInCtaIdxloop only when:
tokenIdx < params.numTokens, andpermutedIdx != -1.In the second loop, you unconditionally use
scale1Arr,scale2Arr,dataX*Arr:for (int tokenInCtaIdx = 0; tokenInCtaIdx < NumTokensPerCta; tokenInCtaIdx++) { float x1 = scale1Arr[tokenInCtaIdx] * dataX1Arr[tokenInCtaIdx]; float x2 = scale2Arr[tokenInCtaIdx] * dataX2Arr[tokenInCtaIdx]; ... }For tokens where
permutedIdx == -1ortokenIdx >= params.numTokens, those array entries are never initialized, so this loop reads uninitialized data. Later loops gate ontokenIdxandpermutedIdxand avoid using the corresponding results, so the unintended values likely don’t affect outputs, but from a C++/CUDA perspective this is still undefined behavior and can lead to unpredictable values feeding into the block-wide reduction.A low-cost fix is to explicitly zero-initialize these arrays (or the relevant entries) when a token is invalid, e.g.:
- float scale1Arr[NumTokensPerCta]; + float scale1Arr[NumTokensPerCta] = {}; ... - int permutedIdx = params.expandedIdxToPermutedIdx[expandedIdx]; - permutedIdxArr[tokenInCtaIdx] = permutedIdx; - if (permutedIdx == -1) - { - continue; - } + int permutedIdx = params.expandedIdxToPermutedIdx[expandedIdx]; + permutedIdxArr[tokenInCtaIdx] = permutedIdx; + if (permutedIdx == -1) + { + // Leave scale/data arrays at zero for this lane. + continue; + }and similarly zero-initialize
dataX*Arr,outArr, andabsOutArrto keep the reduction well-defined for all lanes.
223-231: Coupling between BlockReduce block size and launch configuration should be documented or asserted.
activationDeepSeekKernelinstantiates:using BlockReduce = cub::BlockReduce<PackedType, 128>; __shared__ typename BlockReduce::TempStorage tempStorage;and
run()hardcodes:int const NUM_THREADS_PER_CTA = 128; ... LAUNCH_ACTIVATION(data, activationDeepSeekKernel, numTokensPerCta, grid, NUM_THREADS_PER_CTA, 0, stream);So the
BlockReducetemplate parameter and the actualblockDim.xare implicitly tied through this constant. That’s fine now, but it’s brittle: a future tweak toNUM_THREADS_PER_CTA(or a different call site) would silently break the reduction.Consider adding a small runtime
assert(blockDim.x == 128);at the start of the kernel (debug-only) or a comment inrun()stating thatBlockReduceassumesblockDim.x == 128, so changes must be synchronized.Also applies to: 365-401
365-401: DeepSeek activation scheduling heuristic and LAUNCH_ACTIVATION usage look reasonable; just a small nit.The heuristic choice of
numTokensPerCtabased onnumCtasvs.numSmsand the resultinggridSizeYandgridcomputation are clear and bounded, and usingLAUNCH_ACTIVATIONfor both DeepSeek and non-DeepSeek paths keeps the dtype/NumTokensPerCta handling centralized.Only minor nit: the comment
// FIXME: This is heruistic based on very short benchmark.has a typo (“heruistic”). If this remains a long-term heuristic, consider either removing the FIXME or adding a brief note on which workloads it was tuned for, but that’s purely cosmetic.Also applies to: 408-409
📜 Review details
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Pro
📒 Files selected for processing (2)
cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu(3 hunks)cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.h(2 hunks)
🧰 Additional context used
🧠 Learnings (15)
📓 Common learnings
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.
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.
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.
📚 Learning: 2025-09-23T14:58:05.372Z
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/DevKernel.hcpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu
📚 Learning: 2025-08-14T21:04:50.248Z
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/DevKernel.h
📚 Learning: 2025-11-14T11:22:03.711Z
Learnt from: nzmora-nvidia
Repo: NVIDIA/TensorRT-LLM PR: 9163
File: tensorrt_llm/_torch/auto_deploy/custom_ops/quant.py:107-113
Timestamp: 2025-11-14T11:22:03.711Z
Learning: In TensorRT-LLM AutoDeploy custom ops, when adding hardware capability checks to select between kernel implementations (e.g., cuBLAS vs. CUDA kernel), use descriptive variable names that identify the specific GPU architectures or families being targeted (e.g., `is_blackwell_geforce_or_ada`) rather than generic names like `enable_cuda_core`. This makes it clear that the code is selecting an implementation path based on hardware capabilities, not enabling/disabling hardware features.
Applied to files:
cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.hcpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu
📚 Learning: 2025-08-15T06:46:54.897Z
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/DevKernel.hcpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu
📚 Learning: 2025-09-23T15:01:00.070Z
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/DevKernel.hcpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu
📚 Learning: 2025-08-21T02:39:12.009Z
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/DevKernel.hcpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu
📚 Learning: 2025-08-09T20:57:04.084Z
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/DevKernel.h
📚 Learning: 2025-09-19T21:28:13.751Z
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/DevKernel.hcpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu
📚 Learning: 2025-08-08T22:03:40.707Z
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/DevKernel.hcpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu
📚 Learning: 2025-09-02T13:42:44.885Z
Learnt from: pcastonguay
Repo: NVIDIA/TensorRT-LLM PR: 7455
File: tensorrt_llm/_torch/pyexecutor/py_executor.py:1852-1860
Timestamp: 2025-09-02T13:42:44.885Z
Learning: In MPI communication within TensorRT-LLM pipeline parallelism, different communication types (tokens, logits, termination sync) must use disjoint tag namespaces to avoid message routing collisions when using the same source/destination patterns.
Applied to files:
cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.h
📚 Learning: 2025-08-22T01:54:35.850Z
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/trtllmGenKernels/blockScaleMoe/DevKernel.cu
📚 Learning: 2025-08-19T03:35:20.866Z
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/DevKernel.cu
📚 Learning: 2025-09-23T15:13:48.819Z
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/trtllmGenKernels/blockScaleMoe/DevKernel.cu
📚 Learning: 2025-08-14T15:36:37.610Z
Learnt from: MatthiasKohl
Repo: NVIDIA/TensorRT-LLM PR: 6904
File: cpp/tensorrt_llm/kernels/mlaKernels.cu:436-439
Timestamp: 2025-08-14T15:36:37.610Z
Learning: CUDA kernels prioritize performance and should avoid runtime bounds checking or conditional operations that cause branching/warp divergence. Input validation should be done at the host level before kernel launch, not per-thread in the kernel.
Applied to files:
cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu
⏰ 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 (1)
cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.h (1)
280-285: Verified: macro expansion ensures compatibility with three-parameter struct signature.The two-parameter instantiations at lines 52-61 are not actually incompatible. When
typesis passed throughLAUNCH_ESC(type, 4)(or similar variants), the preprocessor defers its expansion until the template instantiation site. DuringKernelParams<types, true>expansion,LAUNCH_ESCre-expands, yieldingKernelParams<type, 4, true>— providing the required three parameters.All verified call paths through
LAUNCH_NUM_TOKENS_PER_CTA(lines 88, 92, 96) correctly supply two tokens viaLAUNCH_ESC, which combined with the implicit third parameter (true/false), match the new signature. No direct instantiations ofactivation::KernelParamswith the old two-parameter form remain in the codebase.
|
/bot run --disable-fail-fast |
|
PR_Github #24595 [ run ] triggered by Bot. Commit: |
rosenrodt
left a comment
rosenrodt
left a comment
There was a problem hiding this comment.
Thanks! I have left a quesiton inline around perf considerations. Otherwise looks good to me. Tagging @ChristinaZ for insights for past work on finalize kernel optimization, as activation kernel seems similar to finalize kernel
|
PR_Github #24595 [ run ] completed with state |
cpp/tensorrt_llm/kernels/trtllmGenKernels/blockScaleMoe/DevKernel.cu
Outdated
Show resolved
Hide resolved
Signed-off-by: Nikita Korobov <14355239+nekorobov@users.noreply.github.com>
|
/bot run |
|
PR_Github #25358 [ run ] triggered by Bot. Commit: |
|
PR_Github #25358 [ run ] completed with state |
4 participants
Summary by CodeRabbit
New Features
Performance
Description
Similar to flashinfer-ai/flashinfer#2063
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
Update tava architecture diagram if there is a significant design change in PR.
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.