[TRTLLM-7963][feat] Cold L2 cache when doing autotune benchmarking.

[hyukn]

When running ops in models, the L2 cache is usually cold. But in autotune, we don't clear L2 cache when benchmarking. The perf results of some kernels could be affected by the warm/cold L2 cache easily. To achieve more precise profiling results, we clear the L2 cache with the circular buffer method for better benchmarking during AutoTuner profiling (almost the same as #7587).

Summary by CodeRabbit

• New Features

  • Added use_cold_l2_cache configuration option to control multi-batch profiling behavior during autotuning.
  • Enabled batch-based profiling for more comprehensive kernel performance analysis.

• Improvements

  • Increased profiling iterations for more thorough autotuning results.
  • Enhanced CUDA error handling and reporting during profiling.

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.

Details

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 the stage-list parameter 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.md
and the scripts/test_to_stage_mapping.py helper.

kill

kill

Kill all running builds associated with pull request.

skip

skip --comment COMMENT

Skip 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-pipeline

Reuse 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.

[hyukn]

/bot run --disable-fail-fast

[coderabbitai]

📝 Walkthrough

Walkthrough

Extended the autotuner with CUDA driver integration for L2 cache operations, added use_cold_l2_cache configuration flag, increased profiling repeat count from 10 to 30, implemented multi-batch input preparation for profiling, and propagated tuning configuration through profiling pipelines.

Changes

Cohort / File(s)

Change Summary

CUDA Integration and Batch-based Profiling tensorrt_llm/_torch/autotuner.py

Added CUDA driver bindings import and error handling utilities (_checkCudaErrors, _cudaGetErrorEnum, _get_l2_cache_size_in_bytes). Extended TuningConfig with new use_cold_l2_cache: bool = True field. Increased AutoTuner.__init__ repeat parameter from 10 to 30. Introduced _prepare_input_tensors_with_batches method to generate multiple input batches for L2 cache warming during profiling. Updated _profile_runners and _profile_single_kernel to accept and propagate tuning_config. Modified profiling logic to iterate across batches during timing with warmup on final batch. Elevated certain profiling logging from debug to info level.

Sequence Diagram

sequenceDiagram
    participant Caller
    participant AutoTuner
    participant CUDA as CUDA Driver
    participant Profiler as Profiling Engine

    Caller->>AutoTuner: profile_runners(..., tuning_config)
    
    AutoTuner->>AutoTuner: _prepare_input_tensors_with_batches(inputs, tuning_config)
    
    alt use_cold_l2_cache enabled
        AutoTuner->>CUDA: _get_l2_cache_size_in_bytes()
        CUDA-->>AutoTuner: l2_cache_size
        AutoTuner->>AutoTuner: compute batch_count from L2 size
        AutoTuner-->>AutoTuner: return batched_inputs (List[List[Tensor]])
    else use_cold_l2_cache disabled
        AutoTuner-->>AutoTuner: return [inputs] (single batch)
    end
    
    loop for each tactic
        AutoTuner->>AutoTuner: _profile_single_kernel(..., tuning_config)
        
        loop batches (repeat=30)
            loop for each input batch
                Profiler->>Profiler: warm-up with last batch
                Profiler->>Profiler: measure kernel execution
            end
        end
        
        Profiler-->>AutoTuner: profiling_time
    end
    
    AutoTuner-->>Caller: best_tactic

Loading

Estimated code review effort

🎯 3 (Moderate) | ⏱️ ~25 minutes

• CUDA error handling implementation: Review correctness of _checkCudaErrors, _cudaGetErrorEnum, and _get_l2_cache_size_in_bytes for proper error propagation and device interaction.
• Batch computation logic: Verify L2 cache size calculation and batch count derivation align with profiling intent.
• Profiling loop modifications: Ensure batch rotation during timing loop and warmup-on-last-batch logic preserve measurement accuracy.
• Tuning configuration propagation: Confirm tuning_config threading through _profile_runners and _profile_single_kernel is complete and consistent.

Pre-merge checks and finishing touches

❌ Failed checks (2 warnings)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 27.27% which is insufficient. The required threshold is 80.00%.	You can run @coderabbitai generate docstrings to improve docstring coverage.
Description Check	⚠️ Warning	The PR description provides a clear explanation of the issue and solution in the opening paragraph: it describes why L2 cache clearing is necessary (L2 is typically cold in production but warm during benchmarking, affecting kernel selection accuracy) and what approach is used (circular buffer method for L2 cache clearing). However, the description is incomplete against the required template structure. The "Description" section within the template is left empty, and critically, the "Test Coverage" section is entirely blank with no mention of relevant tests that safeguard these changes. Additionally, while the PR Checklist is shown with one item checked at the end, the individual checklist items are not addressed or verified in the description, making it unclear whether the PR follows coding guidelines, includes test cases for new code paths, or whether dependencies were scanned and documentation was updated as required.

✅ Passed checks (1 passed)

Check name	Status	Explanation
Title Check	✅ Passed	The PR title "[TRTLLM-7963][feat] Cold L2 cache when doing autotune benchmarking." is fully related to the main change and follows the required format with a valid JIRA ticket, type indicator, and clear summary. The title accurately reflects the primary objective of the PR—implementing L2 cache clearing during autotuning using a circular buffer approach—and is specific enough for a developer scanning history to understand the core change without ambiguity.

