fix(autotuner): time tactics with %globaltimer under Confidential Computing - #3870
Conversation
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
📝 WalkthroughWalkthroughAdds cached confidential-compute detection and JIT globaltimer support, selects the autotuner timing backend, updates profiling measurements, and adds CUDA tests comparing globaltimer with CUDA events. ChangesGlobaltimer-based autotuner timing
Estimated code review effort: 4 (Complex) | ~45 minutes Sequence Diagram(s)sequenceDiagram
participant AutoTuner
participant is_confidential_compute
participant get_globaltimer_kernel
participant GPUWork
AutoTuner->>is_confidential_compute: detect confidential-compute state
AutoTuner->>get_globaltimer_kernel: acquire globaltimer recorder
AutoTuner->>GPUWork: record start and execute profiled work
AutoTuner->>GPUWork: record end and synchronize
GPUWork-->>AutoTuner: elapsed timing result
Suggested reviewers: 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 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 |
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
There was a problem hiding this comment.
Code Review
This pull request introduces a %globaltimer timing backend for the autotuner to support reliable performance measurements under NVIDIA Confidential Computing (CC) environments. It adds CC detection via NVML, a JIT-compiled CUDA kernel to read the %globaltimer register, and corresponding tests to validate its accuracy against standard CUDA events. The review feedback highlights a critical issue in the JIT compilation tag generation: using the process ID as a fallback when LOCAL_RANK is unset prevents PyTorch from caching the compiled extension across runs, causing unnecessary compilation overhead. Using a stable fallback name like 'single' is recommended to preserve caching.
Important
The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.
| from torch.utils.cpp_extension import load_inline | ||
|
|
||
| # per-process name so the 4 TP ranks don't race on the JIT build dir | ||
| tag = os.environ.get("LOCAL_RANK") or str(os.getpid()) |
There was a problem hiding this comment.
Using str(os.getpid()) as a fallback when LOCAL_RANK is not set completely defeats PyTorch's JIT compilation caching across runs. Since the process ID (PID) changes on every execution, a new build directory (e.g., flashinfer_globaltimer_<pid>) will be created and compiled from scratch every time the program is run. This introduces a 5-10 second compilation delay on every run and leaks compiled .so files in ~/.cache/torch_extensions over time.\n\nSince single-process runs do not have concurrent ranks to race with, we can safely use a stable fallback name like "single". PyTorch's load_inline also has built-in file locking to handle any concurrent processes using the same name.\n\nConsider using a stable fallback name instead of os.getpid().
tag = os.environ.get(\"LOCAL_RANK\") or os.environ.get(\"RANK\") or \"single\"![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
🧹 Nitpick comments (1)
flashinfer/autotuner.py (1)
873-882: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick winCC fallback silently reverts to the unreliable timer.
When CC is auto-detected but the globaltimer kernel fails to build,
_use_global_timeris flipped back toFalse, so profiling usescudaEvent— precisely the path that this PR documents as unreliable under Confidential Computing (can rank tactics wrong and poison the cache).get_globaltimer_kernel()logs a generic warning, but here thelogger.debughides the CC-specific implication. Consider emitting alogger.warningon this branch when CC was detected, so operators know autotuning results may be untrustworthy.♻️ Suggested clearer warning under CC
if self._use_global_timer: self._record_global_timer = get_globaltimer_kernel() if self._record_global_timer is None: # Fallback to cudaEvent if the globaltimer kernel build failed self._use_global_timer = False + if timer_env != "cuda_event": + logger.warning( + "[Autotuner] globaltimer kernel unavailable; falling back " + "to cudaEvent timing. Under Confidential Computing this " + "timing may be unreliable and can degrade tactic selection." + ) else: self._record_global_timer = None🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/autotuner.py` around lines 873 - 882, The fallback in autotuner._use_global_timer currently silently switches to cudaEvent when get_globaltimer_kernel() fails, which hides the Confidential Computing-specific risk. Update the branch in Autotuner logic to emit a logger.warning when CC was auto-detected and _use_global_timer is forced back to False, so the message clearly states autotuning may be unreliable; keep the existing logger.debug only for non-CC status reporting.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Nitpick comments:
In `@flashinfer/autotuner.py`:
- Around line 873-882: The fallback in autotuner._use_global_timer currently
silently switches to cudaEvent when get_globaltimer_kernel() fails, which hides
the Confidential Computing-specific risk. Update the branch in Autotuner logic
to emit a logger.warning when CC was auto-detected and _use_global_timer is
forced back to False, so the message clearly states autotuning may be
unreliable; keep the existing logger.debug only for non-CC status reporting.
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: 4c57929c-658d-4a43-b405-942061820b74
📥 Commits
Reviewing files that changed from the base of the PR and between 5b8da12 and 364e8b95379f01693b2e18f529dbdb00d24f8d47.
📒 Files selected for processing (3)
flashinfer/autotuner.pyflashinfer/utils.pytests/autotuner/test_global_timer.py
|
cc @leejnau for SGLang request tracking. thanks |
|
/bot run tests/autotuner |
|
[FAILED] Pipeline #57282953: 14/20 passed |
|
/bot run tests/autotuner |
|
[SUCCESS] Pipeline #57352239: 14/20 passed |
971ba8a to
d87ce76
Compare
|
/bot run tests/autotuner |
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@flashinfer/utils.py`:
- Around line 1384-1386: Update the environment override handling in the
confidential-compute detection function containing
FLASHINFER_CONFIDENTIAL_COMPUTE: accept only the documented values "0" and "1",
returning False or True respectively, and raise an appropriate error for any
other non-null value instead of silently treating it as disabled.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: cd92a73d-e4c9-4ae4-9706-bf9c6c6e7480
📥 Commits
Reviewing files that changed from the base of the PR and between 971ba8a974435967377898ea260be3324b40e3f6 and d87ce76479e201afc88c3cd4665892b721ae025a.
📒 Files selected for processing (3)
flashinfer/autotuner/autotuner.pyflashinfer/utils.pytests/autotuner/test_global_timer.py
🚧 Files skipped from review as they are similar to previous changes (2)
- flashinfer/autotuner/autotuner.py
- tests/autotuner/test_global_timer.py
|
[FAILED] Pipeline #57477181: 12/20 passed |
| cuda_sources=_GLOBALTIMER_KERNEL_CU, | ||
| functions=["get_globaltimer_timestamp"], | ||
| verbose=False, | ||
| ) |
There was a problem hiding this comment.
Per-rank name is counterproductive and still races across nodes.
Suggest dropping the tag and using a fixed name (flashinfer_globaltimer) so torch's built-in lock does the right thing. If isolation is truly needed, key it on a globally-unique id (hostname + global RANK), not LOCAL_RANK.
There was a problem hiding this comment.
Updated to use the fixed name flashinfer_globaltimer
| # "globaltimer" -> force globaltimer | ||
| # "cuda_event" -> force cuda events | ||
| # unset/default -> auto-detect via is_confidential_compute() | ||
| timer_env = os.getenv("FLASHINFER_AUTOTUNE_TIMER", "").lower() |
There was a problem hiding this comment.
the import hunk in this file only touches the flashinfer.utils import, no import os is added. If flashinfer/autotuner/autotuner.py doesn't already import os at module level, the first AutoTuner.get() will raise NameError. Please double-check and add import os if missing.
There was a problem hiding this comment.
os is already imported in flashinfer/autotuner/autotuner.py
d87ce76 to
8a14bdb
Compare
There was a problem hiding this comment.
Actionable comments posted: 4
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@flashinfer/autotuner/autotuner.py`:
- Around line 1085-1091: Update the timer selection logic in the autotuner
initialization around FLASHINFER_AUTOTUNE_TIMER to accept only the documented
globaltimer and cuda_event values (case-insensitively); reject any other
non-empty value explicitly instead of falling back to is_confidential_compute(),
while preserving auto-detection when the environment variable is unset or empty.
- Around line 2241-2246: Rename the generator variable `input` in the
`one_buffer_bytes` calculation to a non-shadowing name, and update its
`torch.Tensor` check and `numel()`/`element_size()` references consistently.
In `@flashinfer/utils.py`:
- Around line 1387-1401: Update is_confidential_compute() so failures to import
pynvml, initialize/query NVML, or otherwise complete confidential-compute
detection are treated as unknown rather than returning False; return the safe
globaltimer-oriented result expected by AutoTuner (or propagate a clear failure
requiring the override). Preserve False only for a successful check that
confirms CC is disabled.
- Line 22: Update the decorator on get_globaltimer_kernel to use functools.cache
instead of lru_cache(maxsize=1), and adjust the import accordingly so the module
follows the established caching convention.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: 2cfa65ee-59fc-49db-9cc9-b1fc2963e8f2
📥 Commits
Reviewing files that changed from the base of the PR and between d87ce76479e201afc88c3cd4665892b721ae025a and 8a14bdbbd62b986c9721bf97f014d96d44f66ffb.
📒 Files selected for processing (3)
flashinfer/autotuner/autotuner.pyflashinfer/utils.pytests/autotuner/test_global_timer.py
🚧 Files skipped from review as they are similar to previous changes (1)
- tests/autotuner/test_global_timer.py
There was a problem hiding this comment.
Caution
Inline review comments failed to post. This is likely due to GitHub's internal server error or limits when posting large numbers of comments. If you are seeing this consistently it is likely a permissions issue. Please check "Moderation" -> "Code review limits" under your organization settings.
Actionable comments posted: 4
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@flashinfer/autotuner/autotuner.py`:
- Around line 1085-1091: Update the timer selection logic in the autotuner
initialization around FLASHINFER_AUTOTUNE_TIMER to accept only the documented
globaltimer and cuda_event values (case-insensitively); reject any other
non-empty value explicitly instead of falling back to is_confidential_compute(),
while preserving auto-detection when the environment variable is unset or empty.
- Around line 2241-2246: Rename the generator variable `input` in the
`one_buffer_bytes` calculation to a non-shadowing name, and update its
`torch.Tensor` check and `numel()`/`element_size()` references consistently.
In `@flashinfer/utils.py`:
- Around line 1387-1401: Update is_confidential_compute() so failures to import
pynvml, initialize/query NVML, or otherwise complete confidential-compute
detection are treated as unknown rather than returning False; return the safe
globaltimer-oriented result expected by AutoTuner (or propagate a clear failure
requiring the override). Preserve False only for a successful check that
confirms CC is disabled.
- Line 22: Update the decorator on get_globaltimer_kernel to use functools.cache
instead of lru_cache(maxsize=1), and adjust the import accordingly so the module
follows the established caching convention.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: 2cfa65ee-59fc-49db-9cc9-b1fc2963e8f2
📥 Commits
Reviewing files that changed from the base of the PR and between d87ce76479e201afc88c3cd4665892b721ae025a and 8a14bdbbd62b986c9721bf97f014d96d44f66ffb.
📒 Files selected for processing (3)
flashinfer/autotuner/autotuner.pyflashinfer/utils.pytests/autotuner/test_global_timer.py
🚧 Files skipped from review as they are similar to previous changes (1)
- tests/autotuner/test_global_timer.py
🛑 Comments failed to post (4)
flashinfer/autotuner/autotuner.py (2)
1085-1091: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
Reject unsupported
FLASHINFER_AUTOTUNE_TIMERvalues.Line 1090 treats typos as auto-detection, silently ignoring an explicit backend request. Validate the documented values.
Proposed fix
timer_env = os.getenv("FLASHINFER_AUTOTUNE_TIMER", "").lower() if timer_env == "globaltimer": self._use_global_timer = True elif timer_env == "cuda_event": self._use_global_timer = False - else: + elif not timer_env: self._use_global_timer = is_confidential_compute() + else: + raise ValueError( + "FLASHINFER_AUTOTUNE_TIMER must be 'globaltimer' or 'cuda_event'" + )📝 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.timer_env = os.getenv("FLASHINFER_AUTOTUNE_TIMER", "").lower() if timer_env == "globaltimer": self._use_global_timer = True elif timer_env == "cuda_event": self._use_global_timer = False elif not timer_env: self._use_global_timer = is_confidential_compute() else: raise ValueError( "FLASHINFER_AUTOTUNE_TIMER must be 'globaltimer' or 'cuda_event'" )🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/autotuner/autotuner.py` around lines 1085 - 1091, Update the timer selection logic in the autotuner initialization around FLASHINFER_AUTOTUNE_TIMER to accept only the documented globaltimer and cuda_event values (case-insensitively); reject any other non-empty value explicitly instead of falling back to is_confidential_compute(), while preserving auto-detection when the environment variable is unset or empty.
2241-2246: 📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win
Rename
inputto avoid failing Ruff A001.The generator variable shadows Python’s
inputbuiltin.Proposed fix
one_buffer_bytes = sum( - input.numel() * input.element_size() - if isinstance(input, torch.Tensor) + item.numel() * item.element_size() + if isinstance(item, torch.Tensor) else 0 - for input in inputs + for item in inputs )📝 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.one_buffer_bytes = sum( item.numel() * item.element_size() if isinstance(item, torch.Tensor) else 0 for item in inputs )🧰 Tools
🪛 Ruff (0.15.21)
[error] 2245-2245: Variable
inputis shadowing a Python builtin(A001)
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/autotuner/autotuner.py` around lines 2241 - 2246, Rename the generator variable `input` in the `one_buffer_bytes` calculation to a non-shadowing name, and update its `torch.Tensor` check and `numel()`/`element_size()` references consistently.Source: Linters/SAST tools
flashinfer/utils.py (2)
22-22: 📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash set -euo pipefail # Locate the target file and the relevant symbols/usages. git ls-files flashinfer/utils.py printf '\n-- lru_cache/cache occurrences --\n' rg -n "lru_cache|functools\.cache|get_globaltimer_kernel|JIT module|module loader" flashinfer/utils.py printf '\n-- file outline --\n' ast-grep outline flashinfer/utils.py --view expanded || true printf '\n-- relevant line ranges --\n' sed -n '1,80p' flashinfer/utils.py printf '\n----\n' sed -n '1380,1435p' flashinfer/utils.pyRepository: flashinfer-ai/flashinfer
Length of output: 8322
Use
functools.cacheforget_globaltimer_kernel
Replace@lru_cache(maxsize=1)here with@functools.cacheso this JIT loader follows the module-cache convention.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/utils.py` at line 22, Update the decorator on get_globaltimer_kernel to use functools.cache instead of lru_cache(maxsize=1), and adjust the import accordingly so the module follows the established caching convention.Source: Coding guidelines
1387-1401: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash set -euo pipefail # Inspect the relevant function and nearby context. sed -n '1360,1435p' flashinfer/utils.py # Find all references to the confidential-compute detection helper and related timing selection. rg -n "Confidential-compute|ccFeature|nvmlSystemGetConfComputeState|globaltimer|cuda-event|AutoTuner|conf.*compute|confcompute|conf_compute|confidential" flashinferRepository: flashinfer-ai/flashinfer
Length of output: 16666
🏁 Script executed:
#!/bin/bash set -euo pipefail sed -n '1070,1105p' flashinfer/autotuner/autotuner.pyRepository: flashinfer-ai/flashinfer
Length of output: 1960
Treat CC-detection failures as unknown
is_confidential_compute()returnsFalseon missingpynvml, permission errors, or NVML exceptions, soAutoTunersilently selects CUDA-event timing on a CC host unlessFLASHINFER_AUTOTUNE_TIMER=globaltimeris set. Default to globaltimer (or fail loudly and require the override) when detection cannot complete.🧰 Tools
🪛 Ruff (0.15.21)
[warning] 1399-1399: Do not catch blind exception:
Exception(BLE001)
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/utils.py` around lines 1387 - 1401, Update is_confidential_compute() so failures to import pynvml, initialize/query NVML, or otherwise complete confidential-compute detection are treated as unknown rather than returning False; return the safe globaltimer-oriented result expected by AutoTuner (or propagate a clear failure requiring the override). Preserve False only for a successful check that confirms CC is disabled.
8a14bdb to
0632d30
Compare
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (1)
flashinfer/utils.py (1)
1404-1405: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueUse
@functools.cacheper repo guidelines instead oflru_cache(maxsize=1).
functoolsis already imported (Line 19);functools.cacheis the documented pattern for caching JIT module loading in this codebase.As per coding guidelines,
flashinfer/**/*.py: "Cache Python-level JIT module loading with@functools.cacheto avoid repeated compilation/loading."♻️ Proposed fix
-@lru_cache(maxsize=1) +@functools.cache def get_globaltimer_kernel():🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@flashinfer/utils.py` around lines 1404 - 1405, Update the get_globaltimer_kernel decorator from functools.lru_cache(maxsize=1) to functools.cache, preserving the existing single-entry caching behavior and function implementation.Source: Coding guidelines
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@flashinfer/utils.py`:
- Around line 1377-1436: Document the user-facing environment variables
FLASHINFER_CONFIDENTIAL_COMPUTE and FLASHINFER_AUTOTUNE_TIMER in CLAUDE.md and
the relevant .claude/skills guide. Describe their purpose and accepted
configuration values based on the behavior implemented by
is_confidential_compute and the autotune timer configuration, without changing
the runtime code.
---
Nitpick comments:
In `@flashinfer/utils.py`:
- Around line 1404-1405: Update the get_globaltimer_kernel decorator from
functools.lru_cache(maxsize=1) to functools.cache, preserving the existing
single-entry caching behavior and function implementation.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: 6fd8ec05-650e-47a2-b977-6eac05530960
📥 Commits
Reviewing files that changed from the base of the PR and between 8a14bdbbd62b986c9721bf97f014d96d44f66ffb and 0632d301787a8d727aafbdc9c69bbc0ba23f5303.
📒 Files selected for processing (3)
flashinfer/autotuner/autotuner.pyflashinfer/utils.pytests/autotuner/test_global_timer.py
🚧 Files skipped from review as they are similar to previous changes (2)
- tests/autotuner/test_global_timer.py
- flashinfer/autotuner/autotuner.py
Under Confidential Computing, cudaEventElapsedTime is unreliable (can return negative values on the bounce-buffer path), so AutoTuner.choose_one's min(measured_time) ranking picks a near-random tactic per rank and bakes it into the tuning cache. We instead time candidate tactics with the GPU's %globaltimer register (read from a tiny stamp kernel), which is CC-safe and consistent with how TRT-LLM times tactics(NVIDIA/TensorRT-LLM#11657). CC detection (is_confidential_compute) and the %globaltimer stamp kernel (get_globaltimer_kernel) live in flashinfer/utils.py -- both lru_cache'd and shareable. AutoTuner.__init__ selects the backend once and pure_profile brackets the run through a unified record_start/record_end/elapsed_time path (globaltimer stamps vs cuda events). Controlled by FLASHINFER_AUTOTUNE_TIMER: - "globaltimer" — force globaltimer - "cuda_event" — force the legacy cudaEvent timer - Others (default) — globaltimer iff CC is detected CC detection is via NVML; can override with FLASHINFER_CONFIDENTIAL_COMPUTE=1/0. Co-Authored-By: spethe <spethe@nvidia.com> Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…R, FLASHINFER_CONFIDENTIAL_COMPUTE) Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
0632d30 to
1e55fda
Compare
|
/bot run tests/autotuner |
|
[SUCCESS] Pipeline #57999668: 14/20 passed |
5 participants
Problem
CUDA event timing (cudaEventElapsedTime) is unreliable when confidential compute (CC) is enabled. This adds an alternative timing backend that uses a small CUDA kernel reading %globaltimer before and after the profiled work, then computes elapsed time on the host.
Fix
Time candidate tactics with the GPU's
%globaltimerregister (a tiny JIT stamp kernel) instead of CUDA events — monotonic and CC-safe. The return value/signature of the timed path is unchanged, sochoose_oneand the tuning-cache format are untouched.flashinfer/utils.py:is_confidential_compute()(NVML CC detection) andget_globaltimer_kernel()(lazyload_inlinestamp kernel), both@lru_cache'd and reusable.flashinfer/autotuner.py:AutoTuner.__init__selects the timing backend once;pure_profilebrackets the run through a unifiedrecord_start/record_end/elapsed_timepath (globaltimer stamps vs CUDA events).Off-CC behavior is unchanged — by default the
%globaltimerpath is used only when CC is detected; otherwise the legacycudaEventtimer is kept. If the stamp kernel fails to build, it falls back tocudaEvent.Mirrors TensorRT-LLM PR #11657.
Control
FLASHINFER_AUTOTUNE_TIMER:globaltimer— force%globaltimercuda_event— force the legacy cudaEvent timer%globaltimeriff CC is detectedCC detection is via NVML; override with
FLASHINFER_CONFIDENTIAL_COMPUTE=1/0.Test
tests/autotuner/test_global_timer.py:test_global_timer_vs_cuda_event[shape-mode]— times identical GEMM work with both timers, interleaved over 6 trials, in eager and CUDA-graph modes; asserts both means > 0 and that they agree withinmax(0.01 ms, 5% · event_mean, 3 · combined_SEM).test_globaltimer_monotonic— sanity: two stamps around real work are strictly increasing, and more work reads a larger delta.Passing on H100:
Summary by CodeRabbit