Background / Symptom
When evaluating kernels produced/optimized with Atrex, we consistently observe:
The exact same code (byte-identical, same hash) yields noticeably different scores across evaluation runs.
The jitter is not uniform — it is concentrated on large GEMM shapes:
- Large-N / large-M GEMMs: SOL score swings of roughly ~10%–20% across repeated runs of identical code.
- Small shapes (small M / small N): very stable, typically within ±1%.
In other words, the instability comes almost entirely from a few long-running, tiling-sensitive large GEMMs rather than from all kernels.
Suspected root cause
We believe it stems from two compounding factors:
- Non-deterministic autotune tiling selection — for the same shape, autotune may pick a different "best" config (block sizes,
num_warps, num_stages, …) on each run. This is most likely when the top candidates are close together and measurement noise is high, which is exactly the case for large GEMMs.
- Clock / thermal throttling — large GEMMs run long enough to be affected by GPU frequency and thermal throttling, which amplifies the selection instability above.
Factor 2 is environmental and hard to eliminate. But factor 1 (autotune selection non-determinism) can be addressed on the kernel side: once a good config has been found for a given shape, that result should be fixed and reused for the same shape on subsequent runs, instead of being re-autotuned every time.
Request
Provide a way to make an already-tuned best config a deterministic, reproducible result:
- For the same shape in the same environment, the kernel should produce a deterministic execution config rather than potentially re-selecting a different one via autotune each run.
- This lets good results reproduce reliably, instead of depending on submitting multiple times to "get lucky" and taking the max.
I don't have a strong preference on the exact mechanism (e.g. baking the best config into the kernel, removing/narrowing @triton.autotune, or introducing a per-shape config cache that is persisted and reused) — I'll leave the best approach to the maintainers.
Impact
- Results are not reproducible, making it hard to tell whether a change is a real improvement or just evaluation noise.
- Aggregating "each kernel's historical best score" systematically overstates achievable performance versus what actually reproduces.
- Unnecessary repeated autotuning overhead on every run.
Background / Symptom
When evaluating kernels produced/optimized with Atrex, we consistently observe:
The exact same code (byte-identical, same hash) yields noticeably different scores across evaluation runs.
The jitter is not uniform — it is concentrated on large GEMM shapes:
In other words, the instability comes almost entirely from a few long-running, tiling-sensitive large GEMMs rather than from all kernels.
Suspected root cause
We believe it stems from two compounding factors:
num_warps,num_stages, …) on each run. This is most likely when the top candidates are close together and measurement noise is high, which is exactly the case for large GEMMs.Factor 2 is environmental and hard to eliminate. But factor 1 (autotune selection non-determinism) can be addressed on the kernel side: once a good config has been found for a given shape, that result should be fixed and reused for the same shape on subsequent runs, instead of being re-autotuned every time.
Request
Provide a way to make an already-tuned best config a deterministic, reproducible result:
I don't have a strong preference on the exact mechanism (e.g. baking the best config into the kernel, removing/narrowing
@triton.autotune, or introducing a per-shape config cache that is persisted and reused) — I'll leave the best approach to the maintainers.Impact