where the gates land · rows marked predicted or computed are roofline-derived · cpu-interpret and simulated rows are verified off-chip · hardware rows replace both · L-rows come from the path labs, and one of them is retracted in public
Every number, with provenance.
Nothing quantitative appears on this site outside an instrument, and every value states what it is. Each record carries chip, dtype, shapes, date, and the lab that produced it; gate verdicts on the chapter pages resolve to rows here. The L-rows are what the JAX, XLA, and PyTorch labs measured, including one row that says RETRACTED, because a number this site published and later disproved belongs in the same table as the ones that held.
39
The record
| id | kernel | metric | value | baseline | chip | dtype | shapes | date |
|---|---|---|---|---|---|---|---|---|
| P-001 | matmul NxNxN | predicted floor (roofline) | 697.7 µs | chip constants | v5e (predicted) | bf16 | 4096³ | 2026-07-26 |
| P-002 | matmul skinny 8xNxN | predicted floor (roofline) | 41.1 µs | chip constants | v5e (predicted) | bf16 | 8x4096x4096 | 2026-07-26 |
| P-003 | elementwise add | predicted floor (roofline) | 122.8 µs | chip constants | v5e (predicted) | bf16 | 4096² | 2026-07-26 |
| P-004 | softmax rows | predicted floor (roofline) | 81.8 µs | chip constants | v5e (predicted) | bf16 | 4096² | 2026-07-26 |
| P-005 | naive attention spill | computed HBM round-trip | 268 MB | score matrix size | any (computed) | bf16 | seq 8192 | 2026-07-26 |
| C-001 | flash forward (lab 3.2) | max err vs reference | ≤ 1e-4 | jax.nn.softmax @ v | cpu interpret | f32 | 512x1024x64 | 2026-07-26 |
| C-002 | attention custom_vjp grads (lab 3.3) | max err vs jax.grad | ≤ 1e-5 | autodiff of reference | cpu interpret | f32 | 128x256x64 | 2026-07-26 |
| C-003 | causal flash, blocks skipped (lab 3.4) | max err vs masked reference | ≤ 1e-4 | dense masked attention | cpu interpret | f32 | 512x512x64 | 2026-07-26 |
| C-004 | ring all-gather (lab 4.1) | match vs lax.all_gather | exact | collective | 8 simulated devices | f32 | 8 shards x 4x128 | 2026-07-26 |
| C-005 | ring attention (lab 4.2) | max err vs full attention | ≤ 1e-5 | single-device reference | 8 simulated devices | f32 | 128x256x64 sharded | 2026-07-26 |
| M-001 | matmul NxNxN | measured vs roofline floor | 362.5 µs vs 149.4 µs floor (2.4x) | v6e constants | v6e-1 (TPU v6 lite) | bf16 | 4096³ | 2026-07-26 |
| M-002 | matmul skinny 8xNxN | measured vs roofline floor | 157.1 µs vs 21.1 µs floor (7.5x) | v6e constants | v6e-1 (TPU v6 lite) | bf16 | 8x4096x4096 | 2026-07-26 |
| M-003 | elementwise add | measured vs roofline floor | 225.3 µs vs 62.9 µs floor (3.6x) | v6e constants | v6e-1 (TPU v6 lite) | bf16 | 4096² | 2026-07-26 |
| M-004 | softmax rows | measured vs roofline floor | 248.4 µs vs 41.9 µs floor (5.9x) | v6e constants | v6e-1 (TPU v6 lite) | bf16 | 4096² | 2026-07-26 |
| M-005 | reduce_sum | measured vs roofline floor | 162.4 µs vs 21.0 µs floor (7.7x) | v6e constants | v6e-1 (TPU v6 lite) | bf16 | 4096² | 2026-07-26 |
| M-006 | pallas tiled matmul (lab 1.2) | best of 5 block shapes vs XLA | 597.5 µs vs 361.3 µs (1.65x), best block 512x1024x512 | jnp matmul | v6e-1 (TPU v6 lite) | bf16 | 4096³ | 2026-07-26 |
| M-007 | pallas fused softmax (lab 1.3) | vs unfused chain and XLA fused | 443.4 µs vs 349.9 µs unfused, 242.5 µs XLA | unfused max/exp/sum chain | v6e-1 (TPU v6 lite) | bf16 | 32768x512 | 2026-07-26 |
| M-008 | naive attention spill | fusions carrying the full score matrix | 3 fusion ops carry bf16[8192,8192]; 414.4 µs end to end | compiled HLO text | v6e-1 (TPU v6 lite) | bf16 | seq 8192, d 128 | 2026-07-26 |
| M-009 | flash forward (lab 3.2) | vs jax.nn.dot_product_attention and naive | 491.7 µs vs 730.0 µs reference (0.67x); naive 414.8 µs | reference attention | v6e-1 (TPU v6 lite) | bf16 | seq 8192, d 128 | 2026-07-26 |
| M-010 | attention custom_vjp grads (lab 3.3) | max |diff| vs jax.grad, both bf16 | q 0.163 · k 0.203 · v 0.0 | autodiff of reference | v6e-1 (TPU v6 lite) | bf16 | 512x128 q, 1024x128 kv | 2026-07-26 |
| M-011 | causal flash, blocks skipped (lab 3.4) | dense vs causal, same kernel family | 473.4 µs dense vs 330.0 µs causal (1.43x) | dense flash | v6e-1 (TPU v6 lite) | bf16 | seq 8192, d 128 | 2026-07-26 |
| M-012 | pallas tiled matmul, retuned | dimension_semantics + block sweep vs XLA | 384.6 µs vs 353.0 µs (1.09x), block 2048x1024x512 | jnp matmul | v6e-1 (TPU v6 lite) | bf16 | 4096³ | 2026-07-26 |
| M-013 | pallas fused softmax, retuned | rows=1024 blocks vs unfused and XLA fused | 222.7 µs vs 379.5 µs unfused, 267.6 µs XLA fused | unfused max/exp/sum chain | v6e-1 (TPU v6 lite) | bf16 | 32768x512 | 2026-07-26 |
| M-014 | flash vs naive, scaling curve | speedup by sequence length | 1.05x at 4k, 1.22x at 8k, 1.23x at 16k, 1.35x at 32k; 4.15x vs reference at 32k | naive attention | v6e-1 (TPU v6 lite) | bf16 | seq 4k-32k, d 128 | 2026-07-26 |
| M-015 | windowed attention, mask as loop bound | vs XLA masked and dense flash | 272.2 µs vs 1326.9 µs XLA masked (4.87x), 834.2 µs dense (3.06x) | XLA masked attention | v6e-1 (TPU v6 lite) | bf16 | seq 16384, window 1024 | 2026-07-26 |
| M-016 | ragged batch, lengths bound the loop | vs padded XLA at mean length 41% of max | 397.6 µs vs 788.8 µs (1.98x) | padded XLA attention | v6e-1 (TPU v6 lite) | bf16 | 8 x 4096, d 128 | 2026-07-26 |
| M-017 | int8 dequant fused in kernel | vs XLA dequant matmul and bf16 weights | 197.9 µs vs 187.3 µs XLA (0.95x, XLA keeps this one), 230.3 µs bf16 | XLA dequant matmul | v6e-1 (TPU v6 lite) | bf16/int8 | 8x4096x14336 | 2026-07-26 |
| M-018 | naive attention byte accounting | round-trip estimate vs XLA cost analysis | est. 276.8 MB vs 155.2 MB bytes accessed (0.56x): the models disagree; profiler confirm stays open | 2 x S^2 x 2 + io | v6e-1 (TPU v6 lite) | bf16 | seq 8192, d 128 | 2026-07-27 |
| M-019 | flash forward differential, f32 pipeline | max |flash − reference| | 2.15e-6 (bar 1e-3); bf16 pipeline 0.38% relative vs f32 reference | f32 attention reference | v6e-1 (TPU v6 lite) | f32 / bf16 | seq 8192, d 128 | 2026-07-27 |
| M-020 | attention custom_vjp grads, precision-pinned | max |custom_vjp − autodiff|, f32 | q 6.5e-6 · k 6.7e-6 · v 2.4e-7 (bar 1e-2); unpinned the same test reads 0.15, which is the MXU lesson | jax.grad of reference | v6e-1 (TPU v6 lite) | f32 | 512x128 q, 1024x128 kv | 2026-07-27 |
| M-021 | naive attention, profiled on device | the spill, hardware-timed | S-producing fusion 89.6 µs + %online-softmax custom kernel 128.0 µs per iter; XLA:TPU pattern-matched the softmax into its own online-softmax tpu_custom_call (tiling 1024x1024) | XProf device timeline | v6e-1 (TPU v6 lite) | bf16 | seq 8192, d 128 | 2026-07-27 |
| M-022 | achieved HBM copy bandwidth | 1.07 GB pure copy, median of 10 | 1287.5 GB/s = 80.5% of nameplate | 1.6e12 B/s datasheet | v6e-1 (TPU v6 lite) | bf16 | 16384x32768 | 2026-07-27 |
| M-023 | the spill, byte-confirmed | hardware time x measured bandwidth vs estimate | 217.6 µs x 1287.5 GB/s = 280.2 MB vs 276.8 MB estimated: 1.2% agreement on a 20% bar | round-trip estimate | v6e-1 (TPU v6 lite) | bf16 | seq 8192, d 128 | 2026-07-27 |
| L-001 | attention 2048x128 (xla path, LAB·X4) | median of 20, across six identical processes | 147.6 to 200.4 µs (35% spread) | the same program, no flags, same chip | v6e (Colab) | bf16 | 2048x128 | 2026-07-27 |
| L-002 | fusion flag A/B (xla path, LAB·X2) | RETRACTED: the flag was a no-op, both modules identical | 187.5 vs 139.3 µs | superseded by L-001's process spread | v6e (Colab) | bf16 | 2048x128 | 2026-07-27 |
| L-003 | HBM refusal (xla path, LAB·X2) | compile-time demand against capacity | 335.31G temporaries vs 31.24G HBM | the chip's own report | v6e (Colab) | f32 | 300000² | 2026-07-27 |
| L-004 | MLP training loop (pytorch path, LAB·P4) | final loss after 200 steps, torch_xla bridge | 0.00370 | 0.00372 on CPU torch 2.2.2, same seed | v6e (Colab) | f32 | 256x8 to 32 to 1 | 2026-07-27 |
| L-005 | checkpoint and resume (pytorch path, LAB·P4) | loss at the kill step, then the first resumed step | 0.01456 at step 99 to 0.01437 at step 100 | a fresh model on this problem starts near 1.0 | v6e (Colab) | f32 | 256x8 to 32 to 1 | 2026-07-27 |
| L-006 | attention 2048x128, tpu_fusion on against off (LAB·X4) | twelve alternating pairs, analysed as pairs | fusion faster in 10 of 12; median paired difference 16.7 µs | tpu_fusion disabled; marginal ranges overlap, pairs do not | v6e (Colab) | bf16 | 2048x128 | 2026-07-27 |
Your chip against the record
Ran a lab? Paste the results blob its final cell printed. Parsed in your browser; nothing you paste is sent anywhere.
waiting for a results blob ·
Parsed in your browser. Nothing you paste here is sent anywhere.