Paper 01Draft
GPU Mechanism
Same Architecture, Same SM Count Is a Sufficient Condition for One Kernel Family and Neither Necessary Nor Sufficient for Another: Measuring, Naming and Removing the Device Dependence of GPU Training Bytes
In this paper52 min read
Abstract
cuBLAS scopes its bit-wise reproducibility to “the same architecture and the same number of
SMs”. We measured that clause against PyTorch fp32 and bf16 training on eleven NVIDIA
instances spanning seven (architecture, SM count) classes under one PyTorch build, then bought
the two pairs that broke it again with the instruments needed to say why. The clause holds
exactly for fp32: identical bytes across four H100 PCIe cards from three providers, two host
CPU vendors and three driver builds, and across an L40 and an L40S; different bytes across
any two SM counts, including the MIG slices of one card. It is neither necessary nor
sufficient for bf16: an A10 (72 SMs) and an A10G (80) agree, an L40 and an L40S (both 142)
disagree. We name both keys. The fp32 key is cuBLASLt’s smCountTarget: asked for the same
fifteen GEMM shapes at 114, 46 and 14 SMs on one H100, the heuristic changes tile and split-K
for ten and eleven of them, and the trajectories split at the first parameter update. The
bf16 key is a per-product table: the two Ada cards’ traces choose identical algorithms for all
fifteen fp32 shapes and different ones for twelve of fifteen bf16 shapes; locking both cards to
the same clock changes nothing; changing the clock CUDA reports by 52 % changes nothing; and
NVIDIA’s own standalone heuristic library, whose descriptors for the two products differ in
the peak clock alone, returns identical configurations for both. A second, non-cuBLAS root —
ATen’s reduction kernel splitting bias-gradient sums across blocks by SM count — was found by
a per-op trace on the 14-SM slice, and the device random generator’s dependence on SM count
was traced to one std::min and reproduced bit for bit on a CPU for three cards. Two fixes
follow. An exact-integer (Ozaki-scheme) GEMM makes the fp32 program byte-identical across
Ampere, Ada and Hopper and across 14 to 142 SMs at 12–16× the step time; forcing the
reference card’s own cuBLASLt algorithm choices through the public algorithm API makes it
byte-identical across an L40 and an A100 and across the A100’s 14- and 42-SM MIG slices (a
slice’s heuristic is the whole card’s asked for its SM count: 30 of 30 shapes on the A100,
23–26 of 30 on an H100). One configuration accepted by every card — algoId 20 without
split-K — gives one fp32 trajectory across Ampere, Ada and Hopper and five SM counts (A100,
L40, L40S, H100 and the H100’s 46- and 14-SM slices); a reference card’s table of heuristic
choices is portable within an architecture and only partly across (6 of 22 entries admissible
on Hopper), and a card’s own learned table costs 0.94–0.99× the reference. Implemented
in C++ at the CUDA dispatch key, the forcing is bit-identical to the Python instrument and
runs at 0.96–1.07× the cuBLAS reference for the pinned program and 0.97–1.23× for the stock
one (two clean runs each; the reference itself moves 7–15 % between its own runs), where the
Python dispatch mode had cost 2.5–2.8×. The bf16 key is then neutralised, not only
named: with the L40’s learned bf16 algorithm table forced on the L40S, bf16 training is
byte-identical across the two products with either attention backend through 300 steps,
while the stock programs split at the first update in the same sessions; the two cards’
heuristics disagree on 11 of 19 bf16 GEMM shapes in the fields that name the kernel and its
summation order (17 of 19 counting the CTA-swizzle flag). The deterministic
rung itself costs 1.05–1.19×, not the 2–5× of folklore. Every number in this paper is rendered
from a result file; every prediction was registered before its instance existed.
1. Introduction
A verifier who replays a training segment and compares checkpoints byte for byte needs a
promise about hardware: on which machines will the same program produce the same bytes?
For CPUs we answered this earlier by measurement (companion paper, ship1.md): with the
right envelope, six x86 microarchitectures of both vendors agree. For GPUs the only public
promise is cuBLAS’s reproducibility clause — same architecture, same number of SMs, same
toolkit — and cuDNN’s disclaimer. Proof-of-Learning [PoL] and its attacks [Zhang21, Fang22]
live in the tolerance δ that a verifier must allow because that promise was assumed too
weak to make δ zero; the two published routes to exact cross-GPU replication change the
training procedure (rounding logs [SAB24]; replaced operators [RepOps]).
This paper asks the cheaper question: what does the promise actually cover, what does it
key on, and can the keys be removed without rewriting the training loop? We proceed as an
experimentalist would with a black box: measure the map (which pairs of devices agree, per
kernel family), locate the split (which step, which op), read the vendor’s own dispatch
record (cuBLASLt’s trace, which prints the algorithm it chose and the SM count it was asked
to target), vary the candidate keys one at a time (running clock, reported clock, SM count
through MIG and through smCountTarget), and finally intervene (replace the GEMM by an exact
one; force the algorithm choice). Predictions were registered before each instance was
bought (python/recipe_checkpoints/PREREGISTRATION.md §6–§8) and are settled there.
Contributions.
- A measured map of GPU training reproducibility over seven (architecture, SM count) classes and two precisions under one PyTorch build, showing that the fp32 and bf16 kernel families induce different, incomparable equivalence relations over the same silicon (§4.1).
- The fp32 key named in the vendor’s words:
smCountTarget→ tile and split-K → summation order, with the heuristic’s complete SM-count map obtained from a single card (§4.2, §4.6). - The bf16 key named by elimination — a per-product table, with the clock excluded three ways, including by changing the clock the runtime reports (§4.3).
- A second root outside cuBLAS (ATen’s reduction split) and the device generator’s line, both reproduced: the first by a per-op trace on a 14-SM slice, the second bit-for-bit on a CPU for three cards (§4.4, §4.5).
- Two removals: an exact GEMM (device-independent bits across three architectures and five SM counts, 12–16×) and forced algorithms through the public API — one fp32 trajectory across Ampere, Ada and Hopper and five SM counts on a single configuration, a learned table portable within an architecture, and, implemented in C++ at the CUDA dispatch key, at the cuBLAS reference’s price — 0.96–1.23× across two cards, two program variants and paired runs, against a reference whose own runs differ by 7–15 % (§5.2–5.3).
- The bf16 key neutralised: the L40’s bf16 algorithm table forced on the L40S makes bf16 training byte-identical across the two products whose heuristics disagree on 11 of 19 shapes (kernel-naming fields; 17 of 19 with CTA swizzle), with either attention backend, while the stock programs split at step 1 in the same sessions (§5.4).
- The price of the deterministic rung across seven classes: 1.05–1.19× (§4.7).
- The mechanism read on a report we did not make: pytorch/pytorch#105726’s fp16 A100 ≠ 3090 = L4 divergence reproduced bit for bit on the reporter’s inputs, keyed to the heuristic’s stage count (K-slab 64 vs 32) as moved by the SM count, and removed by a configuration both cards accept — either card produces either output when told which (§5.5).
2. Background and related work
cuBLAS’s clause. “By design, all cuBLAS API routines from a given toolkit version,
generate the same bit-wise results at every run when executed on GPUs with the same
architecture and the same number of SMs” [cuBLAS-doc §2.1.4]; the guarantee is withdrawn
across toolkit versions, under multiple streams sharing a workspace, and under fixed-point
emulation. cuBLASLt exposes the heuristic and the algorithm selection this paper uses:
cublasLtMatmulAlgoGetHeuristic, CUBLASLT_MATMUL_DESC_SM_COUNT_TARGET,
cublasLtMatmulAlgoInit / cublasLtMatmulAlgoConfigSetAttribute / cublasLtMatmulAlgoCheck
/ cublasLtMatmul, and a logger (CUBLASLT_LOG_LEVEL) that prints the chosen algoId,
tile, stages, reduction scheme and split-K count for every call [cuBLAS-doc §3.1.3].
PyTorch’s determinism note [PT-repro] documents use_deterministic_algorithms,
CUBLAS_WORKSPACE_CONFIG, and per-op caveats; it does not state that the CUDA generator’s
output depends on the device’s SM count (we filed that: [PT-issue]).
Verifiable training. Proof-of-Learning [PoL] verifies a training transcript by replaying segments and accepting if the replayed weights land within a tolerance δ calibrated to hardware noise; the published attacks [Zhang21, Fang22] construct transcripts inside δ. Exact replication across GPUs has been achieved by altering the procedure — Srivastava, Arora and Boneh train in higher precision and log rounding decisions [SAB24]; Gensyn’s RepOps replaces CUDA operators to fix operation order [RepOps]. We instead measure the stock kernels’ promise and then remove its keys through public APIs.
Invariance in inference. He et al. name batch invariance — kernels whose output does
not depend on the batch size — as the missing property behind non-reproducible LLM inference
and build batch-invariant RMSNorm, matmul and attention [He25]. Zhang et al. extend the
property to the tensor-parallel size with tree-based invariant kernels (TBIK): matrix
multiplication and reduction primitives whose intra- and inter-GPU reduction orders follow one
hierarchical binary tree, so the bits do not depend on how many ranks the computation is cut
into [Zhang25], after the same group measured that batch size, GPU count and GPU version change
greedy decoding outputs [Yuan25]. Veit’s kernel contracts supply the specification language
for such properties — identifier, scope, precondition, postcondition, tolerance, reference
oracle, measurement protocol, violation signature — and twelve contract classes spanning
precision, ordering, compiler-induced and exceptional-value failures [Veit26]. Our result is a
companion statement for training, and it sits beside TBIK’s: rank count there, SM count here —
kernel families differ in which device property their output depends on, and the map is per
family. In Veit’s terms, §4 gives the ordering failure class two named violation signatures
(smCountTarget, the per-product bf16 table) and §5 two reference-conforming
implementations.
Exact GEMM by error-free transformation. The Ozaki scheme [Ozaki12] splits floating-point operands into slices whose products accumulate exactly; GEMMul8 [Uchino24] carries it to integer tensor cores. We use it as an instrument: integer accumulation has no order, so a GEMM built on it cannot depend on the SM count.
Philox. The counter-based generator [Salmon11] is what cuRAND and PyTorch’s CUDA distributions use; its outputs are deterministic functions of (key, counter). What varies by device is the mapping of tensor elements to counters (§4.5).
3. Method
Programs. The reference loop of the companion paper (TinyTransformer: 4 layers, d = 256,
4 heads, block 128, 0.4 M parameters, byte-level Shakespeare, batch 16) in two variants:
pinned (fixed-order embedding backward, an AdamW written out with torch.sqrt, loss as
log-softmax + gather) and stock (nn.Embedding, torch.optim.AdamW, CrossEntropyLoss).
Parameters are initialised on the CPU from a fixed seed and copied, so every device starts
from identical bytes; a later divergence is training arithmetic.
Envelope rungs. G0: seeds only. G1: torch.use_deterministic_algorithms(True),
CUBLAS_WORKSPACE_CONFIG=:4096:8, cuDNN deterministic. G2: G1 + TF32 off + math attention.
Dtype fp32 or bf16 autocast; attention backend flash, math or memory-efficient where the
dtype allows.
Measurement. The full resumable state (parameters and optimizer state) is digested with the project’s canonical encoding at steps 0, 1, 2, 3, 5, 10, 20, 50, 100, 200, 300 (the campaign used 0/100/200/300). Runs are separate processes; two or three per cell. Per-step wall time is the median over steps 10 onwards after a device synchronise.
Instruments (added for the mechanism pass). (i) The kernel record: torch.profiler’s
ordered list of every kernel launched in one step, hashed. (ii) cuBLASLt’s own trace
(CUBLASLT_LOG_LEVEL=5): the chosen algorithm per GEMM, parsed per (shape, layout, compute
type). (iii) A per-op trace at step 0: digests of every module output, the loss, every
gradient and every parameter after the update. (iv) Device attributes as CUDA reports them
(cudaDeviceGetAttribute) and as nvidia-smi reports them (running clocks, power,
temperature). (v) The heuristic queried directly through ctypes, with and without
CUBLASLT_MATMUL_DESC_SM_COUNT_TARGET. (vi) Clock control: nvidia-smi -lgc (running
clock) and -ac (applications clock, which changes cudaDevAttrClockRate). (vii) A device
RNG sweep: rand/randn at 2^12 … 2^22 from a seeded device generator, bytes hashed.
Interventions. (a) ExactGEMM: a TorchDispatchMode that routes every fp32 mm,
addmm, bmm, baddbmm (forward and autograd backward) through an Ozaki-scheme GEMM —
four 7-bit integer slices per operand scaled by powers of two from frexp, slice products
by an exact integer GEMM (torch._int_mm, or an fp32/fp64 GEMM where every partial sum is an
exact integer), fixed-order fp64 recombination, one rounding; ozaki2 also routes reductions
over all-but-the-last dimension (bias gradients). (b) PinnedGEMM: the same interception,
but each GEMM runs through cublasLtMatmul with an algorithm forced by
cublasLtMatmulAlgoInit + ConfigSetAttribute — either one configuration for every shape or
a per-shape table learned from a reference card’s heuristic.
Hardware. NVIDIA Brev instances, one at a time; each instance backed up in full and
deleted; worst-case spend at hourly rounding is recorded per instance. The provider handed
out the same physical L40 (identical GPU UUID) for the l40-3 and l40-4 sessions; every
other row is a distinct card.
| instance | product | arch (CC) | SMs | threads/SM | driver | PCI id | host |
|---|---|---|---|---|---|---|---|
| a10-1 | NVIDIA A10 | 8.6 | 72 | — | 570.148.08 | — | Intel(R) Xeon(R) Platinum 8358 CPU |
| a100-2 | NVIDIA A100-SXM4-80GB | 8.0 | 108 | — | 580.126.09 | — | Intel(R) Xeon(R) Platinum 8462Y+ |
| a10g-1 | NVIDIA A10G | 8.6 | 80 | — | 595.91.07 | — | AMD EPYC 7R32 |
| h100-1 | NVIDIA H100 PCIe | 9.0 | 114 | — | 570.195.03 | — | AMD EPYC 9554 64-Core Processor |
| h100-2 | NVIDIA H100 PCIe | 9.0 | 114 | — | 570.195.03 | — | AMD EPYC 9554 64-Core Processor |
| h100-3 | NVIDIA H100 PCIe | 9.0 | 114 | — | 570.148.08 | — | Intel(R) Xeon(R) Platinum 8480+ |
| l4-1 | NVIDIA L4 | 8.9 | 58 | — | 580.178.04 | — | Intel(R) Xeon(R) CPU @ 2.20GHz |
| l40-1 | NVIDIA L40 | 8.9 | 142 | — | 580.126.09 | — | Intel(R) Xeon(R) Platinum 8352Y CP |
| l40s-1 | NVIDIA L40S | 8.9 | 142 | — | 580.126.09 | — | Intel(R) Xeon(R) Platinum 8462Y+ |
| t4-3 | Tesla T4 | 7.5 | 40 | — | 580.178.04 | — | Intel(R) Xeon(R) CPU @ 2.30GHz |
| t4g-1 | NVIDIA T4G | 7.5 | 40 | — | 595.91.07 | — | — |
| a10g-2 | NVIDIA A10G | 8.6 | 80 | 1536 | 595.91.07 | 0x223710DE/0x152F10DE | AMD EPYC 7R32 |
| h100-4 | NVIDIA H100 PCIe | 9.0 | 114 | 2048 | 580.126.09 | 0x233110DE/0x162610DE | Intel(R) Xeon(R) Platinum 8352Y CP |
| l40-2 | NVIDIA L40 | 8.9 | 142 | 1536 | 580.126.09 | 0x26B510DE/0x169D10DE | AMD EPYC 7773X 64-Core Processor |
| l40s-2 | NVIDIA L40S | 8.9 | 142 | 1536 | 580.126.09 | 0x26B910DE/0x185110DE | Intel(R) Xeon(R) Platinum 8462Y+ |
| a100-3 | NVIDIA A100-SXM4-80GB | 8.0 | 108 | 2048 | 580.126.09 | 0x20B210DE/0x147F10DE | Intel(R) Xeon(R) Platinum 8462Y+ |
| h100-5 | NVIDIA H100 PCIe | 9.0 | 114 | 2048 | 580.126.09 | 0x233110DE/0x162610DE | Intel(R) Xeon(R) Platinum 8352Y CP |
| l40-3 | NVIDIA L40 | 8.9 | 142 | 1536 | 580.126.09 | 0x26B510DE/0x169D10DE | Intel(R) Xeon(R) Platinum 8352Y CP |
| l40-4 | NVIDIA L40 | 8.9 | 142 | 1536 | 580.126.09 | 0x26B510DE/0x169D10DE | Intel(R) Xeon(R) Platinum 8352Y CP |
| l40s-3 | NVIDIA L40S | 8.9 | 142 | 1536 | 580.126.09 | 0x26B910DE/0x185110DE | Intel(R) Xeon(R) Platinum 8462Y+ |
4. Results
4.1 The map: two kernel families, two equivalence relations
Eleven instances, seven classes, one PyTorch build (2.13.0+cu129). Every card reproduces itself run to run under G1 in every cell. Across cards:
| pair | fp32 pinned G1 | fp32 stock G1 | fp32 pinned G0 | bf16 + flash pinned G1 |
|---|---|---|---|---|
| H100 PCIe (hyperstack, AMD host, 570.195) vs H100 PCIe (Lambda, Intel host, 570.148) | identical through 300 | identical through 300 | identical through 300 | identical through 300 |
| L40 vs L40S (both 142 SMs, AD102) | identical through 300 | identical through 300 | identical through 300 | differ (split by 100) |
| A10 (72 SMs) vs A10G (80 SMs), GA102 | differ (split by 100) | differ (split by 100) | differ (split by 100) | identical through 300 |
| L4 (58) vs L40 (142), Ada | differ (split by 100) | differ (split by 100) | differ (split by 100) | differ (split by 100) |
| A100 SXM4 (108) vs H100 PCIe (114) | differ (split by 100) | differ (split by 100) | differ (split by 100) | differ (split by 100) |
The fp32 rows follow cuBLAS’s clause exactly: identity iff same architecture and same SM count. The bf16 row does not: A10 = A10G across a different SM count, L40 ≠ L40S on the same die and SM count. Since bf16 + math (no flash kernel) splits L40/L40S too (§4.3), the bf16 family here is cuBLASLt’s bf16 GEMM path. The two families induce different partitions of the same hardware. The split is at the first parameter update in every differing pair:
| pair | configuration | last identical step | first differing step |
|---|---|---|---|
| h100-4 vs a10g-2 | core-fine:pinned/G1 | 0 | 1 |
| h100-4 vs a10g-2 | core-fine:stock/G1 | 0 | 1 |
| h100-4 vs a10g-2 | bf16-flash-fine:pinned/G1 | 0 | 1 |
| h100-4 vs a10g-2 | bf16-math-fine:pinned/G1 | 0 | 1 |
| h100-4 vs a10g-2 | ozaki-pinned:pinned/G1 | 100 | never (identical through 100) |
| l40-2 vs a10g-2 | core-fine:pinned/G1 | 0 | 1 |
| l40-2 vs a10g-2 | core-fine:stock/G1 | 0 | 1 |
| l40-2 vs a10g-2 | bf16-flash-fine:pinned/G1 | 0 | 1 |
| l40-2 vs a10g-2 | bf16-math-fine:pinned/G1 | 0 | 1 |
| l40-2 vs a10g-2 | ozaki-pinned:pinned/G1 | 100 | never (identical through 100) |
| l40-2 vs h100-4 | core-fine:pinned/G1 | 0 | 1 |
| l40-2 vs h100-4 | core-fine:stock/G1 | 0 | 1 |
| l40-2 vs h100-4 | bf16-flash-fine:pinned/G1 | 0 | 1 |
| l40-2 vs h100-4 | bf16-math-fine:pinned/G1 | 0 | 1 |
| l40-2 vs h100-4 | ozaki-pinned:pinned/G1 | 100 | never (identical through 100) |
| l40-2 vs l40s-2 | core-fine:pinned/G1 | 300 | never (identical through 300) |
| l40-2 vs l40s-2 | core-fine:stock/G1 | 300 | never (identical through 300) |
| l40-2 vs l40s-2 | bf16-flash-fine:pinned/G1 | 0 | 1 |
| l40-2 vs l40s-2 | bf16-math-fine:pinned/G1 | 0 | 1 |
| l40-2 vs l40s-2 | ozaki-pinned:pinned/G1 | 100 | never (identical through 100) |
| l40s-2 vs a10g-2 | core-fine:pinned/G1 | 0 | 1 |
| l40s-2 vs a10g-2 | core-fine:stock/G1 | 0 | 1 |
| l40s-2 vs a10g-2 | bf16-flash-fine:pinned/G1 | 0 | 1 |
| l40s-2 vs a10g-2 | bf16-math-fine:pinned/G1 | 0 | 1 |
| l40s-2 vs a10g-2 | ozaki-pinned:pinned/G1 | 100 | never (identical through 100) |
| l40s-2 vs h100-4 | core-fine:pinned/G1 | 0 | 1 |
| l40s-2 vs h100-4 | core-fine:stock/G1 | 0 | 1 |
| l40s-2 vs h100-4 | bf16-flash-fine:pinned/G1 | 0 | 1 |
| l40s-2 vs h100-4 | bf16-math-fine:pinned/G1 | 0 | 1 |
| l40s-2 vs h100-4 | ozaki-pinned:pinned/G1 | 100 | never (identical through 100) |
4.2 The fp32 key: smCountTarget
cuBLASLt’s trace prints every heuristic query with smCountTarget=<SMs>; the SM count is an
explicit input. Across the two Ada cards the fp32 algorithms agree for every shape; across
architectures they differ for every shape; and — the single-variable statement — on one H100
the same fifteen shapes at 46 and 14 SMs (MIG slices) get a different algorithm for ten and
eleven of them:
| traces compared | precision | shapes in common | shapes with a different algorithm |
|---|---|---|---|
| L40 vs L40S | fp32 | 15 | 0 |
| L40 vs L40S | bf16 (flash run) | 15 | 12 |
| L40 vs L40S | bf16 (math run) | 21 | 16 |
| L40 vs H100 | fp32 | 15 | 15 |
| L40 vs A10G | fp32 | 15 | 13 |
| H100 vs A10G | fp32 | 15 | 15 |
| L40: seeds-only vs deterministic rung | fp32 | 15 | 0 |
| H100: seeds-only vs deterministic rung | fp32 | 15 | 0 |
| H100 configuration | SMs | fp32 shapes with a different algorithm than the whole card (of 15) |
|---|---|---|
| MIG 3g.40gb | 46 | 10 of 15 |
| MIG 1g.10gb | 14 | 11 of 15 |
Asking the heuristic on one card (an L40) for the same shapes at eleven SM-count targets
gives the whole map without buying the cards (§4.6). CUBLAS_WORKSPACE_CONFIG changed no
algorithm on any card: the deterministic rung’s effect on cuBLAS is nil for this program;
what it changes is the index-put/scatter kernels and the memory-efficient attention’s
split-key count.
4.3 The bf16 key: a per-product table
The L40 and L40S share die, SM count, L2, memory clock and threads per SM; CUDA reports peak clocks of 2,490 and 2,520 MHz and different PCI device ids. Three experiments exclude the clock. Locking both cards to the same running clock (1500 MHz, then 900) leaves every digest and kernel record unchanged; changing the clock the runtime reports by 52 % with the applications clock leaves all thirty heuristic choices unchanged on two cards; and NVIDIA’s standalone heuristics library, whose predefined descriptors for the two products differ in exactly one field (the peak SM clock), returns identical configurations for both:
| card | locked SM clock | fp32 G1 digests equal unlocked | bf16 + flash equal unlocked | bf16 + math equal unlocked |
|---|---|---|---|---|
| a10g-2 (NVIDIA A10G) | 1500 MHz | True | True | True |
| a10g-2 (NVIDIA A10G) | 900 MHz | True | True | True |
| h100-4 (NVIDIA H100 PCIe) | 1500 MHz | True | True | True |
| h100-4 (NVIDIA H100 PCIe) | 900 MHz | True | True | True |
| l40-2 (NVIDIA L40) | 1500 MHz | True | True | True |
| l40-2 (NVIDIA L40) | 900 MHz | True | True | True |
| l40s-2 (NVIDIA L40S) | 1500 MHz | True | True | True |
| l40s-2 (NVIDIA L40S) | 900 MHz | True | True | True |
Reported peak clock changed by nvidia-smi -ac on a100-3: 1410000 → 1200000 kHz; heuristic choices that changed: 0 of 30.
Reported peak clock changed by nvidia-smi -ac on l40-3: 2490000 → 1200000 kHz; heuristic choices that changed: 0 of 30.
Descriptor bytes differing between the predefined L40 and L40S (of 136): 8, in words at offsets [8, 56, 128]; the numeric one is the SM clock, 2490000 vs 2520000 kHz.
| comparison (top configuration per shape) | fp32 | bf16 |
|---|---|---|
| L40 vs L40S | 0 of 15 shapes differ | 0 of 15 shapes differ |
| L40 vs L4 | 10 of 15 shapes differ | 7 of 15 shapes differ |
| L40 vs A10 | 12 of 15 shapes differ | 9 of 15 shapes differ |
| L40 vs RTX_6000_ADA | 0 of 15 shapes differ | 0 of 15 shapes differ |
| A100_SXM vs H100_PCIE | 15 of 15 shapes differ | 15 of 15 shapes differ |
What remains is the product identity: cuBLASLt’s bf16 heuristic distinguishes an L40S from an L40 by a per-SKU table. Every public knob is excluded; the table itself is closed.
4.4 A second root, outside cuBLAS
With every GEMM made exact (§5.1), the pinned program on the H100’s 14-SM slice still differed
from the whole card at step 1 — and the step-0 trace showed every forward output equal, the
loss equal, and exactly four of 53 gradients different: the four linear1.bias gradients,
grad.sum(0) over 2,048 rows. ATen’s reduce_kernel (Reduce.cuh, setReduceConfig) splits
each output’s input across thread blocks by multiProcessorCount; 14 SMs give a different
split from 114, 46 and 142. Routing those sums through the exact GEMM closed it (§5.1).
4.5 The device generator
calc_execution_policy (DistributionTemplates.h) caps the launch grid at
multiProcessorCount × (maxThreadsPerMultiProcessor / 256); the grid-stride kernel maps
element li to Philox counter (k, idx) through the grid. Above one wave the map depends on
the SM count. A CPU implementation of Philox4x32-10 with cuRAND’s counter layout and its
integer-to-float conversion reproduces each card’s rand bytes exactly from that card’s two
numbers:
| card | SMs × blocks/SM (cap) | CPU predicts rand at 2^12…2^22 |
rand ≤ 2^16 equals the other cards |
rand ≥ 2^18 equals the L40 |
|---|---|---|---|---|
| a10g-2 (NVIDIA A10G) | 80 × 6 = 480 | 6 of 6 | yes | no |
| h100-4 (NVIDIA H100 PCIe) | 114 × 8 = 912 | 6 of 6 | yes | no |
| l40-2 (NVIDIA L40) | 142 × 6 = 852 | 6 of 6 | (reference) | (reference) |
| l40s-2 (NVIDIA L40S) | 142 × 6 = 852 | 6 of 6 | yes | yes |
Campaign: distinct randn(2^20) digests by (product, SM count): 8 devices, 7 distinct digest sets.
4.6 Seeds alone, at n = 20 and on small slices
| configuration | program | runs | identical |
|---|---|---|---|
| a10g-2 g0x10 | pinned | 10 | yes |
| a10g-2 g0x10 | stock | 10 | yes |
| h100-4 g0x20 | pinned | 20 | yes |
| h100-4 g0x20 | stock | 20 | yes |
| h100-4 mig-1g.10gb-g0x10 | pinned | 10 | yes |
| h100-4 mig-1g.10gb-g0x10 | stock | 10 | NO — 2 distinct trajectories at step 300 |
| h100-4 mig-3g.40gb-g0x10 | pinned | 10 | yes |
| h100-4 mig-3g.40gb-g0x10 | stock | 10 | yes |
| l40-2 g0x10 | pinned | 10 | yes |
| l40-2 g0x10 | stock | 10 | yes |
| l40s-2 g0x10 | pinned | 10 | yes |
| l40s-2 g0x10 | stock | 10 | yes |
The race that the T4 and L4 showed with seeds alone is reproduced on the same silicon by shrinking it to 14 SMs — rarely (one run in ten), for the stock program, and not in cuBLAS (the seeds-only and deterministic-rung traces choose identical algorithms; no in-place reduction appears). The heuristic’s SM-count map from one card:
Heuristic default equals the traced choice for 27 of 30 shapes (the exceptions are the three bias-epilogue addmm calls).
| precision | shapes whose configuration changes with smCountTarget |
distinct configurations per shape (min–max) | vocabulary weight gradient split-K by target |
|---|---|---|---|
| fp32 | 15 of 15 | 2–10 | 14→9, 42→14, 46→15, 58→19, 72→12, 80→13, 98→12, 108→13, 114→14, 142→17 |
| bf16 | 15 of 15 | 2–9 | 14→7, 42→21, 46→23, 58→29, 72→36, 80→32, 98→32, 108→20, 114→20, 142→35 |
4.7 The price of the deterministic rung
| instance | product (SMs) | pinned program G1/G0 | stock program G1/G0 |
|---|---|---|---|
| a10-1 | NVIDIA A10 (72) | 1.09× | 1.07× |
| a100-2 | NVIDIA A100-SXM4-80GB (108) | 1.07× | 1.12× |
| a10g-1 | NVIDIA A10G (80) | 1.07× | 1.13× |
| h100-1 | NVIDIA H100 PCIe (114) | 1.09× | 1.12× |
| h100-2 | NVIDIA H100 PCIe (114) | 1.19× | 1.17× |
| h100-3 | NVIDIA H100 PCIe (114) | 1.14× | 1.14× |
| l4-1 | NVIDIA L4 (58) | 1.14× | 1.10× |
| l40-1 | NVIDIA L40 (142) | 1.08× | 1.19× |
| l40s-1 | NVIDIA L40S (142) | 1.14× | 1.29× |
| t4-3 | Tesla T4 (40) | 1.07× | 1.05× |
| t4g-1 | NVIDIA T4G (40) | — | — |
5. Removing the keys
5.1 An exact GEMM: device-independent bits at 12–16×
| card | Ozaki identical run-to-run | step ms (Ozaki / cuBLAS) | price | max abs loss delta vs cuBLAS |
|---|---|---|---|---|
| a10g-2 | True | 248.9 / 20.2 | 12.29× | 2.3e-05 |
| h100-4 | True | 223.7 / 16.3 | 13.72× | 1.3e-05 |
| l40-2 | True | 229.7 / 18.0 | 12.78× | 4.3e-05 |
| l40s-2 | True | 168.6 / 11.3 | 14.9× | 4.3e-05 |
| pair | Ozaki (pinned program) | Ozaki (stock program) | ozaki2 (bias sums exact too) |
|---|---|---|---|
| h100-4 vs a10g-2 | identical through 100 | identical through 100 | identical through 100 |
| l40-2 vs a10g-2 | identical through 100 | identical through 100 | — |
| l40-2 vs h100-4 | identical through 100 | identical through 100 | — |
| l40-2 vs l40s-2 | identical through 100 | identical through 100 | — |
| l40s-2 vs a10g-2 | identical through 100 | identical through 100 | — |
| l40s-2 vs h100-4 | identical through 100 | identical through 100 | — |
| H100 MIG configuration | cuBLAS fp32 vs whole card | Ozaki vs whole card | Ozaki trace: gradients differing at step 0 | ozaki2 vs whole card |
|---|---|---|---|---|
| MIG 3g.40gb | differ at 1 | identical through 100 | none | identical through 100 (see §8 settlement) |
| MIG 1g.10gb | differ at 1 | differ at 1 | blocks.layers.0.linear1.bias, blocks.layers.1.linear1.bias, blocks.layers.2.linear1.bias, blocks.layers.3.linear1.bias | identical through 100 (see §8 settlement) |
5.2 Forced algorithms: the reference card’s choices, everywhere
- a100-3 (NVIDIA A100-SXM4-80GB, cuBLASLt 120901): self-test 14 of 14 cases, max relative error vs torch 1.756387877326564e-06, identical twice True.
- reference 12.59 ms/step; passthrough dispatch mode 20.83 ms (overhead 8.24 ms); pinned-pinned 28.7 ms = 2.28× (1.63× net of the dispatch mode); pinned-table 28.02 ms = 2.23× (1.57× net of the dispatch mode)
- mig-1g.10gb-pinned-table_vs_whole: identical through 300; gradients differing at step 0: []
- mig-1g.10gb-pinned_vs_whole: identical through 300; gradients differing at step 0: []
- mig-3g.40gb-pinned-table_vs_whole: identical through 300; gradients differing at step 0: []
- mig-3g.40gb-pinned_vs_whole: identical through 300; gradients differing at step 0: []
- the heuristic on the 1g.10gb slice (14 SMs) makes the whole card’s choice for
smCountTarget=14on 30 of 30 shapes - the heuristic on the 3g.40gb slice (42 SMs) makes the whole card’s choice for
smCountTarget=42on 30 of 30 shapes
- l40-3 (NVIDIA L40, cuBLASLt 120901): self-test 7 of 7 cases, max relative error vs torch 2.0512320184339374e-06, identical twice True.
- reference 14.34 ms/step; passthrough dispatch mode 27.2 ms (overhead 12.86 ms); pinned-pinned 43.16 ms = 3.01× (2.11× net of the dispatch mode); pinned-table 41.69 ms = 2.91× (2.01× net of the dispatch mode)
- l40-3 vs a100-3, pinned-pinned: identical through 300; same configurations forced: False; first divergent forward op None, gradients differing 0
- l40-3 vs a100-3, pinned-stock: identical through 300; same configurations forced: False; first divergent forward op None, gradients differing 0
- l40-3 vs a100-3, pinned-table: identical through 300; same configurations forced: True; first divergent forward op None, gradients differing 0
- l40-3 vs a100-3, pinned-table-stock: identical through 300; same configurations forced: True; first divergent forward op None, gradients differing 0
The forced program’s GPU time is the heuristic’s by construction (the same kernels run); the remaining cost above is the Python dispatch mode (measured with a mode that intercepts every op and changes nothing) and the Python/ctypes call path. §5.3 removes both.
5.3 The same forcing in C++: the price is the reference’s
recipe_checkpoints/pinned_gemm.cpp overrides aten::mm, addmm, bmm, baddbmm and
sum.dim_IntList at the CUDA dispatch key (TORCH_LIBRARY_IMPL) with kernels that run
cuBLASLt under the forced configuration — the per-shape table first, the device’s own
heuristic when learning, an ordered pin list last — for fp32 and bf16; every call the kernel
does not take falls through to PyTorch’s own CUDA entry point, so nothing else in the process
changes. Forward and autograd backward reach these kernels directly; there is no Python on
the hot path. The arithmetic is the Python wrapper’s, deliberately, so the two could be
compared bit for bit. The extension is built with g++ alone from the CUDA headers in the
nvidia wheels and compiled on the instance in 18–24 s.
- h100-5: extension built on the instance in 20.3 s (torch 2.13.0+cu129, cuBLASLt 120901); self-test 7 of 7 fp32 cases within 2.0987185080050912e-06 of
torch.matmul, 7 of 7 bit-identical to the Python wrapper’s output for the same pin; 7 of 7 bf16 cases, 7 equal to cuBLAS’s own bf16 bits.- price without instrumentation: reference 15.71 ms/step, C++ forced L40 table 18.39 ms = 1.171×; stock variant 9.77 → 10.98 ms (1.124×); this card’s own learned table 14.74 ms (0.938×); the Python dispatch mode in the same session 49.06 ms (3.123×)
- the L40 table on this card: configuration sources per shape {‘learned (table entry refused here)’: 16, ‘table’: 6} of 22; refused entries 16
- algorithm ids this card offers: fp32 [0, 1, 10, 11, 12, 13, 14, 15, 16, 20, 55, 57, 76]; bf16 [5, 6, 11, 12, 13, 14, 16, 21, 23, 24, 30, 31, 39, 41, 42, 58, 66, 70, 71]
- mig-1g.10gb-cpp-native_vs_whole: identical through 300; gradients differing at step 0: 0
- mig-1g.10gb-cpp-pin_vs_whole: identical through 300; gradients differing at step 0: 0
- mig-1g.10gb-cpp-ref_vs_whole: differ at step 1; gradients differing at step 0: 50
- mig-1g.10gb-cpp-table_vs_whole: differ at step 1; gradients differing at step 0: 53
- mig-3g.40gb-cpp-native_vs_whole: identical through 300; gradients differing at step 0: 0
- mig-3g.40gb-cpp-pin_vs_whole: identical through 300; gradients differing at step 0: 0
- mig-3g.40gb-cpp-ref_vs_whole: differ at step 1; gradients differing at step 0: 49
- mig-3g.40gb-cpp-table_vs_whole: differ at step 1; gradients differing at step 0: 53
- profiled step: GEMM-like launches 83 (C++) vs 79 (reference); total launches 1414 vs 1286 (the deterministic-mode fills for the wrapper’s outputs, the un-fused bias adds, and
gemvlaunches replacing ATen’s reduction for the bias gradients) - cpp-table: 2 run(s), run-to-run identical True, 21.43 ms/step, intercepted {‘addmm’: 3600, ‘bmm’: 7200, ‘mm’: 11700, ‘sum->mm’: 4800}, let through {}, refused {‘128x128@128x64|NN|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘128x128@128x64|NT|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘128x64@64x128|TN|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x1024@1024x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x1024@1024x256|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x1024|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x1024|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x256|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x65|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x768|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x768@768x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘256x2048@2048x1024|NT|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘256x2048@2048x256|NT|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘64x128@128x128|NT|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘65x2048@2048x256|NT|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’}
- cpp-table-stock: 1 run(s), run-to-run identical True, 14.88 ms/step, intercepted {‘addmm’: 3600, ‘bmm’: 7200, ‘mm’: 11700, ‘sum->mm’: 4800}, let through {}, refused {‘128x128@128x64|NN|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘128x128@128x64|NT|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘128x64@64x128|TN|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x1024@1024x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x1024@1024x256|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x1024|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x1024|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x256|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x65|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x256@256x768|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘2048x768@768x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘256x2048@2048x1024|NT|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘256x2048@2048x256|NT|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘64x128@128x128|NT|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘65x2048@2048x256|NT|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’}
- bf16-cpp-math: 2 run(s), run-to-run identical True, 22.36 ms/step, intercepted {‘addmm’: 3600, ‘bmm’: 7200, ‘mm’: 11700}, let through {}, refused {‘128x128@128x64|NN|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘128x64@64x128|TN|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x1024@1024x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x1024@1024x256|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x256@256x1024|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x256@256x1024|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x256@256x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x256@256x256|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x768@768x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:256x2048@2048x256|NT|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:64x128@128x128|NT|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’}
- bf16-cpp-flash: 2 run(s), run-to-run identical True, 18.41 ms/step, intercepted {‘addmm’: 3600, ‘mm’: 11700}, let through {}, refused {‘bf16:2048x1024@1024x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x1024@1024x256|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x256@256x1024|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x256@256x1024|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x256@256x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x256@256x256|TN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:2048x768@768x256|NN|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘bf16:256x2048@2048x256|NT|b1 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’}
- l40-4: extension built on the instance in 24.2 s (torch 2.13.0+cu129, cuBLASLt 120901); self-test 7 of 7 fp32 cases within 2.0512320184339374e-06 of
torch.matmul, 7 of 7 bit-identical to the Python wrapper’s output for the same pin; 7 of 7 bf16 cases, 7 equal to cuBLAS’s own bf16 bits.- price without instrumentation: reference 14.38 ms/step, C++ forced L40 table 13.73 ms = 0.955×; stock variant 9.49 → 9.29 ms (0.979×); the Python dispatch mode in the same session 39.63 ms (2.756×)
- the L40 table on this card: configuration sources per shape {‘table’: 22} of 22
- profiled step: GEMM-like launches 75 (C++) vs 75 (reference); total launches 1414 vs 1290 (the deterministic-mode fills for the wrapper’s outputs, the un-fused bias adds, and
gemvlaunches replacing ATen’s reduction for the bias gradients) - cpp-table: 2 run(s), run-to-run identical True, 18.38 ms/step, intercepted {‘addmm’: 3600, ‘bmm’: 7200, ‘mm’: 11700, ‘sum->mm’: 4800}, let through {}, refused {}
- cpp-table-stock: 1 run(s), run-to-run identical True, 12.54 ms/step, intercepted {‘addmm’: 3600, ‘bmm’: 7200, ‘mm’: 11700, ‘sum->mm’: 4800}, let through {}, refused {}
- bf16-cpp-math: 2 run(s), run-to-run identical True, 21.86 ms/step, intercepted {‘addmm’: 3600, ‘bmm’: 7200, ‘mm’: 11700}, let through {}, refused {}
- bf16-cpp-flash: 2 run(s), run-to-run identical True, 16.64 ms/step, intercepted {‘addmm’: 3600, ‘mm’: 11700}, let through {}, refused {}
- l40s-3: extension built on the instance in 18.1 s (torch 2.13.0+cu129, cuBLASLt 120901); self-test 7 of 7 fp32 cases within 2.0512320184339374e-06 of
torch.matmul, 7 of 7 bit-identical to the Python wrapper’s output for the same pin; 7 of 7 bf16 cases, 7 equal to cuBLAS’s own bf16 bits.- price without instrumentation: reference 10.29 ms/step, C++ forced L40 table 10.99 ms = 1.068×; stock variant 6.33 → 7.78 ms (1.229×); the Python dispatch mode in the same session 25.72 ms (2.5×)
- the L40 table on this card: configuration sources per shape {‘table’: 22} of 22
- profiled step: GEMM-like launches 75 (C++) vs 75 (reference); total launches 1414 vs 1290 (the deterministic-mode fills for the wrapper’s outputs, the un-fused bias adds, and
gemvlaunches replacing ATen’s reduction for the bias gradients) - cpp-table: 2 run(s), run-to-run identical True, 12.96 ms/step, intercepted {‘addmm’: 3600, ‘bmm’: 7200, ‘mm’: 11700, ‘sum->mm’: 4800}, let through {}, refused {}
- cpp-table-stock: 1 run(s), run-to-run identical True, 9.1 ms/step, intercepted {‘addmm’: 3600, ‘bmm’: 7200, ‘mm’: 11700, ‘sum->mm’: 4800}, let through {}, refused {}
- bf16-cpp-math: 2 run(s), run-to-run identical True, 16.36 ms/step, intercepted {‘addmm’: 3600, ‘bmm’: 7200, ‘mm’: 11700}, let through {}, refused {}
- bf16-cpp-flash: 2 run(s), run-to-run identical True, 13.3 ms/step, intercepted {‘addmm’: 3600, ‘mm’: 11700}, let through {}, refused {}
- h100-5 bf16-cpp-math vs l40-4 pybf16-cpp-math (C++ vs Python, bf16): differ at step 1 (11 compared steps); same configurations forced: False; gradients differing at step 0: 53
- h100-5 cpp-pin vs a100-3 pinned-pinned (C++ vs Python, single pin): identical through step 300 (11 compared steps); same configurations forced: False; gradients differing at step 0: 0
- h100-5 cpp-pin vs l40-3 pinned-pinned (C++ vs Python, single pin): identical through step 300 (11 compared steps); same configurations forced: True; gradients differing at step 0: 0
- h100-5 cpp-table vs a100-3 pinned-table (C++ vs Python, same forced table): differ at step 1 (11 compared steps); same configurations forced: False; gradients differing at step 0: 53
- h100-5 cpp-table vs l40-3 pinned-table (C++ vs Python, same forced table): differ at step 1 (11 compared steps); same configurations forced: False; gradients differing at step 0: 53
- h100-5 vs l40-4, cpp-pin: identical through step 300; same configurations forced: True; gradients differing at step 0: 0
- h100-5 vs l40-4, cpp-table: differ at step 1; same configurations forced: False; gradients differing at step 0: 53
- h100-5 vs l40s-3, cpp-pin: identical through step 300; same configurations forced: True; gradients differing at step 0: 0
- h100-5 vs l40s-3, cpp-table: differ at step 1; same configurations forced: False; gradients differing at step 0: 53
- l40-4 bf16-cpp-math vs l40-4 pybf16-cpp-math (C++ vs Python, bf16): identical through step 300 (11 compared steps); same configurations forced: True; gradients differing at step 0: 0
- l40-4 cpp-pin vs a100-3 pinned-pinned (C++ vs Python, single pin): identical through step 300 (11 compared steps); same configurations forced: False; gradients differing at step 0: 0
- l40-4 cpp-pin vs l40-3 pinned-pinned (C++ vs Python, single pin): identical through step 300 (11 compared steps); same configurations forced: True; gradients differing at step 0: 0
- l40-4 cpp-table vs a100-3 pinned-table (C++ vs Python, same forced table): identical through step 300 (11 compared steps); same configurations forced: True; gradients differing at step 0: 0
- l40-4 cpp-table vs l40-3 pinned-table (C++ vs Python, same forced table): identical through step 300 (11 compared steps); same configurations forced: True; gradients differing at step 0: 0
- l40-4 vs l40s-3, cpp-pin: identical through step 300; same configurations forced: True; gradients differing at step 0: 0
- l40-4 vs l40s-3, cpp-table: identical through step 300; same configurations forced: True; gradients differing at step 0: 0
- l40s-3 bf16-cpp-math vs l40-4 pybf16-cpp-math (C++ vs Python, bf16): identical through step 300 (11 compared steps); same configurations forced: True; gradients differing at step 0: 0
- l40s-3 cpp-pin vs a100-3 pinned-pinned (C++ vs Python, single pin): identical through step 300 (11 compared steps); same configurations forced: False; gradients differing at step 0: 0
- l40s-3 cpp-pin vs l40-3 pinned-pinned (C++ vs Python, single pin): identical through step 300 (11 compared steps); same configurations forced: True; gradients differing at step 0: 0
- l40s-3 cpp-table vs a100-3 pinned-table (C++ vs Python, same forced table): identical through step 300 (11 compared steps); same configurations forced: True; gradients differing at step 0: 0
- l40s-3 cpp-table vs l40-3 pinned-table (C++ vs Python, same forced table): identical through step 300 (11 compared steps); same configurations forced: True; gradients differing at step 0: 0
Three things settle. The C++ program is the Python program: bit-identical self-test against
the Python pin on both cards, and the fp32 trajectory under the L40 table equal to the Python
mode’s run of the previous session at every compared step (the same physical card, so this
is a comparison of implementations; the cross-card comparisons are L40 → L40S and L40 →
A100). The price is the reference’s: with two clean runs each, paired by run index, the
pinned program costs 0.955 / 0.972× the reference on the L40 and 1.068 / 0.999× on the L40S,
the stock program 0.979 / 0.978× and 1.229 / 0.972× — the reference itself moves 7–15 %
between its own two runs on the L40S, and on the L40 the C++ program is below both reference
runs. The residual is a fixed ~120 small launches per step: deterministic-mode fills for the
outputs the wrapper allocates (+123), un-fused bias adds (+12), sixteen gemv launches
replacing sixteen of ATen’s seventeen reduce_kernel launches — the sixteen linear-bias
gradients; the seventeenth is the positional embedding’s gradient, a sum over the batch
dimension alone, which stays on ATen’s reduction kernel (16 inputs per output, so no
cross-block split; see §7) — and four fewer memsets. The kernel record has the same number of
GEMM launches in both programs, 75, of which 71 are the same kernel: the four bias-epilogue
addmm shapes run ampere_sgemm_64x64_tn in place of the fused-epilogue
ampere_sgemm_128x32_tn, and the forced table adds four split-K reduction launches (21
against 17). The 2.2–2.9× of §5.2 was the instrument; the fix costs the reference’s time.
Hopper (registered as §10 before the instance existed). Hopper’s fp32 heuristic chooses
algoId 55 where Ada chooses 0, 1 and 20, so the first question was admissibility: the H100
lists Ada’s ids (its fp32 ids are 0, 1, 10–16, 20, 55, 57, 76) but cublasLtMatmulAlgoCheck
accepts the L40’s table for only 6 of 22 shapes — every algoId 0 and 1 configuration is
refused, the 13 and 20 entries pass — and the extension learns Hopper’s own choice for the
rest (recorded per shape), so that program’s bits differ from the L40’s from step 1, as the
registration’s second branch said. One configuration is accepted everywhere: algoId 20, stages 8×5, splitK 1, and the C++ program under it is byte-identical across the H100
(114 SMs), the L40 (142), the L40S (142), the A100 (108) and the H100’s 46- and 14-SM MIG
slices through step 300 — three architectures, five SM counts, one trajectory, on a stock
PyTorch build through a public API. The A100 point had been pinned to the same algorithm
with tile 128×64 (the other sessions used 64×64; the pair’s “same configurations: False” in
the regions records this), so the identity also says that without split-K the output tile
does not alter algoId 20’s accumulation order. The pin’s price on the H100 is 1.05×
(instrumented); a Hopper-native learned table costs 0.94–0.99× the reference (14.74 / 14.04
against 15.71 / 14.16 ms), is run-to-run identical, and is byte-identical across the H100’s
MIG slices while the reference splits at step 1. Tables travel within an architecture (L40 →
L40S; L40 → A100 for these shapes), the single pin travels across, and the two are not in
tension: a deployment pins its reference card’s table and keeps the cross-architecture pin
for the shapes a foreign card refuses. One limit surfaced: on the A100 the slice’s heuristic
equalled the whole card’s asked for the slice’s SM count on 30 of 30 shapes; on the H100 on 26
and 23 of 30 (46 and 14 SMs) — smCountTarget is the dominant key on Hopper, not the only one
(the slice’s L2 share and bandwidth are the untested candidates); the forced program does not
consult the heuristic and is unaffected.
5.4 bf16 under forced algorithms: the table neutralised
§4.3 named the bf16 key by elimination. Forcing lets it be confirmed directly: if the L40 and
the L40S differ in bf16 only because their heuristics consult a per-product table, then
forcing the L40’s choices on the L40S must make the two trajectories one, and forcing nothing
must leave them apart. On the L40 the bf16 program’s GEMM shapes were learned into a table
(19 shapes under math attention, 15 under flash — attention’s bmms are inside the flash
kernel); on the L40S that table was forced.
- h100-5 vs l40-4 (compute capability 9.0 vs 8.9 — different architectures, no bf16 identity claimed): every bf16 cell differs from step 1, as expected with different kernels forced
- h100-5 vs l40s-3 (compute capability 9.0 vs 8.9 — different architectures, no bf16 identity claimed): every bf16 cell differs from step 1, as expected with different kernels forced
- l40-4 vs l40s-3:
- bf16-cpp-flash: identical through step 300; gradients differing at step 0: 0, first divergent forward op None
- bf16-cpp-math: identical through step 300; gradients differing at step 0: 0, first divergent forward op None
- bf16-learn-flash learned tables: the two cards’ own heuristics choose a different configuration for 7 of 15 bf16 shapes in the fields that name the kernel and its summation order (algorithm id, tile, stages, split-K, reduction scheme); 13 of 15 if the CTA-swizzle flag is counted; fields differing: {‘id’: 4, ‘reduction’: 2, ‘splitk’: 3, ‘stages’: 4, ‘swizzle’: 8, ‘tile’: 5}
- bf16-learn-math learned tables: the two cards’ own heuristics choose a different configuration for 11 of 19 bf16 shapes in the fields that name the kernel and its summation order (algorithm id, tile, stages, split-K, reduction scheme); 17 of 19 if the CTA-swizzle flag is counted; fields differing: {‘id’: 5, ‘reduction’: 2, ‘splitk’: 3, ‘stages’: 8, ‘swizzle’: 9, ‘tile’: 8}
- bf16-ref-flash: differ at step 1 (identical at 0); gradients differing at step 0: 53, first divergent forward op {‘a’: ‘int16/16x128x65/d78e4073813b6fe4’, ‘b’: ‘int16/16x128x65/aa3def1cec0c638d’, ‘index’: 40, ‘module’: ‘head’}
- bf16-ref-math: differ at step 1 (identical at 0); gradients differing at step 0: 53, first divergent forward op {‘a’: ‘int16/16x128x65/0bb683f42be310b2’, ‘b’: ‘int16/16x128x65/4e3fd01e81a5f433’, ‘index’: 40, ‘module’: ‘head’}
- h100-5 bf16-learn-math_learned: 19 shapes; refused {‘128x128@128x64|NN|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’, ‘128x64@64x128|TN|b64 (table)’: ‘AlgoCheck refused: NOT_SUPPORTED’}
- h100-5 bf16-learn-flash: 15 bf16 GEMM shapes in a step, all intercepted, none refused
- l40-4 bf16-learn-math: 19 bf16 GEMM shapes in a step, all intercepted, none refused
- l40-4 bf16-learn-flash: 15 bf16 GEMM shapes in a step, all intercepted, none refused
- l40s-3 bf16-learn-math: 19 bf16 GEMM shapes in a step, all intercepted, none refused
- l40s-3 bf16-learn-flash: 15 bf16 GEMM shapes in a step, all intercepted, none refused
Both halves hold. With the table forced, bf16 + math and bf16 + flash are byte-identical
across the two cards through step 300, every gradient equal at step 0. Without it — the same
sessions, the same build — the stock bf16 programs split at step 1 with either attention
backend, the first divergent forward output being the vocabulary projection, with 53
gradients differing at step 0. The table’s content, in the wrapper’s own words: the two
cards’ heuristics choose a different configuration for 11 of 19 bf16 shapes in the fields
that name the kernel and its summation order — algorithm id (5 shapes), tile (8), stages
(8), split-K (3), reduction scheme (2) — and for 17 of 19 if the CTA-swizzle flag, which
changes the order blocks are scheduled in but not the arithmetic, is counted (the trace
comparison of §4.3, 12 of 15, counts every printed field including that flag). For the
128×128 @ 128×64 attention batch in NN layout, the L40 takes algo 21, tile 64×64, stages 32×6 and the L40S algo 21, tile 32×32, stages 64×1. The flash kernel is not a second
per-product key for this pair. The table is still closed source and has not been read; it has been named and then
neutralised, which is the confirmation §4.3 could not give. A stronger check than digest equality: the profiled
step’s kernel record under the forced table is the same sequence of kernels on both cards —
1,755 launches with identical names, counts and order for bf16 + math, 1,495 for bf16 + flash
— where under the stock heuristics the two cards’ kernel-name sets differ (a splitKreduce
kernel runs on the L40S alone). Same selection, same kernels, same bytes: the table is a
selector, and the alternative — a per-product code path inside kernels of the same name — did
not occur here.
5.5 A divergence that was not ours: pytorch/pytorch#105726
The mechanism should read on a report we did not make. pytorch/pytorch#105726 (open since July
2023; module: half) is a batched fp16 torch.matmul — attention scores × values, batch 32,
116×116 @ 116×128 — whose output differs between an RTX 3090 (82 SMs), two A100s (108) and an L4
(58); the two A100s agreed with each other; the maintainers’ replies were “two different GPUs
running two different (likely nondeterministic) matmul kernels” and “a different order of
operations”. The reporter attached the inputs and every card’s output — the L4’s equals the
3090’s. We ran the reporter’s inputs on an A100-SXM4-80GB and an L4 (torch 2.13.0+cu129;
PREREGISTRATION §11, registered before the instances existed):
| card | SMs | stock torch.matmul sha256 |
equals the reporter’s file |
|---|---|---|---|
| NVIDIA A100-SXM4-80GB | 108 | 64ee87e23355fa00 |
42_output.npy |
| NVIDIA L4 | 58 | 9b234fb11ea41f3d |
38_output.npy, L4_output.npy |
smCountTarget |
NVIDIA A100-SXM4-80GB | NVIDIA L4 |
|---|---|---|
| default | algoId 23, tile 32x32, stages 64x1, splitK 1, reduction NONE | algoId 23, tile 32x32, stages 32x1, splitK 1, reduction NONE |
| 58 | algoId 23, tile 256x128, stages 32x3, splitK 1, reduction NONE | algoId 23, tile 32x32, stages 32x1, splitK 1, reduction NONE |
| 82 | algoId 23, tile 256x128, stages 32x3, splitK 1, reduction NONE | algoId 23, tile 32x32, stages 32x1, splitK 1, reduction NONE |
| 108 | algoId 23, tile 32x32, stages 64x1, splitK 1, reduction NONE | algoId 23, tile 32x32, stages 32x1, splitK 1, reduction NONE |
| 114 | algoId 23, tile 32x32, stages 64x1, splitK 1, reduction NONE | algoId 23, tile 32x32, stages 32x1, splitK 1, reduction NONE |
| 142 | algoId 23, tile 32x32, stages 64x1, splitK 1, reduction NONE | algoId 23, tile 32x32, stages 32x1, splitK 1, reduction NONE |
Configurations accepted on every card: 3; computing identical bits on every card: 3.
algoId 23, tile 32x32, stages 64x1, splitK 1: identical across cards,64ee87e23355fa00= the reporter’s 42_output.npy; max abs error vs fp64 9.57e-04, 9.57e-04algoId 23, tile 32x32, stages 128x1, splitK 1: identical across cards,f15c8913bfc20b4f; max abs error vs fp64 9.57e-04, 9.57e-04algoId 16, tile UNDEFINED, stages UNDEFINED, splitK 1: identical across cards,607b1fa095d63951; max abs error vs fp64 9.57e-04, 9.57e-04
Three things. The reporter’s bytes are reproduced exactly on both cards three years, two CUDA
majors and a torch major later: the divergence is a property of the (device, call), not of
their machines. The key is again the heuristic’s configuration and the SM count moves it — but
in this fp16 tensor-core family the operative field is the stage count (the K-slab width, 64
on the A100 against 32 on the L4 and, when the A100’s heuristic is asked to plan for 58 or 82
SMs, 32 again), not the tile or split-K; there is no split-K anywhere, which is why the
reporter’s allow_fp16_reduced_precision_reduction experiment changed nothing. And the pin
works across architectures for configurations both cards accept: four do, each computing the
same bits on Ampere and on Ada — including the reporter’s A100 output on an L4 and their L4/3090
output on an A100 — at exactly the stock kernels’ error against the fp64 product. Either card
produces either of the reporter’s outputs when told which configuration to run. The result was
posted on the thread with a self-contained ctypes script (docs/upstream/pytorch-105726-*);
#137389 (A5000 / A6000 / A10 — one GA102 die at 64 / 84 / 72 SMs, fp32) is the same mechanism
through the SM count.
6. What a transcript must carry
Product id (as nvidia-smi reports it), SM count, torch build, dtype, attention backend, and
the deterministic flag — and, if the algorithms are forced, the pin table. With the product
id and SM count matching, exact replay across machines held on four H100 PCIe cards (three
providers, two host CPU vendors, three driver builds) and two A10G instances; with the
algorithms forced, across an L40 and an A100 and across the A100’s 14-, 42- and 108-SM
configurations.
With the algorithms forced (§5.2–5.4), the transcript names the forced table — its digest and
the cuBLASLt version that accepted it — in place of the product id, and interchangeability
follows from the table rather than from the SKU: any device on which cublasLtMatmulAlgoCheck
accepts every entry is a candidate replayer, and whether it reproduces the reference bits is a
measurement — here made by the 300-step digest comparison for the L40S (C++), the A100 and its
slices (the Python implementation), against the L40.
7. Threats to validity and what is not claimed
One PyTorch build; one small model at one shape — the reduction split and the GEMM heuristics
are shape-dependent, and the 46-SM slice coinciding with the whole card is luck of the shape,
not a law. n = 2 or 3 per cell except where stated; the bf16 controls, the learned-table runs
and the stock-variant C++ runs of §5.3–5.4 are single runs per card (the L40’s control is
duplicated by the first-pass Python-mode run), the clean price runs are pairs. One reduction in
the forced program — the positional embedding’s gradient — still runs on ATen’s reduction
kernel; it needs no cross-block split at this shape, and the program was identical across 14,
42 and 108 SMs with it in place, but a shape that did split there would not be covered by the
forcing. The Ozaki GEMM’s accuracy is fp32-class,
not cuBLAS’s bits. The 14-SM seeds-only race is located to a short list of kernels, not a
line. The bf16 key is named by elimination of every public knob, not by reading the table.
Forced algorithms were shown across three architectures for fp32 (Ampere, Ada, Hopper) on a
single configuration, across Ada ↔ Ampere for a whole learned table, and across two products
of one architecture for bf16 (L40, L40S, both 142 SMs), under one build; a learned table is
only partly admissible on another architecture (6 of 22 entries on Hopper), and no bf16
statement is made across architectures or SM counts. “The SM count enters through
smCountTarget and nothing else” is an Ampere statement (30 of 30); on Hopper the slice’s
heuristic departs from the SM-count query on 4–7 of 30 shapes for a reason not identified.
The per-product table has been neutralised, not read. Prices are for a
launch-bound 0.4 M-parameter model; the C++ implementation’s residual is a fixed number of
small launches per step, whose share on a larger model is to be measured, not assumed.
8. Reproducibility
Harness: atlas/runners/gpu/ (gpu_probe.py, mech_local.sh, pin_local.sh,
mech_compare.py, pin_compare.py, rng_replay.py, nvmmh_probe.py); the interventions:
python/recipe_checkpoints/exact_gemm.py, cublaslt_pin.py (tests in test_exact_gemm.py),
and the C++ extension pinned_gemm.cpp with its loader pinned_gemm.py; the forced tables
atlas/runners/gpu/pin_table_l40.json and pin_table_l40_bf16.json. Result files:
python/recipe_checkpoints/results/gpu/<instance>/ and the reports
2026-09-10-gpu-campaign.*, 2026-09-10-gpu-mechanism.*, 2026-09-11-nvmmh-descriptor.*,
2026-09-11-gpu-pinning.*, 2026-09-11-gpu-cpp-bf16.*, 2026-09-11-issue105726.* (with
atlas/runners/gpu/issue105726.py and the reporter’s files under issue105726-inputs/). Registration and
settlement: python/recipe_checkpoints/PREREGISTRATION.md §6–§11. cuBLASLt traces are kept beside the
instance backups. This document’s regions are regenerated by
atlas/runners/gpu/mech_paper_numbers.py --render and checked by --check.
Bibliography (verified against primary sources, 2026-09-11)
- [cuBLAS-doc] NVIDIA. cuBLAS Library documentation, §2.1.4 “Results Reproducibility”, §3.1.3 “cuBLASLt Logging”,
cublasLtMatmulDescAttributes_t(CUBLASLT_MATMUL_DESC_SM_COUNT_TARGET),cublasLtReductionScheme_t. https://docs.nvidia.com/cuda/cublas/ (text fetched 2026-09-10). - [PT-repro] PyTorch. Reproducibility (developer note). https://docs.pytorch.org/docs/stable/notes/randomness.html.
- [PT-src] PyTorch v2.13.0 sources:
aten/src/ATen/native/cuda/DistributionTemplates.h(lines 49–61, 65–88),aten/src/ATen/native/cuda/Reduce.cuh(setReduceConfig),aten/src/ATen/native/transformers/cuda/attention_backward.cu(num_splits_key). https://github.com/pytorch/pytorch/tree/v2.13.0. - [PT-issue] pytorch/pytorch#196613, CUDA generator output depends on the device’s SM count, filed by this project 2026-09-10. https://github.com/pytorch/pytorch/issues/196613.
- [PoL] H. Jia, M. Yaghini, C. A. Choquette-Choo, N. Dullerud, A. Thudi, V. Chandrasekaran, N. Papernot. Proof-of-Learning: Definitions and Practice. IEEE S&P 2021; arXiv:2103.05633.
- [Zhang21] R. Zhang, J. Liu, Y. Ding, Z. Wu, Q. Wang, K. Ren. “Adversarial Examples” for Proof-of-Learning. IEEE S&P 2022; arXiv:2108.09454.
- [Fang22] C. Fang, H. Jia, A. Thudi, M. Yaghini, C. A. Choquette-Choo, N. Dullerud, V. Chandrasekaran, N. Papernot. Proof-of-Learning is Currently More Broken Than You Think. IEEE EuroS&P 2023; arXiv:2208.03567.
- [SAB24] M. Srivastava, S. Arora, D. Boneh. Optimistic Verifiable Training by Controlling Hardware Nondeterminism. NeurIPS 2024; arXiv:2403.09603.
- [RepOps] Gensyn. RepOps: reproducible operators (2025) — as cited and verified in
docs/04-VERIFIABLE-TRAINING.md§7. - [He25] H. He et al. (Thinking Machines Lab). Defeating Nondeterminism in LLM Inference. Blog, 10 September 2025. https://thinkingmachines.ai/blog/defeating-nondeterminism-in-llm-inference/.
- [Ozaki12] K. Ozaki, T. Ogita, S. Oishi, S. M. Rump. Error-free transformations of matrix multiplication by using fast routines of matrix multiplication and its applications. Numerical Algorithms 59 (2012); doi:10.1007/s11075-011-9478-1.
- [Uchino24] Y. Uchino, K. Ozaki, T. Imamura. Performance Enhancement of the Ozaki Scheme on Integer Matrix Multiplication Unit. arXiv:2409.13313 (2024).
- [Salmon11] J. K. Salmon, M. A. Moraes, R. O. Dror, D. E. Shaw. Parallel random numbers: as easy as 1, 2, 3. SC ’11; doi:10.1145/2063384.2063405.
- [nvMMH] NVIDIA. nvMatmulHeuristics 0.1.0.27, pip
nvidia-matmul-heuristics(headernvMatmulHeuristics.h, predefined GPU descriptors). - [Veit26] C. Veit. Kernel Contracts: A Specification Language for ML Kernel Correctness Across Heterogeneous Silicon. arXiv:2604.22032 (cs.LG, cs.PL), 23 April 2026. Verified at the arXiv record 2026-09-11.
- [Zhang25] Z. Zhang, X. Ding, J. Yuan, R. Liu, H. Mao, J. Xing, Z. Liu. Deterministic Inference across Tensor Parallel Sizes That Eliminates Training–Inference Mismatch (Tree-Based Invariant Kernels, TBIK). arXiv:2511.17826, v2 29 May 2026; ICML 2026 per the authors’ repository, https://github.com/nanomaoli/llm_reproducibility. Verified 2026-09-11.
- [Yuan25] J. Yuan, H. Li, X. Ding, W. Xie, Y.-J. Li, W. Zhao, K. Wan, J. Shi, X. Hu, Z. Liu. Understanding and Mitigating Numerical Sources of Nondeterminism in LLM Inference. arXiv:2506.09501, v2 24 October 2025; NeurIPS 2025 per the authors’ repository. Verified 2026-09-11.
Pending verification — not citable yet
- TBIK and Kernel Contracts were quarantined here until 2026-09-11 (the nicknames had been searched, not the ids); both are now verified at their arXiv records and cited above.
- Gensyn RepOps’ “30–60 % overhead” is carried from the companion document’s verification, not re-verified here.
Manuscript note
Status: draft, 2026-09-11. Markdown master. Companion to
docs/paper/ship1.md(CPU replay) anddocs/paper/atlas.md(tools). Every load-bearing number lives in a generated region delimited by<!-- MECH:BEGIN key -->/<!-- MECH:END key -->, filled bypython3 atlas/runners/gpu/mech_paper_numbers.py --renderfrom the result files underpython/recipe_checkpoints/results/and gated by the same command with--check. Prose argues; regions count. Citations appear in the bibliography only when verified against a primary source; the rest is quarantined and not cited as established.