The tile shape is part of the number

Two GPU kernels can be deterministic, accept the same tensors, implement the same matrix multiplication, and still disagree on the last few bits. The irritating part is how ordinary the cause can be: a tile shape, a split-K choice, or a compiler lowering changed the tree of floating-point additions.

A paper submitted September 10, 2026, Taming Bitwise Behavior in GPU Kernels with Tensor Core, makes that tree an explicit compiler problem. Ziteng Yang, Nicholas J. Riasanovsky, Warren Deng, and Vivek Sarkar reconstruct cuBLAS arithmetic from the outside, add a balanced reduction mode to Triton, and build static checkers that partition compiled kernels by exact bit behavior before an autotuner starts timing them.

That last part is the one I keep coming back to. Autotuning normally treats configuration knobs as a pile of equivalent ways to compute one function, then asks which pile member is fastest. Floating point makes that story slightly fake. Some knobs are performance choices. Some knobs quietly choose a different arithmetic program.

Deterministic is a smaller promise#

Determinism answers a run-to-run question: if I execute this implementation again under the same promised conditions, do I get the same answer? Bitwise equivalence between implementations asks whether two different kernels produce the same bit pattern for the same input. One does not imply the other.

NVIDIA's current cuBLAS documentation makes the scope concrete. Within one toolkit version, the same GPU architecture, and the same number of SMs, its routines are designed to repeat bitwise; that guarantee does not extend across toolkit versions, and multi-stream execution needs extra conditions. PyTorch is similarly careful: its May 2026 reproducibility note says complete reproducibility is not guaranteed across releases or platforms, while deterministic mode chooses repeatable operators where it can.

None of this is a GPU quirk. IEEE-style floating-point addition does not associate. I ran this tiny NumPy example while checking the article; both expressions contain the same three values, but the parentheses are enough to move the result from 3.14 to 0.0 in float32.

import numpy as npa = np.float32(1e20)b = np.float32(-1e20)c = np.float32(3.14)left = np.float32(np.float32(a + b) + c)right = np.float32(a + np.float32(b + c))print(left)   # 3.14print(right)  # 0.0

A GPU reduction is that example with thousands of leaves and a compiler deciding the parentheses. A thread first combines values in registers, lanes combine partials through warp operations, and warps may combine again through shared memory. Change which element belongs to which lane and the tree can change even when the source-level reduction still reads like the same sum.

The performance knobs are not one species#

Triton makes this tension easy to see because its official matrix multiplication tutorial puts a list of triton.Config values behind @triton.autotune. Block sizes, num_warps, and num_stages are compiled into separate kernels, timed, and cached by shape. From a throughput view they are all candidates. From a bitwise view they are not interchangeable.

The pipeline-depth row is especially useful because it prevents the rule from collapsing into “every optimization changes the bits.” A deeper software pipeline can issue future loads earlier without moving the additions those values eventually feed. The paper makes a similar conditional claim for warp specialization: producer warps can be split from consumer warps without changing the arithmetic, provided the compiler does not repartition the tile over the remaining compute warps.

Split-K sits on the other side. Dividing the contracted K dimension lets separate blocks compute partial output matrices and merge them later. That is excellent parallel work when K is large. It also inserts a new set of cuts into the reduction tree, so where the cuts fall, the precision of the partials, and the merge order become part of the numerical behavior.

A descriptor instead of a vibe#

The paper packages the bit-relevant choices for matrix multiplication into GEMMDesc. The point is subtraction as much as addition: the descriptor records things that decide reduction order and rounding, and leaves out knobs that only decide speed. Two implementations with the same descriptor are making a much stronger claim than “both are GEMM.” They are claiming the same relevant floating-point computation.

This also gives the authors a way to study cuBLAS without its source. They feed carefully chosen numerical probes through the closed library and infer the hidden grouping from the answers. Large positive and negative values can cancel under one grouping but swallow a small residual under another, so the output leaks where the reduction was cut. Repeating those probes over shapes and dtypes turns a black-box kernel selector into a descriptor the Triton implementation can try to match.

