From adf566787ae432c135a9a64cace5667fd71a3cf3 Mon Sep 17 00:00:00 2001 From: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> Date: Fri, 20 Feb 2026 21:10:41 -0500 Subject: [PATCH 1/6] [None][feat] add globaltimer-based timing backend for autotuner profiling CUDA event timing (cudaEventElapsedTime) is unreliable when confidential compute (CC) is enabled (see https://nvbugs/4159316 for details). 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. The timing backend is auto-selected based on CC state, with an env var override (TLLM_PROFILING_TIMER=globaltimer|cuda_event) for manual control. Co-Authored-By: Claude Opus 4.6 (1M context) Signed-off-by: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> --- cpp/tensorrt_llm/kernels/globalTimerKernel.cu | 40 ++++++++++++++ cpp/tensorrt_llm/kernels/globalTimerKernel.h | 30 +++++++++++ .../nanobind/runtime/bindings.cpp | 12 +++++ tensorrt_llm/_torch/autotuner.py | 54 ++++++++++++++++--- 4 files changed, 129 insertions(+), 7 deletions(-) create mode 100644 cpp/tensorrt_llm/kernels/globalTimerKernel.cu create mode 100644 cpp/tensorrt_llm/kernels/globalTimerKernel.h diff --git a/cpp/tensorrt_llm/kernels/globalTimerKernel.cu b/cpp/tensorrt_llm/kernels/globalTimerKernel.cu new file mode 100644 index 000000000000..54e07d8a7608 --- /dev/null +++ b/cpp/tensorrt_llm/kernels/globalTimerKernel.cu @@ -0,0 +1,40 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + * + * 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. + */ + +#include "tensorrt_llm/common/config.h" +#include "tensorrt_llm/common/cudaUtils.h" +#include "tensorrt_llm/kernels/globalTimerKernel.h" + +using namespace tensorrt_llm::common; + +TRTLLM_NAMESPACE_BEGIN + +namespace kernels +{ +__global__ void readGlobalTimerKernel(uint64_t* timestamp) +{ + asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(*timestamp)); +} + +void invokeReadGlobalTimer(uint64_t* d_timestamp, cudaStream_t stream) +{ + readGlobalTimerKernel<<<1, 1, 0, stream>>>(d_timestamp); + check_cuda_error(cudaGetLastError()); +} +} // namespace kernels + +TRTLLM_NAMESPACE_END diff --git a/cpp/tensorrt_llm/kernels/globalTimerKernel.h b/cpp/tensorrt_llm/kernels/globalTimerKernel.h new file mode 100644 index 000000000000..630f2a8294f2 --- /dev/null +++ b/cpp/tensorrt_llm/kernels/globalTimerKernel.h @@ -0,0 +1,30 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + * + * 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. + */ + +#pragma once + +#include "tensorrt_llm/common/config.h" +#include "tensorrt_llm/common/cudaUtils.h" + +TRTLLM_NAMESPACE_BEGIN + +namespace kernels +{ +void invokeReadGlobalTimer(uint64_t* d_timestamp, cudaStream_t stream); +} // namespace kernels + +TRTLLM_NAMESPACE_END diff --git a/cpp/tensorrt_llm/nanobind/runtime/bindings.cpp b/cpp/tensorrt_llm/nanobind/runtime/bindings.cpp index df8120b3ef82..153c1a5dd5e0 100644 --- a/cpp/tensorrt_llm/nanobind/runtime/bindings.cpp +++ b/cpp/tensorrt_llm/nanobind/runtime/bindings.cpp @@ -22,6 +22,7 @@ #include "tensorrt_llm/kernels/communicationKernels/customLowPrecisionAllReduceKernels.h" #include "tensorrt_llm/kernels/customAllReduceKernels.h" #include "tensorrt_llm/kernels/delayStream.h" +#include "tensorrt_llm/kernels/globalTimerKernel.h" #include "tensorrt_llm/nanobind/common/customCasters.h" #include "tensorrt_llm/runtime/cudaEvent.h" #include "tensorrt_llm/runtime/cudaStream.h" @@ -315,6 +316,17 @@ void initBindings(nb::module_& m) tensorrt_llm::kernels::invokeDelayStreamKernel(delay_micro_secs, stream); }, "Delay kernel launch on the default stream"); + m.def( + "record_global_timer", + [](int64_t data_ptr, nb::object py_stream) + { + auto* ptr = reinterpret_cast(data_ptr); + auto stream_ptr = nb::cast(py_stream.attr("cuda_stream")); + cudaStream_t stream = reinterpret_cast(stream_ptr); + nb::gil_scoped_release release; + tensorrt_llm::kernels::invokeReadGlobalTimer(ptr, stream); + }, + "Record GPU global timer value to device memory"); m.def( "max_workspace_size_lowprecision", [](int32_t tp_size) { return tensorrt_llm::kernels::max_workspace_size_lowprecision(tp_size); }, diff --git a/tensorrt_llm/_torch/autotuner.py b/tensorrt_llm/_torch/autotuner.py index 22efbc76877b..f27830232413 100644 --- a/tensorrt_llm/_torch/autotuner.py +++ b/tensorrt_llm/_torch/autotuner.py @@ -19,8 +19,9 @@ import tensorrt_llm from tensorrt_llm._torch.distributed import Distributed -from tensorrt_llm._utils import nvtx_range -from tensorrt_llm.bindings.internal.runtime import delay_kernel +from tensorrt_llm._utils import confidential_compute_enabled, nvtx_range +from tensorrt_llm.bindings.internal.runtime import (delay_kernel, + record_global_timer) from tensorrt_llm.logger import logger from tensorrt_llm.mapping import Mapping @@ -712,6 +713,7 @@ class AutoTuner: stream_delay_micro_secs (int): Delay on CUDA stream before the profiled kernel runs in microseconds (default: 1000) """ _CUDA_GRAPH_DELAY_MICRO_SECS = 100 + _NS_PER_MS = 1e6 _instance = None def __init__(self, warmup=2, repeat=10, stream_delay_micro_secs=1000): @@ -726,6 +728,21 @@ def __init__(self, warmup=2, repeat=10, stream_delay_micro_secs=1000): self.profiling_cache = AutoTunerProfilingCache() self.is_tuning_mode = False + # Timing backend: globaltimer kernel vs cuda events. + # TLLM_PROFILING_TIMER env var overrides auto-detection: + # "globaltimer" -> force globaltimer + # "cuda_event" -> force cuda events + # unset/default -> auto-detect via confidential_compute_enabled() + timer_env = os.getenv("TLLM_PROFILING_TIMER", "").lower() + if timer_env == "globaltimer": + self._use_global_timer = True + elif timer_env == "cuda_event": + self._use_global_timer = False + else: + self._use_global_timer = confidential_compute_enabled() + + logger.debug(f"[Autotuner] use_global_timer: {self._use_global_timer}") + # Add statistics tracking self.stats = AutoTunerStatistics() @@ -1150,10 +1167,33 @@ def _profile_single_kernel( avg_time = float('inf') def pure_profile(stream: torch.cuda.Stream, repeat: int): - start = torch.cuda.Event(enable_timing=True) - end = torch.cuda.Event(enable_timing=True) graph = torch.cuda.CUDAGraph() + if self._use_global_timer: + start_ts = torch.empty(1, dtype=torch.int64, device='cuda') + end_ts = torch.empty(1, dtype=torch.int64, device='cuda') + + def record_start(): + record_global_timer(start_ts.data_ptr(), stream) + + def record_end(): + record_global_timer(end_ts.data_ptr(), stream) + + def elapsed_time(): + return (end_ts.item() - start_ts.item()) / self._NS_PER_MS + else: + start_evt = torch.cuda.Event(enable_timing=True) + end_evt = torch.cuda.Event(enable_timing=True) + + def record_start(): + start_evt.record() + + def record_end(): + end_evt.record() + + def elapsed_time(): + return start_evt.elapsed_time(end_evt) + with torch.cuda.stream(stream): if use_cuda_graph: with torch.cuda.graph(graph): @@ -1176,7 +1216,7 @@ def pure_profile(stream: torch.cuda.Stream, repeat: int): else: delay_kernel(self.stream_delay_micro_secs, stream) - start.record() + record_start() if use_cuda_graph: graph.replay() @@ -1188,10 +1228,10 @@ def pure_profile(stream: torch.cuda.Stream, repeat: int): **kwargs, ) - end.record() + record_end() stream.synchronize() - return start.elapsed_time(end) / repeat + return elapsed_time() / repeat # warm up, no timing for _ in range(self.warmup): From 3787ad5ad079a319aad86d0ee7c39bd87564218f Mon Sep 17 00:00:00 2001 From: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> Date: Mon, 23 Feb 2026 14:19:56 -0500 Subject: [PATCH 2/6] Incorporate some suggestions from coderabbit Signed-off-by: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> --- cpp/tensorrt_llm/nanobind/runtime/bindings.cpp | 2 +- tensorrt_llm/_torch/autotuner.py | 2 ++ 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/cpp/tensorrt_llm/nanobind/runtime/bindings.cpp b/cpp/tensorrt_llm/nanobind/runtime/bindings.cpp index 153c1a5dd5e0..182c85623012 100644 --- a/cpp/tensorrt_llm/nanobind/runtime/bindings.cpp +++ b/cpp/tensorrt_llm/nanobind/runtime/bindings.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 * * Licensed under the Apache License, Version 2.0 (the "License"); diff --git a/tensorrt_llm/_torch/autotuner.py b/tensorrt_llm/_torch/autotuner.py index f27830232413..fc014f2a1018 100644 --- a/tensorrt_llm/_torch/autotuner.py +++ b/tensorrt_llm/_torch/autotuner.py @@ -1180,6 +1180,8 @@ def record_end(): record_global_timer(end_ts.data_ptr(), stream) def elapsed_time(): + # GPU %globaltimer counts in nanoseconds, + # Torch.cuda.Event.elapsed_time() returns ms return (end_ts.item() - start_ts.item()) / self._NS_PER_MS else: start_evt = torch.cuda.Event(enable_timing=True) From 95580831134d11e796fd6cf4e5813c62369667a5 Mon Sep 17 00:00:00 2001 From: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> Date: Wed, 25 Feb 2026 08:22:54 -0500 Subject: [PATCH 3/6] add unit test comparing globaltimer and cuda-event timing backends Profiles existing GemmRunner tactics with both backends and asserts the measured times agree within 5%. Parametrized over use_cuda_graph to cover both graph-replay and loop-based profiling paths. Co-Authored-By: Claude Opus 4.6 (1M context) Signed-off-by: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> --- tests/unittest/_torch/misc/test_autotuner.py | 41 ++++++++++++++++++++ 1 file changed, 41 insertions(+) diff --git a/tests/unittest/_torch/misc/test_autotuner.py b/tests/unittest/_torch/misc/test_autotuner.py index a41fe443a571..106399dadca4 100644 --- a/tests/unittest/_torch/misc/test_autotuner.py +++ b/tests/unittest/_torch/misc/test_autotuner.py @@ -868,3 +868,44 @@ def test_autotuner_distributed_strategy(strategy, mpi_pool_executor): ) for r in results: assert r is True + + +@pytest.mark.parametrize("use_cuda_graph", [False, True]) +def test_global_timer_vs_cuda_event(use_cuda_graph): + """Verify globaltimer and cuda-event backends report times within 5%.""" + runner = GemmRunner() + x = torch.randn(M // 2, 64, device='cuda') + w = torch.randn(64, 128, device='cuda') + tuning_config = TuningConfig(use_cuda_graph=use_cuda_graph) + + tuner = AutoTuner() + + for tactic in runner.get_valid_tactics([x, w], OptimizationProfile()): + # Profile with cuda events + tuner._use_global_timer = False + event_time = tuner._profile_single_kernel( + runner=runner, + inputs=[x, w], + tactic=tactic, + tuning_config=tuning_config, + use_cuda_graph=use_cuda_graph, + ) + + # Profile with globaltimer + tuner._use_global_timer = True + gt_time = tuner._profile_single_kernel( + runner=runner, + inputs=[x, w], + tactic=tactic, + tuning_config=tuning_config, + use_cuda_graph=use_cuda_graph, + ) + + assert event_time > 0, f"cuda event time should be positive, got {event_time}" + assert gt_time > 0, f"globaltimer time should be positive, got {gt_time}" + + rel_diff = abs(gt_time - event_time) / event_time + assert rel_diff < 0.05, ( + f"tactic={tactic}: globaltimer ({gt_time:.4f}ms) and cuda event " + f"({event_time:.4f}ms) differ by {rel_diff * 100:.1f}%, expected < 5%" + ) From 2880762fdc945ab2d31ad6de73ac58f2f2cff2d2 Mon Sep 17 00:00:00 2001 From: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> Date: Wed, 25 Feb 2026 08:40:40 -0500 Subject: [PATCH 4/6] _NS_PER_MS -> constant, since the inline constant is then colocated with the comment that describes its function Signed-off-by: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> --- tensorrt_llm/_torch/autotuner.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tensorrt_llm/_torch/autotuner.py b/tensorrt_llm/_torch/autotuner.py index fc014f2a1018..78b205ae0db7 100644 --- a/tensorrt_llm/_torch/autotuner.py +++ b/tensorrt_llm/_torch/autotuner.py @@ -713,7 +713,6 @@ class AutoTuner: stream_delay_micro_secs (int): Delay on CUDA stream before the profiled kernel runs in microseconds (default: 1000) """ _CUDA_GRAPH_DELAY_MICRO_SECS = 100 - _NS_PER_MS = 1e6 _instance = None def __init__(self, warmup=2, repeat=10, stream_delay_micro_secs=1000): @@ -1182,7 +1181,7 @@ def record_end(): def elapsed_time(): # GPU %globaltimer counts in nanoseconds, # Torch.cuda.Event.elapsed_time() returns ms - return (end_ts.item() - start_ts.item()) / self._NS_PER_MS + return (end_ts.item() - start_ts.item()) / 1e6 else: start_evt = torch.cuda.Event(enable_timing=True) end_evt = torch.cuda.Event(enable_timing=True) From 3550c8179155ca75618fd26151719719c7c9ef8c Mon Sep 17 00:00:00 2001 From: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> Date: Wed, 25 Feb 2026 09:50:07 -0500 Subject: [PATCH 5/6] Make comment more clear Signed-off-by: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> --- tensorrt_llm/_torch/autotuner.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorrt_llm/_torch/autotuner.py b/tensorrt_llm/_torch/autotuner.py index 78b205ae0db7..f14bee0e7552 100644 --- a/tensorrt_llm/_torch/autotuner.py +++ b/tensorrt_llm/_torch/autotuner.py @@ -1179,8 +1179,8 @@ def record_end(): record_global_timer(end_ts.data_ptr(), stream) def elapsed_time(): - # GPU %globaltimer counts in nanoseconds, - # Torch.cuda.Event.elapsed_time() returns ms + # GPU %globaltimer counts in ns; convert to ms to match the + # units of Torch.cuda.Event.elapsed_time() return (end_ts.item() - start_ts.item()) / 1e6 else: start_evt = torch.cuda.Event(enable_timing=True) From 9e451d85451fbb47417671f42fdf59b048329e19 Mon Sep 17 00:00:00 2001 From: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> Date: Thu, 26 Feb 2026 14:03:18 -0500 Subject: [PATCH 6/6] Profile actual GEMMs and make the comparison more statistically rigorous Signed-off-by: Dan Hansen <1+dhansen-nvidia@users.noreply.github.com> --- tests/unittest/_torch/misc/test_autotuner.py | 141 ++++++++++++++----- 1 file changed, 109 insertions(+), 32 deletions(-) diff --git a/tests/unittest/_torch/misc/test_autotuner.py b/tests/unittest/_torch/misc/test_autotuner.py index 106399dadca4..3bacfa74071e 100644 --- a/tests/unittest/_torch/misc/test_autotuner.py +++ b/tests/unittest/_torch/misc/test_autotuner.py @@ -1,7 +1,9 @@ import itertools import json +import math import os import pickle +import statistics import sys import tempfile from typing import Any, List @@ -871,41 +873,116 @@ def test_autotuner_distributed_strategy(strategy, mpi_pool_executor): @pytest.mark.parametrize("use_cuda_graph", [False, True]) -def test_global_timer_vs_cuda_event(use_cuda_graph): - """Verify globaltimer and cuda-event backends report times within 5%.""" - runner = GemmRunner() - x = torch.randn(M // 2, 64, device='cuda') - w = torch.randn(64, 128, device='cuda') - tuning_config = TuningConfig(use_cuda_graph=use_cuda_graph) +def test_global_timer_vs_cuda_event(use_cuda_graph, monkeypatch): + """Verify globaltimer and cuda-event backends are statistically indistinguishable.""" - tuner = AutoTuner() + class PureGemmRunner(TunableRunner): - for tactic in runner.get_valid_tactics([x, w], OptimizationProfile()): - # Profile with cuda events - tuner._use_global_timer = False - event_time = tuner._profile_single_kernel( - runner=runner, - inputs=[x, w], - tactic=tactic, - tuning_config=tuning_config, - use_cuda_graph=use_cuda_graph, - ) + def get_valid_tactics(self, inputs: List[FakeTensor], + profile: OptimizationProfile, + **kwargs) -> List[int]: + return [0] - # Profile with globaltimer - tuner._use_global_timer = True - gt_time = tuner._profile_single_kernel( - runner=runner, - inputs=[x, w], - tactic=tactic, - tuning_config=tuning_config, - use_cuda_graph=use_cuda_graph, - ) + def forward(self, + /, + inputs: List[torch.Tensor], + *, + tactic: int = 0, + **kwargs) -> torch.Tensor: + assert tactic == 0 + return inputs[0] @ inputs[1] - assert event_time > 0, f"cuda event time should be positive, got {event_time}" - assert gt_time > 0, f"globaltimer time should be positive, got {gt_time}" + # Keep full profiling repeats enabled to reduce measurement noise. + monkeypatch.setenv("TLLM_AUTOTUNER_DISABLE_SHORT_PROFILE", "1") - rel_diff = abs(gt_time - event_time) / event_time - assert rel_diff < 0.05, ( - f"tactic={tactic}: globaltimer ({gt_time:.4f}ms) and cuda event " - f"({event_time:.4f}ms) differ by {rel_diff * 100:.1f}%, expected < 5%" + gemm_shapes = [ + (256, 4096, 11008), + (512, 8192, 8192), + ] + num_trials = 6 + rel_tol = 0.05 + stat_zscore = 3.0 + abs_tol_ms = 0.01 + + runner = PureGemmRunner() + tuning_config = TuningConfig(use_cuda_graph=use_cuda_graph) + tuner = AutoTuner() + trial_rows = [] + + for m, k, n in gemm_shapes: + x = torch.randn(m, k, device='cuda', dtype=torch.float16) + w = torch.randn(k, n, device='cuda', dtype=torch.float16) + + event_times = [] + gt_times = [] + + # Interleave both backends to avoid drift effects from neighboring load. + for _ in range(num_trials): + tuner._use_global_timer = False + event_times.append( + tuner._profile_single_kernel( + runner=runner, + inputs=[x, w], + tactic=0, + tuning_config=tuning_config, + use_cuda_graph=use_cuda_graph, + )) + + tuner._use_global_timer = True + gt_times.append( + tuner._profile_single_kernel( + runner=runner, + inputs=[x, w], + tactic=0, + tuning_config=tuning_config, + use_cuda_graph=use_cuda_graph, + )) + + event_ms = event_times[-1] + gt_ms = gt_times[-1] + abs_diff = abs(gt_ms - event_ms) + rel_diff = abs_diff / event_ms if event_ms > 0 else float('inf') + trial_rows.append((m, k, n, len(event_times), event_ms, gt_ms, + abs_diff, rel_diff)) + + event_mean = statistics.fmean(event_times) + gt_mean = statistics.fmean(gt_times) + event_var = statistics.variance(event_times) + gt_var = statistics.variance(gt_times) + mean_diff = abs(gt_mean - event_mean) + rel_diff = mean_diff / event_mean + + # Two-sample mean delta should be small vs a fixed tolerance and + # indistinguishable within sampling noise. + combined_sem = math.sqrt(event_var / num_trials + gt_var / num_trials) + allowed_diff = max(abs_tol_ms, rel_tol * event_mean, + stat_zscore * combined_sem) + + assert event_mean > 0, ( + f"({m},{k},{n}): cuda event mean should be positive, got {event_mean}" + ) + assert gt_mean > 0, ( + f"({m},{k},{n}): globaltimer mean should be positive, got {gt_mean}" ) + assert mean_diff <= allowed_diff, ( + f"({m},{k},{n}): timing backends are distinguishable " + f"(cuda_event_mean={event_mean:.4f}ms, " + f"globaltimer_mean={gt_mean:.4f}ms, " + f"relative_diff={rel_diff * 100:.2f}%, " + f"allowed_diff={allowed_diff:.4f}ms, " + f"combined_sem={combined_sem:.4f}ms, " + f"event_samples={event_times}, gt_samples={gt_times})") + + # Visible with `pytest -s`; otherwise captured by pytest. + print("\nGlobaltimer vs cuda-event trial comparison") + print(f"cuda_graph={use_cuda_graph}, trials_per_shape={num_trials}") + print("-" * 102) + print( + f"{'shape (M,K,N)':>21} {'trial':>5} {'cuda_event (ms)':>16} " + f"{'globaltimer (ms)':>17} {'abs diff (ms)':>14} {'rel diff (%)':>13}") + print("-" * 102) + for m, k, n, trial, event_ms, gt_ms, abs_diff, rel_diff in trial_rows: + print(f"{f'({m},{k},{n})':>21} {trial:>5d} " + f"{event_ms:>16.4f} {gt_ms:>17.4f} " + f"{abs_diff:>14.4f} {rel_diff * 100:>13.2f}") + print("-" * 102)