fix(tensilelite): Reduce unaccounted overhead of ScopedTimer and timing_context (ROCm#6043)

## Motivation
Each `ScopedTimer` / `timing_context` call performed string formatting
and I/O inline, adding significant overhead that inflated parent timers,
especially for bad for high call count but low execution time
operations.

Additionally, `gpu_kernel_execution` and `cpu_reference_gemm` were
placed at the wrong hierarchy level, making the analysis output
misleading.

Follow-up to ROCm#6259.

## Technical Details

Overhead reduction
- **Deferred I/O** (`TimingInstrumentation.hpp`,
`TimingInstrumentation.py`): Replaced inline
`writeLine`/`_timing_logger.info` calls with lightweight buffer appends
during measurement. All formatting and I/O now happens in a single
`flushTimingBuffer()` call after measurement is complete.
- **Overhead calibration** (`TimingInstrumentation.hpp`, `main.cpp`):
Added `calibrateTimingOverhead()` which measures per-call `ScopedTimer`
overhead at startup (warmup + 100k iterations). The calibrated value is
emitted as a `timing_overhead` record and used by `analyze_timing.py` to
subtract instrumentation overhead from parent totals.
- **Buffer pre-allocation** (`TimingInstrumentation.hpp`):
`initTimingBuffer()` reserves 2.5M entries upfront to avoid reallocation
stalls during benchmarking.
- **Analysis improvements** (`analyze_timing.py`): Added per-call
overhead subtraction based on descendant invocation counts, and a new
"Timing Overhead" phase group for `timing_overhead` and
`flush_timing_buffer`.

Fixes
- **Hierarchy fixes** (`main.cpp`, `BenchmarkTimer.cpp`,
`analyze_timing.py`):
- Moved `gpu_kernel_execution` under `benchmark_runs` (was a sibling,
incorrectly double-counted).
- Moved `cpu_data_init` and `cpu_reference_gemm` under `pre_problem`
(where they actually execute).

## Test Plan
Run timing instrumentation end-to-end and verify that parent totals are
consistent with child sums after overhead adjustment.

## Test Result
Timing output shows correct parent-child relationships and
overhead-adjusted totals.

---------

Signed-off-by: Alex Vasile <48962821+Alex-Vasile@users.noreply.github.com>

Author: Alex-Vasile

projects/hipblaslt/tensilelite/Tensile/Common/TimingInstrumentation.py

@@ -37,16 +37,37 @@
     _timing_logger.setLevel(logging.INFO)
     _timing_logger.propagate = False
 
+# Deferred I/O buffer: list of (category, duration_ms) tuples.
+# All formatting and I/O happens in flush_timing_buffer().
+# Raw timings are stored without overhead adjustment; the analysis script
+# (analyze_timing.py) handles overhead subtraction using the hierarchy and
+# the timing_overhead record emitted by the C++ client.
+_timing_buffer = []
+
 
 @contextmanager
 def timing_context(category_name):
-    """Context manager for timing instrumentation."""
+    """Context manager for timing instrumentation.
+
+    Records raw wall-clock time with no overhead subtraction.  Python-side
+    overhead (context-manager protocol, time.time_ns, dict lookup) is not
+    tracked because there are only ~a dozen calls per run — the overhead
+    is negligible relative to the seconds-scale measurements.  C++ overhead
+    is tracked separately via a calibrated timing_overhead record.
+    """
     if globalParameters.get("TimingInstrumentation", False):
         start = time.time_ns()
         try:
             yield
         finally:
-            elapsed_ms = (time.time_ns() - start) / 1_000_000
-            _timing_logger.info(f"TIMING:{category_name}:{elapsed_ms:.3f}")
+            elapsed_ns = time.time_ns() - start
+            _timing_buffer.append((category_name, elapsed_ns / 1_000_000))
     else:
         yield
+
+
+def flush_timing_buffer():
+    """Write all buffered timing records and reset."""
+    for category, duration_ms in _timing_buffer:
+        _timing_logger.info(f"TIMING:{category}:{duration_ms:.3f}")
+    _timing_buffer.clear()

projects/hipblaslt/tensilelite/Tensile/Tensile.py

@@ -45,7 +45,7 @@
 from Tensile.Common.Capabilities import makeIsaInfoMap
 from Tensile.Common.GlobalParameters import globalParameters, assignGlobalParameters, \
                                             restoreDefaultGlobalParameters
-from Tensile.Common.TimingInstrumentation import timing_context
+from Tensile.Common.TimingInstrumentation import timing_context, flush_timing_buffer
 from Tensile.Toolchain.Assembly import AssemblyToolchain, makeAssemblyToolchain
 from Tensile.Toolchain.Source import SourceToolchain, makeSourceToolchain
 from Tensile.Toolchain.Validators import validateToolchain, ToolchainDefaults
@@ -125,6 +125,7 @@ def executeStepsInConfig(
                 buildOnly,
                 solutionPoolFiles,
             )
+        flush_timing_buffer()
         print1("")
 
     if buildOnly:
@@ -155,6 +156,7 @@ def executeStepsInConfig(
                     debugConfig.printIndexAssignmentInfo,
                     isaInfoMap,
                 )
+            flush_timing_buffer()
             print1("")
         else:
             print1("# LibraryLogic already done.")
@@ -178,6 +180,7 @@ def executeStepsInConfig(
                 deviceId,
                 gfxName,
             )
+        flush_timing_buffer()
         print1("")
 
 
@@ -682,6 +685,9 @@ def Tensile(userArgs):
         print("Overriding {0}={1}".format(key, value))
         globalParameters[key] = value
 
+    if "MaxFileName" in globalParameters or "MaxFileName" in config:
+        printWarning("MaxFileName is no longer configurable, it will be automatically set to 64")
+
     executeStepsInConfig(config, outputPath, asmToolchain, srcToolchain, isaInfoMap, cCompiler, debugConfig, device_id, prob_sol_map, buildOnly, solutionPoolFiles)
 
 def TensileConfigPath(*args):

projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/TimingInstrumentation/test_timing_instrumentation_char.py

@@ -12,7 +12,7 @@
 import pytest
 
 from Tensile.Common.GlobalParameters import globalParameters
-from Tensile.Common.TimingInstrumentation import timing_context
+from Tensile.Common.TimingInstrumentation import flush_timing_buffer, timing_context
 
 pytestmark = pytest.mark.unit
 
@@ -40,9 +40,10 @@ def test_timing_context_disabled_yields():
 
 
 def test_timing_context_enabled_times_and_logs():
-    # Flag on -> the timing path; body runs and a TIMING line is emitted. The
-    # "tensile.timing" logger has propagate=False, so attach a capture handler
-    # directly to it (caplog/capsys can't see it).
+    # Flag on -> the timing path; body runs and a TIMING line is buffered until
+    # flush_timing_buffer() emits it. The "tensile.timing" logger has
+    # propagate=False, so attach a capture handler directly to it (caplog/capsys
+    # can't see it).
     import logging
 
     messages = []
@@ -53,13 +54,17 @@ def emit(self, record):
 
     logger = logging.getLogger("tensile.timing")
     handler = _Capture()
+    flush_timing_buffer()
     logger.addHandler(handler)
     try:
         with _timing_flag(True):
             ran = []
             with timing_context("mycat"):
                 ran.append(1)
             assert ran == [1]
+            assert not any("TIMING:mycat:" in m for m in messages)
+            flush_timing_buffer()
     finally:
         logger.removeHandler(handler)
+        flush_timing_buffer()
     assert any("TIMING:mycat:" in m for m in messages)

projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_analyze_timing.py

@@ -250,7 +250,6 @@ def test_nested_timing_totals(self, nested_timing_log):
             + timings["code_object_loading"][0]
             + timings["cpu_data_init"][0]
             + timings["cpu_reference_gemm"][0]
-            + timings["gpu_kernel_execution"][0]
         )
         assert cpp_sum <= timings["python_client_execution"][0]