the path · 0/15
start the path
the bench

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.

records
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
compare

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.