Interleaved Driver Benchmarking

[Micky774]

Description

Provides an alternative driver-based implementation of the micro benchmarking suite. This implementation includes cold-caching and interleaving as orthogonal features. This PR is structured as a clean addition for ease of review.

Each individual benchmark file is <60 LOC with configs handled in a central models.py for housing configs, and a single driver which parses the benchmarks to be run AOT allowing for arbitrary scheduling.

This PR includes the significance testing changes introduced in #614

This suite is designed for detecting performance regressions across a shape sweep, not just point-estimating a single kernel. Samples are collected round-robin across benchmarks and re-shuffled each round, so time-correlated GPU drift (thermal ramp, power/voltage fluctuation, a neighbor on a shared GPU) becomes shared variance instead of a systematic bias between benchmarks.

Type of change

• Documentation change (change only to the documentation, either a fix or a new content)
• Bug fix (non-breaking change which fixes an issue)
• New feature (non-breaking change which adds functionality)
• Breaking change (fix or feature that would cause existing functionality to not work as expected)
• Infra/Build change
• Code refactoring

Changes

Please list the changes introduced in this PR:

• Change A
• Change B

Checklist:

• I have read and followed the contributing guidelines
• The functionality is complete
• I have commented my code, particularly in hard-to-understand areas
• I have made corresponding changes to the documentation
• My changes generate no new warnings
• I have added tests that prove my fix is effective or that my feature works
• New and existing unit tests pass locally with my changes

[Micky774]

To demonstrate the significance of interleaving, I ran a head-to-head comparison against our current benchmark suite.

Methodology

---

Three "families" of benchmarks were evaluated across both suites: normalization, casting, and grouped GEMM. The dev suite and the PR suite were run with these three benchmarks a total of K=8 times. Then pairs were constructed from each run-set of a given family. That is to say, for the 8 normalization kernel benchmark results from the dev suite, we constructed $\binom 8 2$ = 28 run-pairs and compared them via the Brunner Munzel significance test with $\alpha=0.001$ to find how many false-positives were produced.

We consider them false positives due to being introduced by transient noise rather than genuine kernel differences. From the eyes of the test, they are "true" positives in that they are genuinely different, albeit due to noise. This distinction is subtle but worth noting. Note that there should be zero significant results since the runs are meant to be identical.

Both suites were tested with n=12 samples per run, however the amount of "inner" samples used to amortize dispatch latency differed, with dev being automatically determined as per default settings, and with this PR being set to as many as it would take to ensure at least 50ms of runtime. Higher "inner" samples further amortize dispatch time, asymptotically approaching kernel execution time, whereas smaller counts better amortize transient noise due to smaller contiguous runs contributing towards the same config.

While some configs differ between the benchmarks, resulting in different number of configs, this was normalized by calculating the False Positive Rate (FPR) rather than a direct count.

Results

---

The dev branch is colored RED as stats whereas this PR is colored BLUE as asv -- this is just a legacy naming convention and not reflective of any critical part of the suite.

Considerations

---

This was NOT performed under performance determinism mode or pure hardware isolation. Anecdotally, performance determinism mode helps bridge the gap, but interleaving still provides a superior FPR reduction. True hardware isolation is difficult to guarantee in workload environments. Either way, the core message here is that interleaving meaningfully mitigates / improves a genuine existing problem.

[Micky774]

On a newer run with n=12 and inner=20 for both, when letting the PR branch use min(inner runs) as its statistic instead of the mean, we get the following FPR graphs:

[matthiasdiener]

Running a similar methodology, but restricted to GEMMs (to be able to run hipblaslt-bench) for only 3 GEMMs (so total runtime remains manageable) shows the following results:

A few initial points to note:

• Surprisingly, the only way to get an FPR of ~0% is to run samples with a fresh process (whether via hipblaslt-bench or a modified Python driver), other methods (like interleaving/rotating buffers/cache flushing/idle gap/realloc) do not seem to achieve this. Fresh processes cause a significant overhead in the Python driver, however
• adding additional measurement options to interleaving does not seem to meaningfully improve its results
• interleaving adds a bit of within-run noise, as expected, but only a relatively small amount

Overall, I think interleaving/BM are interesting features, but their practical use probably needs a bit more analysis.

Questions that aren't answered yet:

• Is FPR the right metric to evaluate BM?
• Is BM the right way to measure regressions (e.g., compared to an approach like FlyDSL)?
• Are the downsides of interleaving/BM (non-trivial driver changes, memory overhead, diverging from comparable GPU/ML projects) worth the benefits?

I'd recommend that if we include interleaving/BM, we:

• also expose the other measurement options (rotating buffers, cache flushing, ...)
• Not implement these directly within TE, but in some external package