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..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"); @@ -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..f14bee0e7552 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 @@ -726,6 +727,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 +1166,35 @@ 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(): + # 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) + 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 +1217,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 +1229,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): diff --git a/tests/unittest/_torch/misc/test_autotuner.py b/tests/unittest/_torch/misc/test_autotuner.py index a41fe443a571..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 @@ -868,3 +870,119 @@ 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, monkeypatch): + """Verify globaltimer and cuda-event backends are statistically indistinguishable.""" + + class PureGemmRunner(TunableRunner): + + def get_valid_tactics(self, inputs: List[FakeTensor], + profile: OptimizationProfile, + **kwargs) -> List[int]: + return [0] + + def forward(self, + /, + inputs: List[torch.Tensor], + *, + tactic: int = 0, + **kwargs) -> torch.Tensor: + assert tactic == 0 + return inputs[0] @ inputs[1] + + # Keep full profiling repeats enabled to reduce measurement noise. + monkeypatch.setenv("TLLM_AUTOTUNER_DISABLE_SHORT_PROFILE", "1") + + 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)