Expand description
FP16-MIRROR PREFILL (MEMRA_PP_F16=1): cuBLASLt FP16 TN GEMM on a resident fp16 dequant mirror of the Q8_0 trunk weights.
Probe verdict 2026-07-26 (tools/bench_lt_f16.cu on the H100 box): 611-687 TF at the 9B m=512 prefill shapes vs the vendored MMQ per-shape medians = 3.2-3.7x per launch (MMQ = 60% of prime). Why fp16 and not faster int8: the exact wgmma arc proved Q8_0’s per-32-block scale fold serializes Hopper’s warpgroup MMA pipe (ptxas C7514, ledger’d); fp16 f32-accumulate has no mid-loop accumulator reads and streams at tensor-core rate.
NUMERIC CONFIG (new, explicit, opt-in — MEMRA_PP_FP8/GDN-chunked precedent): the int8 part of the dequant is exact in fp16 (7 mantissa bits into 11); rounding enters at d*q products and the activation f32->fp16 cast. run-gen argmax battery + kernel-check tolerance gate arbitrate. Decode (m<16) keeps the Q8_0 dp4a/MMVQ chain untouched — decode==verify law holds.
VRAM: the mirror duplicates every 2D Q8_0 projection at 2 B/w (9B model ~+17GB) — an 80GB H100 lane feature. MEMRA_PP_F16_BUDGET_MB (default 32768) caps the spend, layer-order prefix.
Structs§
- F16Scratch
- Resident scratch (fp8_ffi::Fp8Scratch pattern): fp16 activation (grown to the largest m*k seen) + the cuBLASLt workspace. Single GPU worker; the Mutex guards lazy build/grow only.
Functions§
- pp_
f16_ enabled - MEMRA_PP_F16 gate, read once. DEFAULT ON on the Hopper lane (80GB — the mirror costs 2 B/w, ~17GB on the 9B; the box carries it), opt-in elsewhere; =1/=0 overrides either way. Promotion battery (2026-07-26, H100): kernel-check ALL GREEN (f16 rel <= 6.5e-3, band 1e-2); run-gen argmax MATCH on p1/p2/p3 long prompts; greedy streams IDENTICAL to the MMQ config on all three; pp512 8674 -> 15626 tok/s (+80%, N=5 medians). Decode untouched (m>=16 arm).