Their reported scope matters. For fp16, bf16, and fp8 e4m3 GEMMs on H100, GB200, and GB300, the reconstructed Triton family bit-matched the tested cuBLAS behavior, with the paper explicitly excluding a cuBLAS split-K tail bug from its 100% statement. That is a measured result over the studied configurations, not a promise that every future cuBLAS version has now been reverse-engineered forever.

Make the compiler build the tree on purpose#

Reverse-engineering a vendor library is one half of the paper. The other half asks what a compiler should do when a stable order is the requirement rather than compatibility with someone else's order. The authors add an “inner tree” mode to Triton's backend that lowers reductions into a balanced dependence tree.

Balanced matters for more than neat diagrams. A left fold over n values has a dependence chain of n - 1 additions. A balanced tree has depth around log2(n). So fixing the association does not inherently mean choosing the least parallel tree. The actual performance problem is layout: the chosen tree has to line up with where values live in lanes and registers, or the GPU pays for shuffles and awkward movement.

Their data-layout optimization attacks that movement. Across 27 measured kernels on GB300 and H100, 19 land within 10% of the free-order mode. Five of the ten H100 cases in the paper are faster than the free-order baseline after the layout work. There are still losses — one column-sum epilogue is reported as much worse — but the result is more interesting than a blanket “reproducibility costs 20%.” The cost depends on whether the compiler is allowed to optimize around the fixed semantic constraint.

Autotune inside an equivalence class#

Once a compiler can say which configurations are bit-equivalent, autotuning gets a cleaner contract. Pick a reference behavior, throw away candidates that compute different bits, then benchmark the survivors. The shape below is conceptual pseudocode rather than a Triton API, but it is the programming model the paper is pushing toward.

# Conceptual pseudocode, not a current Triton API.compiled = [compile(config) for config in configs]reference = compiled[0]same_bits = [    kernel for kernel in compiled    if bitwise_equivalent(reference, kernel)]winner = min(same_bits, key=benchmark_time)

The clever bit is that the equivalence test does not rerun thousands of random tensors and hope. The paper builds static checkers over compiled NVIDIA PTX and AMD GCN. For a symbolic thread, the checker tracks a map from registers to expression nodes. Loads create data nodes; arithmetic instructions create operation nodes; lane exchanges and stores connect the pieces. The resulting dependence structure can be compared between kernels.

It is intentionally one-sided. A sound checker is allowed to fail to prove that two equivalent kernels are equal; that only leaves some speed on the table. It must not merge two kernels that can disagree. In the paper's evaluation, every pair certified equivalent across 51,152 GB300 configurations from 47 Triton kernels and five formats matched bytes on the generated test inputs. The checker exactly recovers the partition for their GEMM family, while some reductions and normalizations are over-split by up to 2.5x.

That is a useful compiler trade. False separation costs search space. False equivalence breaks the promise. The checker is designed to be conservative in the direction you can debug.

The configuration is part of the semantics now#

Most programmers do not need byte-identical floating-point answers across kernel implementations. Higher precision, tolerances, and ordinary numerical analysis are usually the better tools. Even PyTorch's current numerical-accuracy note says mathematically identical floating-point computations are not generally promised to be bitwise identical across platforms and releases. Chasing exact bits everywhere would turn an optimization problem into paperwork.

But some systems really do need a stable numerical identity: regression tests, batch-invariant serving, matching a training and inference engine, or moving a kernel between a vendor library and generated code without changing the model's arithmetic underneath it. In those cases, “deterministic” is too coarse and “same formula” is too weak.

The limitation is also where this work gets interesting. The static checker sees PTX or AMD GCN, not the final machine code after every possible driver decision. Its strongest evaluation is around reductions and GEMM-like kernels, and exact equivalence can still be sensitive to instruction choice, partial precision, fusion, and rounding. A later compiler or library version can move those choices again.

Still, the paper gives the autotuner a concept it was missing. A tile shape is not always just a faster or slower way to run the same program. Sometimes it is part of the number the program computes. Once the compiler admits that, it can optimize aggressively inside the right equivalence class instead of pretending the class never existed.