Glyd
GitHubGet started

On the GPU

What is on the GPU

A bf16 weight is a sign bit, 8 exponent bits and 7 mantissa bits. In a trained model the sign and mantissa are noise and the exponent carries about 2.6 bits of information (why). Every format keeps the sign and mantissa byte as it is and codes the exponent:

Format The exponent Bits a weight Decode
mma, the tiered layout 2-bit digits in tiers: the tensor’s 3 commonest exponents, the next 3, then the next 3, then the exponent itself, in the order the tensor cores take their operand 10.80 in registers, straight into the tensor cores
mma12, the 12-bit layout a 4-bit code into the tensor’s 15 commonest exponents, the rest in a step’s exception list 12.04 three byte permutes per four weights
fast a 3-bit code into the tensor’s 7 commonest exponents, an escape to the exponent itself 11.25 bit operations, every weight in parallel
huffman a per-tensor prefix code, 32 streams a tile 10.88 a lane decodes its stream in turn

Bits a weight on Qwen2.5. No code that sees each tensor’s exponents on their own gets below about 10.6 bits a weight. The two layouts are the ones this site measures; fast is e2e.py’s default format.

For a generated token (a few tokens a step) the product is fused: the kernel reads the packed weights, decodes them in registers and multiplies, and bf16 weights are never written out. Prompts run a tiled kernel that decodes each weight once for many tokens. On Hopper, steps of 17 to 128 tokens in the 12-bit layout go through the copy engine (TMA) and wgmma; on GDDR Ampere and Ada, steps of 17 to 64 tokens through the same plan with that generation’s instructions. Every product sums in a fixed order: the same result every run.

Two layouts, and which GPU runs which

The tiered layout takes the most off and costs the most arithmetic to decode. Where memory is the limit, as on an RTX 4080 SUPER at a few tokens a step, it is also the faster. Where the GPU’s memory outruns the decode (an H100’s HBM3, or many tokens a step), the 12-bit layout keeps up. GPU time a generated token, Qwen2.5-7B-Instruct, q, k, v and gate, up merged, at 1 / 8 / 32 / 64 sequences a step:

GPU bf16 tiered (mma, 10.80 bits) 12-bit (mma12, 12.04 bits)
RTX 4080 SUPER 16 GB 21.97 / 22.85 / 26.19 / 27.67 ms 16.52 / 17.37 / 18.80 / 25.14 17.48 / 18.25 / 19.67 / 21.66
A10 24 GB 33.71 / 34.77 / 35.50 / 37.98 ms 24.33 / 25.61 / 35.99 / 49.66 24.40 / 25.76 / 29.03 / 33.05
A100 40 GB 14.18 / 14.87 / 16.18 / 17.80 ms 15.34 / 16.00 / 23.99 / 28.81 11.84 / 13.77 / 17.16 / 19.95
H100 SXM 80 GB 6.90 / 7.50 / 8.02 / 8.55 ms 6.55 / 7.24 / 8.69 / 9.99

On Ada the tiered layout, the smallest, is also the fastest to 32 sequences; on an A10 the 12-bit one is 13 to 28% under bf16’s time at every count; on an A100 and an H100 the 12-bit one is under bf16’s time to 8 sequences and 6 to 17% over at 32 and 64, where their small matrices cost more than bf16’s. glyd_gpu.best_layout(), and e2e.py --format auto, take the tiered layout on Ada and wherever only it fits, the 12-bit one elsewhere. The site’s benchmarks show each GPU with the layout it picks.

e2e.py, option by option

python e2e.py MODEL_DIR [options], from the repository’s gpu/ folder:

Option What it does
--format fast (the default), huffman, mma (the tiered layout, up to 64 tokens a step multiplied straight from it), mma12 (the 12-bit layout) or auto (the one for this GPU)
--fused multiply straight from the packed weights
--baseline run the model in bf16 too, in the same process, and compare
--batch 1,8,32 generate for this many copies of the prompt at once, each size measured (default 1)
--tokens 128 new tokens to generate (default 128)
--merge the Linears that take the same input (q, k, v; gate, up) as one product each, for bf16 and Glyd alike, as serving engines run them
--profile 16 GPU time by kernel over this many generated tokens at each --batch size: a step after the prompt, and the prompt’s
--prefill 64,512,2048 prompt lengths to time one forward pass at
--ppl FILE perplexity over windows of --ppl-window tokens (default 64) from the file’s 10th MB, and how often the next-token choice is bf16’s
--mmlu 300 MMLU accuracy over this many questions, 0-shot, by the answer letter’s logit
--smi PREFIX once the model is on its GPUs, nvidia-smi’s report into PREFIX-bf16.txt and PREFIX-glyd.txt
--gpus 2 spread the layers over this many GPUs (bf16 by accelerate’s device map, Glyd balanced by packed size)
--gpu-mem 40 GiB of bf16 weights a GPU may hold in the baseline’s device map (default all but 2 GiB)
--kv 1024,4096 prompt lengths (tokens of enwik8, needs --ppl) to generate after with the KV cache compressed, against the plain cache
--gemm-max 64 with fast, steps of up to this many tokens multiply straight from the packed weights

Repeating a benchmark

Every number on this site comes from a run of these scripts, and every run’s log is in the repository: benchmarks/gpu. The Qwen3.8 27B run on an H100 PCIe, for one, holds the machine (machine.txt: the GPU, its memory, the driver), the environment (setup.txt: PyTorch and CUDA), the commands (run.log) and each script’s output. It ran:

python e2e.py models/Qwen3.8-27B --fused --tokens 64 --batch 1,8,32 --prefill 64,128,512,2048 \
  --ppl enwik8 --mmlu 300 --merge --format auto --baseline --gpus 1 --profile 16 \
  --smi results/smi-Qwen3.8-27B-auto-x1
  • GPU time a token (benchmarks): --fused --merge --baseline --batch 1,8,32,64 --profile 16, with --format for the layout. The time is a step’s after the prompt: 17 steps less 1, over 16.
  • Memory in use (Qwen3-32B on one 48 GB GPU): --smi PREFIX, with --gpus N where bf16 needs more than one GPU.
  • Quality (the nineteen models): --baseline --ppl enwik8 --mmlu 300.
  • Sizes: python sizes.py MODEL_DIR ..., every Linear layer’s matrix in both layouts, bit for bit.
  • The products alone: python gemm.py MODEL_DIR 1,16,64 times every product of one layer against bf16’s; python shapes.py MODEL_DIR the fused product; python check.py model.safetensors packs, unpacks and compares every tensor; python kv.py checks the compressed KV cache against the plain one.

The rented-GPU runs use scripts/gpu_lambda.sh: one Lambda Cloud instance a run, terminated at the end, its logs copied into benchmarks/gpu/.