Commit 5258b80
committed
docs+scratch(fp8): FP8 activations e2e design + gb10 fp8-vs-bf16 GEMM microbench
Track 2 (fp8 activations end-to-end) — design-only round, nothing wired into
the live pipeline.
docs/FP8-ACTIVATIONS-PATHC.md: the thesis (hold activations in fp8 e4m3, run all
GEMMs/kernels in fp8, in-place fwd+bwd), tensor-precision map, scaling strategy
(MXFP8 E4M3 block-32 primary + per-tensor delayed-amax secondary, E5M2-grads
rejected), in-place fwd+bwd choreography, memory-halving math incl the
bs4-fp8==bs1-fp16 cross-track identity, numerical guardrails (fp32 accumulate +
the §17 dD aligned-gold lesson generalized to fp8), bs4 dominant-GEMM shape table
(M=16384), microbench spec, phased plan, sources.
scratch/fp8_gemm_microbench.py: standalone gb10 probe measuring fp8(e4m3) vs bf16
TFLOPs at the four prod bs4 transformer GEMMs + one SSD tile. R1 = TE
te.Linear under MXFP8BlockScaling/DelayedScaling(E4M3) in te.fp8_autocast; R2 =
in-house fp8_matmul_path_c cooperative T.gemm (records the Metal-only CUDA-emission
gap honestly). Reuses fp8_amax.py (fp8_amax_tilelang/fp8_quantize_tilelang) for
the per-tensor quantize reference. Emits a machine-parseable RESULT_JSON line
(per-shape bf16/fp8 TFLOPs + speedup + fp8-vs-bf16 byte sizes + every route incl
recorded gaps) plus a pretty block.
RULE #1: --prod required (no fake shape); RAISES on NaN/Inf output and zero-norm
ref; absent/blocked/Metal-only routes are RECORDED with where+what, never silently
skipped; no fp8->bf16 masquerade. Local check: py_compile + ast.parse OK, --prod
gate exits 2 without --prod and fails-loud under --prod off-CUDA, all 3 imported
cppmega_mlx symbols resolve, JSON schema round-trips.1 parent 29fe04f commit 5258b80
2 files changed
Lines changed: 740 additions & 0 deletions
0 commit comments