the numbers

Benchmarks, and how to read them

Measured against triton-windows (real Triton) and torch (which calls NVIDIA's hand-tuned cuBLAS/cuDNN libraries) on the same machine: an RTX PRO 5000 Blackwell laptop GPU.

Identical kernel source (modulo nl/tl) and identical config sweeps; medians over 50 reps with the L2 cache flushed between reps. The machine is a 110 W laptop that thermally throttles under sustained load, so suites start from a similar thermal state and within-run columns are the fair comparison. Rerun everything with python benchmarks/bench.py --cooldown 300; the raw tables live in benchmarks/results.md.

the bandwidth roofMemory-bound kernels: parity

softmax, layernorm and vector add bandwidth: newt, triton and torch within a few percent of each other — There is nothing left to win: softmax reads and writes each byte once, and coalesced + vectorized + fused saturates the bus. Any correct compiler ties. This is the roofline model's bandwidth roof.

the compute roofTensor-core matmul

fp16 matmul sustained TFLOP/s across sizes: newt 67 to 83, triton 81 to 119, torch 68 to 103

the journey at 4096 cubed: 63.3 to 79.5 to 81.7 sustained, 109.7 cold vs triton about 120

matmul fp16 (TFLOP/s)	1024³	2048³	4096³	8192³
newt v0.1 (WMMA, sync staging)	39.1	69.1	63.3	62.8
+ cross-iteration cp.async ring	63.7	70.6	79.5	70.1
+ mma.sync / ldmatrix / swizzle + N stages	67.2	82.7	81.7	77.0
triton-windows (same run)	81.2	101.1	119.0	100.8
torch / cuBLAS (same run)	67.8	103.1	100.3	95.0

That is 76-83% of Triton sustained; cold-start single-kernel runs reach 109.7 TFLOP/s, ~92% of Triton's own cold numbers. The remaining gap is Triton's finest-grained machinery: strength-reducing address computations across loop iterations and specializing warps into producer/consumer roles. Both are known, documented, and out of scope for a nano on purpose. tf32 (a 19-bit float format used for fp32 matmuls on tensor cores) still uses NVIDIA's higher-level WMMA API and sits near 45%: mechanical future work.

scopeWhat was deliberately left out

Random number generation inside kernels, device-side printing, calling one @jit function from another, non-NVIDIA backends, fp8 formats, Helion's larger search space (loop reordering, persistent kernels), and the last few percent of matmul scheduling described above. Each omission is documented where a user would hit it. None of them changes the ideas this project exists to demonstrate: the modern GPU kernel stack, from tile-level Python down to tensor-core machine code, fits in four thousand readable lines once you know which problems are essential and which are incidental.

How newt gets these numbers Get started + docs