✨ Finishing touches

• 📝 Generate docstrings

🧪 Generate unit tests (beta)

• Create PR with unit tests
• Post copyable unit tests in a comment

---

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.

❤️ Share

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands and usage tips.}

[coderabbitai]

Actionable comments posted: 2

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/autotuner.py (1)

1-1: Add required NVIDIA Apache-2.0 copyright header.

Per coding guidelines, all source files must include the NVIDIA Apache-2.0 copyright header with the current year (2025) at the top of the file.

As per coding guidelines

🧹 Nitpick comments (4)

tensorrt_llm/_torch/autotuner.py (4)

952-983: LGTM with optional memory consideration.

The batch preparation logic correctly computes the number of buffers needed to flush L2 cache by multiplying cache size by 3 and dividing by buffer size. The defensive check for zero-sized tensors is good.

Optional: For very large input tensors, creating num_buffers clones could consume significant GPU memory. Consider documenting this memory trade-off in a comment, or adding a maximum memory threshold check if memory exhaustion becomes a concern in practice.

---

750-757: LGTM with minor formatting note.

The batch preparation and warmup logic correctly implements L2 cache clearing by using multiple input batches.

Nitpick: Lines 751-752 contain extra blank lines that could be removed for consistency with the rest of the codebase.

---

1004-1010: Consider using current CUDA device instead of device_id=0.

The method defaults to device_id=0, but it should query the L2 cache size for the currently active CUDA device to ensure accurate profiling when using multi-GPU setups.

Apply this diff to use the current device:

-    def _get_l2_cache_size_in_bytes(self, device_id: int = 0) -> int:
-        device = self._checkCudaErrors(driver.cuDeviceGet(device_id))
+    def _get_l2_cache_size_in_bytes(self) -> int:
+        device_id = torch.cuda.current_device()
+        device = self._checkCudaErrors(driver.cuDeviceGet(device_id))
         return self._checkCudaErrors(
             driver.cuDeviceGetAttribute(
                 driver.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE,
                 device,
             ))

---

1026-1034: Use TypeError for invalid type and move import to top.

Two improvements:

1. Per static analysis (TRY004) and Python best practices, use TypeError instead of RuntimeError for type validation errors.
2. Move the nvrtc import to the top of the file for consistency.

Apply this diff to address both points:

At the top of the file with other imports:

 import torch
 from cuda.bindings import driver
+from cuda.bindings import nvrtc

In the method:

     def _cudaGetErrorEnum(self, error) -> str:
-        from cuda.bindings import nvrtc
         if isinstance(error, driver.CUresult):
             err, name = driver.cuGetErrorName(error)
             return name if err == driver.CUresult.CUDA_SUCCESS else "<unknown>"
         elif isinstance(error, nvrtc.nvrtcResult):
             return nvrtc.nvrtcGetErrorString(error)[1]
         else:
-            raise RuntimeError("Unknown error type: {}".format(error))
+            raise TypeError(f"Unknown error type: {type(error)}")

📜 Review details

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 6b755fd and dbd4a64.

📒 Files selected for processing (1)

• tensorrt_llm/_torch/autotuner.py (10 hunks)

🧰 Additional context used

📓 Path-based instructions (3)

**/*.{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:

• tensorrt_llm/_torch/autotuner.py

**/*.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/autotuner.py

