[None][feat] add globaltimer-based timing backend for autotuner profi…

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

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

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<uint64_t*>(data_ptr);
+            auto stream_ptr = nb::cast<int64_t>(py_stream.attr("cuda_stream"));
+            cudaStream_t stream = reinterpret_cast<cudaStream_t>(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); },

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):

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)