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
• In terms of absolute performance, interleaving and sequential are close to each other in each "band"

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

[Micky774]

Thanks for the analysis!

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

Indeed, this is how ASV operates for exactly the goal of maximizing reproducibility and stability. Like you said however, overhead makes it infeasible for TE -- specifically the loading time for our .so objects means import takes too long and each process incurs a flat overhead in seconds that would make benchmarking take hours/days.

• adding additional measurement options to interleaving does not seem to meaningfully improve its results

Makes sense, since interleaving largely accomplishes e.g. cold caching on its own. Not sure of the others you included, but I suspect it's just the case that interleaving is a strict "superset" of other isolation mechanisms you've tried.

• interleaving adds a bit of within-run noise, as expected, but only a relatively small amount

Interestingly enough that's most likely a sign of it working in a sense, since it's trading bias for variance.

Questions that aren't answered yet:

• Is FPR the right metric to evaluate BM?

FPR and power are generally the metrics used for validating statistical methods, however in this case I wanted to hone in FPR for a specific invariance: identical kernels should result in identical measurements, implying an FPR of exactly $\alpha$.

Edit: I also ran a simulated analysis of power using the same dataset as earlier for the FPR investigation, except artificially scaling the "candidate" values by a delta percentage to mimic regressions. I marked the FPR (corresponding to the asymptotic limit as $\delta\rightarrow 0$) as dotted lines. Here are the power curves at the experimental configs (i.e. n=12, inner=20):

These curves will look different for different choices of $\alpha$ and n, so we can use the curves as a way to heuristically select a value of n -- i.e. determine a target power for a specific regression percent, such as "A 90% chance to catch a 5% regression" and scan choices of n until we find one that arrives at those odds.

Note

The higher power of the dev implementation (stats) might likely be due to the higher sensitivity to noise, and not necessarily a rigorous detection based on underlying kernel measurement changes. There is an intrinsic tradeoff between $\alpha$/FPR and power. The highest power test is one that says "yes" to everything all the time.

• Is BM the right way to measure regressions (e.g., compared to an approach like FlyDSL)?

Not sure what FlyDSL does, but in general BM is one of the best non-parametric tests for determining whether there is an order difference in samples from two distributions. It's essentially a more robust Mann–Whitney U test. With that being said, even a general permutation test results in the same outcomes for the empirical studies I've conducted, so as far as statistical significance tests go, there's not too much to worry about there IMO. Again, not sure what FlyDSL does so maybe we could discuss that separately.

• Are the downsides of interleaving/BM (non-trivial driver changes, memory overhead, diverging from comparable GPU/ML projects) worth the benefits?

Downsides include:

• Requiring AOT knowledge on sample order (this is free for a driver-style benchmark solution, less so for a more modular solution).
• Memory overhead proportional to the interleaving group size used, since if you interleave M configs, you need to be able to host M instances of the setup in memory simultaneously
• I don't view divergence from other projects as a significant downside assuming the measurements aren't affected -- which from your experiments seems to be the case, hence I don't worry much.

Upsides include:

• Preserving a critical invariance (same kernel ==> same measurements)
• Improving robustness of regression tracking
• Improving confidence in observed performance gains
• Increases functional power of the tests to determine true regressions and gains (since it is less sensitive to noise-induced bias)
• Improving statistical quality of results (really a rehashing of the above, but worth considering for communication purposes)

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

• also expose the other measurement options (rotating buffers, cache flushing, ...)

Agreed

• Not implement these directly within TE, but in some external package

I'm not sure that's necessary. We don't intend to make it generic, so I'm not quite sure what the benefit therein is. Keeping it as a part of TE will make it easier for folks to contribute to it, and will be consistent with how other open-source projects tend to maintain performance benchmarking utilities.

Recommended Next Steps

The goal isn't perfect FPR nor perfect power. I think we should think carefully about the following:

1. What percent of the time can we tolerate getting a false-positive "change" in kernel performance across identical kernels (FPR)
2. To what percent should we pick up a regression e.g. down to a >5% regression, or >1%, or >10% (threshold)
3. What percent of the time should we be able to detect a true regression at our selected threshold (power)

With these in mind we can establish the feasibility of different methods at delivering on those targets, and even determine explicitly things like the required sample size. Maybe we're okay without interleaving. Maybe we need process isolation. Maybe we are fine with a few samples, or maybe we need many samples. The point is that these should be determined by our goals, whatever they are.

Ideally, we have FPR = $\alpha$ but in practice that seems infeasible. I think "As low as computationally reasonable" is my personal stance there -- so something like interleave + realloc from Matthias's analysis. I think a 5% regression threshold is reasonable. I believe we should aim for a power of 90-95% for CI regression testing, and be okay with a power of 80-90% for prototyping and local benchmarking.