**/*.{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:

• tensorrt_llm/_torch/autotuner.py

🧠 Learnings (3)

📚 Learning: 2025-08-26T09:37:10.463Z

Learnt from: jiaganc
PR: NVIDIA/TensorRT-LLM#7031
File: tensorrt_llm/bench/dataclasses/configuration.py:90-104
Timestamp: 2025-08-26T09:37:10.463Z
Learning: In TensorRT-LLM's bench configuration, the `get_pytorch_perf_config()` method returns `self.pytorch_config` which is a Dict[str, Any] that can contain default values including `cuda_graph_config`, making the fallback `llm_args["cuda_graph_config"]` safe to use.

Applied to files:

• tensorrt_llm/_torch/autotuner.py

📚 Learning: 2025-09-16T09:30:09.716Z

Learnt from: tongyuantongyu
PR: NVIDIA/TensorRT-LLM#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:

• tensorrt_llm/_torch/autotuner.py

📚 Learning: 2025-08-27T14:23:55.566Z

Learnt from: ixlmar
PR: NVIDIA/TensorRT-LLM#7294
File: tensorrt_llm/_torch/modules/rms_norm.py:17-17
Timestamp: 2025-08-27T14:23:55.566Z
Learning: The TensorRT-LLM project requires Python 3.10+ as evidenced by the use of TypeAlias from typing module, match/case statements, and union type | syntax throughout the codebase, despite some documentation still mentioning Python 3.8+.

Applied to files:

• tensorrt_llm/_torch/autotuner.py

🧬 Code graph analysis (1)

tensorrt_llm/_torch/autotuner.py (1)

tensorrt_llm/logger.py (2)

• info (138-139)
• error (126-127)

🪛 Ruff (0.14.2)

tensorrt_llm/_torch/autotuner.py

1016-1017: Avoid specifying long messages outside the exception class

(TRY003)

---

1034-1034: Prefer TypeError exception for invalid type

(TRY004)

⏰ 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 (6)

tensorrt_llm/_torch/autotuner.py (6)

16-16: LGTM!

The CUDA driver import is correctly added to support L2 cache size queries.

---

627-629: LGTM!

Elevating this log to INFO level provides better visibility into the autotuning process.

---

690-690: LGTM!

Correctly propagates tuning_config to enable L2 cache clearing in the profiling function.

---

732-732: LGTM!

Signature update correctly accepts the tuning_config parameter.

---

767-772: LGTM!

The timing loop correctly rotates through input batches using modulo indexing, ensuring L2 cache is flushed between profiling iterations when multiple batches are available.

---

1012-1024: LGTM!

The CUDA error checking implementation correctly handles CUDA API result tuples and provides informative error messages.

[tensorrt-cicd]

PR_Github #23008 [ run ] triggered by Bot. Commit: dbd4a64

[hyukn]

/bot run --disable-fail-fast

[limin2021]

LGTM.

[tensorrt-cicd]

PR_Github #23016 [ run ] triggered by Bot. Commit: 43ff5b1

[tensorrt-cicd]

PR_Github #23008 [ run ] completed with state ABORTED. Commit: dbd4a64
LLM/main/L0_MergeRequest_PR #17346 (Blue Ocean) completed with status: ABORTED

[tensorrt-cicd]

PR_Github #23016 [ run ] completed with state SUCCESS. Commit: 43ff5b1
/LLM/main/L0_MergeRequest_PR pipeline #17352 completed with status: 'FAILURE'

[hyukn]

/bot run --disable-fail-fast

[hyukn]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #23156 [ run ] triggered by Bot. Commit: 8ee7b5e

[tensorrt-cicd]

PR_Github #23157 [ run ] triggered by Bot. Commit: 8ee7b5e

[tensorrt-cicd]

PR_Github #23156 [ run ] completed with state ABORTED. Commit: 8ee7b5e

[tensorrt-cicd]

PR_Github #23157 [ run ] completed with state SUCCESS. Commit: 8ee7b5e
/LLM/main/L0_MergeRequest_PR pipeline #17462 completed with status: 'FAILURE'

[hyukn]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #23303 [ run ] triggered by Bot. Commit: 8ee7b5e

[tensorrt-cicd]

PR_Github #23303 [ run ] completed with state ABORTED. Commit: 8ee7b5e
LLM/main/L0_MergeRequest_PR #17558 (Blue Ocean) completed with status: ABORTED

[hyukn]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #23336 [ run ] triggered by Bot. Commit: 46ee028

[hyukn]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #23366 [ run ] triggered by Bot. Commit: 46ee028

[tensorrt-cicd]

PR_Github #23366 [ run ] completed with state SUCCESS. Commit: 46ee028
/LLM/main/L0_MergeRequest_PR pipeline #17607 completed with status: 'FAILURE'

[hyukn]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #23492 [ run ] triggered by Bot. Commit: 33a6f64

[tensorrt-cicd]

PR_Github #23492 [ run ] completed with state SUCCESS. Commit: 33a6f64
/LLM/main/L0_MergeRequest_PR pipeline #17683 completed with status: 'FAILURE'

[hyukn]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #23628 [ run ] triggered by Bot. Commit: 33a6f64

[tensorrt-cicd]

PR_Github #23628 [ run ] completed with state FAILURE. Commit: 33a6f64
/LLM/main/L0_MergeRequest_PR pipeline #17778 completed with status: 'FAILURE'

[hyukn]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #23958 [ run ] triggered by Bot. Commit: 9defb28

[tensorrt-cicd]

PR_Github #23958 [ run ] completed with state SUCCESS. Commit: 9defb28
/LLM/main/L0_MergeRequest_PR pipeline #18041 completed with status: 'FAILURE'

[hyukn]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #24079 [ run ] triggered by Bot. Commit: 9defb28

[tensorrt-cicd]

PR_Github #24079 [ run ] completed with state SUCCESS. Commit: 9defb28
/LLM/main/L0_MergeRequest_PR pipeline #18146 completed with status: 'FAILURE'

[hyukn]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #25568 [ run ] triggered by Bot. Commit: a8dc71c

[tensorrt-cicd]

PR_Github #25568 [ run ] completed with state SUCCESS. Commit: a8dc71c
/LLM/main/L0_MergeRequest_PR pipeline #19365 completed with status: 'SUCCESS'
Pipeline passed with automatic retried tests. Check the rerun report for details.

[hyukn]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #25643 [ run ] triggered by Bot. Commit: 605b2ab

[tensorrt-cicd]

PR_Github #25643 [ run ] completed with state SUCCESS. Commit: 605b2ab
/LLM/main/L0_MergeRequest_PR pipeline #19430 completed with status: 'SUCCESS'
Pipeline passed with automatic retried tests. Check the rerun report for details.