H100-KERNEL-5X-ROADMAP — how fak’s own CUDA kernel gets 5–10× faster on Hopper

The honest frame up front. Every speedup number below is a projection gated on a measured Hopper run — this host has no NVIDIA GPU / CUDA toolkit, so the figures here are derived from the measured H100 baseline + the code that is already in tree, not from a new run. The baseline itself is real and measured (GCP-H100-RESULTS.md). This document PLANS the next steps and ships the first executable one (the apples-to-apples Q8 device-decode bench row); the kernel changes it scopes are tracked, GPU-gated follow-ons.

The measured baseline (Qwen2.5-3B-Instruct Q8_0, single-stream, 1× H100 80GB, sm_90)

Engine Precision Prefill tok/s Decode tok/s
llama.cpp CUDA Q8_0 19,310.5 361.6
fak-cuda f32 51.0 96.3
fak-cpu (pure-Go) Q8_0 109.7 15.7

Two gaps, with very different shapes:

Why the decode gap is bandwidth, not arithmetic

Single-stream decode is a stack of GEMVs (one activation row × each weight matrix). GEMV is memory-bandwidth-bound: the cost is dominated by streaming the weight bytes, not the multiply-adds. fak-cuda decode runs f32 weights (4 bytes/weight); llama.cpp runs Q8_0 (1 int8 code + a thin per-block f32 scale ≈ ~1.03 bytes/weight). So fak streams ~3.9× more bytes per token than llama.cpp — which lands almost exactly on the measured 3.75× decode gap. The repo’s own CPU head-to-head reaches the same verdict from the other side: “decode is fundamentally memory-bandwidth-bound (streaming Q8 weights)” (LLAMACPP-HEADTOHEAD-RESULTS.md).

The correction this roadmap makes to the record: the H100 results note and the bench harness both said fak-cuda runs f32 because “the cuda backend does not advertise UploadDtype.” That is stale. The CUDA backend advertises UploadDtype: true (internal/compute/cuda.go:450) and already implements native Q8_0, Q4_K, F16, and AWQ device GEMMs (uploadQ8/uploadQ8Resident/uploadQ4K/uploadF16 in cuda.go; k_q8_gemm/k_q4k_gemm in internal/compute/cuda_kernels.cu). The f32 H100 number was not a missing capability — the bench simply never requested Q8. modelbench routes the HAL through the Q8 device path whenever -quant/-lean is set against a backend that advertises UploadDtype (cmd/modelbench/main.go:852, gate at :447).

The levers, ranked

Ranked by (expected multiplier × confidence ÷ cost). File anchors are exact.

Lever 1 — Q8 device weights for decode · ~3.9× decode · HIGH confidence · SHIPPED (wiring); GPU-run pending

What. Run fak-cuda decode on resident Q8_0 weights (int8 codes + per-block f32 scales, native k_q8_gemm GEMV) instead of f32. Streams ~1 byte/weight instead of 4.

Where. Already implemented end-to-end: cuda.go uploadQ8Resident → k_q8_gemm (cuda_kernels.cu:336), HAL Q8 routing at cmd/modelbench/main.go:852. Off-GPU cosine witnesses exist (#485, cuda_quant_test.go, tools/run_485_acceptance_on_gpu.sh, floor cudaQ8CosineMin = 0.999).

Shipped here. tools/gcp_bench.py now has a fak-cuda-q8 engine (-lean -backend cuda) — the apples-to-apples row vs llama.cpp Q8_0. It is opt-in until a green Hopper run witnesses the device Q8 GEMV, then it promotes into the default all set.

Next checkable step. On a 1× H100: python tools/gcp_bench.py --tier a3-high-h100-1g --spot --engine llama,fak-cuda,fak-cuda-q8. Expectation: fak-cuda-q8 decode ≈ 300–375 tok/s (≈ llama.cpp Q8 parity), i.e. ~3.9× over the 96.3 f32 row — and a first on-hardware correctness pass of the device Q8 GEMV. This run also tells us whether decode is now launch-overhead-bound (Lever 2).

Lever 2 — Reusable “replay-many” CUDA graph for decode · ~1.5–2× on top of Q8 · MED-HIGH · tracked (#35/#3), GPU-gated

What. A decode step issues ~500–700 kernel launches (≈30 layers × {3–5 GEMVs, 2 RMSNorm, RoPE, flash-attn, SwiGLU, 3–4 adds, 2 KV writes} + argmax). Capturing that op stream into a CUDA graph and launching it once per token collapses ~600 launches into

  1. The catch already learned in-tree: per-token re-capture is a no-win — re-recording a ~600-node graph every token costs about what the 600 launches cost (cuda.go:44). The real win is capture once, replay many.

What already exists. The instantiate-once machinery is built: fcuda_graph_end_launch keeps g_exec and uses cudaGraphExecUpdate rather than re-instantiating (cuda_kernels.cu:943-968); the KV write is a scalar-offset kernel (k_copyrow) so the exec is patchable as the cache grows (cuda_kernels.cu:535); #969 pool pre-warm makes capture allocation-free; FAK_CUDA_GRAPH=1 gates it; cudaKVMaxPos fixes KV capacity so no realloc happens during capture.

The missing piece. Make the decode graph length-agnostic so it is captured ONCE per session and cudaGraphLaunch-ed every subsequent token with zero per-token CPU/capture work. That needs pos/nPos to be device-resident scalars the kernels read, instead of host launch-params that change every token (which is what forces re-capture today). The kernels to convert: k_rope (reads pos), k_flash_attention (reads nPos), the KV-write offset, k_argmax. Anchor for the position counter: cuda_kernels.cu:535-544 (the scalar-offset write is already the template).

Why it compounds with Lever 1. Once Q8 cuts weight traffic ~4×, the per-launch overhead becomes the dominant residual: ~600 launches × ~3–5 µs ≈ 2–3 ms/token of pure launch latency, i.e. a ~330–500 tok/s ceiling sitting right at llama.cpp’s number. Removing it is what takes fak past parity, not just to it. The repo’s own RTX-4070 note already shows fak reaching decode parity with llama.cpp Q8 using a reusable graph, at f32 precision (LLAMACPP-HEADTOHEAD-RESULTS.md intro; GPU.md §3b) — so “graph + Q8” is the combination that should clear it on Hopper.

Next checkable step. First, Lever 1’s run with FAK_CUDA_GRAPH=1 to re-confirm the per-token-capture no-win on H100 (cheap, no code). Then implement device-resident pos/nPos and re-measure; success = decode tok/s rising materially above the Q8-only number with the forward still bit-faithful to cpuref.

Lever 3 — Fix the prefill amortization defect · large (10–100× on prefill) · MED · needs a phase profile first

What. fak-cuda prefill (51 tok/s) being slower per token than decode (96) is a structural defect: prefill should stream each weight once for all P=512 tokens. The GEMMs already use cublasSgemm/cublasGemmEx (cuda_kernels.cu:210, :274), so the defect is not GEMM quality. The likely culprits, in order: per-op overhead with no prefill graph (every op a separate launch + a devTr cudaMalloc), per-op stream serialization, or a prefill attention path whose cost is not being amortized across the batch.

Next checkable step (diagnose before fixing). The flag already exists: modelbench-cuda -gguf <q8> -lean -backend cuda -phase-profile emits per-phase ms (cmd/modelbench/main.go runPrefill → phaseTable). Run it on H100, read which phase dominates a 512-token prefill, and fix that phase specifically. This is the single largest raw-number headroom and matters most for long-prompt agentic workloads (big system prompts, tool outputs).

Lever 4 — Tensor-core / TF32 prefill · large on prefill · HIGH (that it helps) · SHIPPED (TF32 wiring); GPU-run pending

What. The f32 SGEMM runs on Hopper’s FP32 CUDA cores, leaving the tensor cores idle. The compute-bound prefill phase is exactly where routing those GEMMs onto the tensor cores pays.

Shipped here (TF32). A FAK_CUDA_TF32=1 toggle now routes the existing f32 SGEMM (fcuda_matmul_f32 / cublasSgemm, cuda_kernels.cu) through the tensor cores at TF32 input precision with F32 accumulation — cublasSetMathMode(g_blas, CUBLAS_TF32_TENSOR_OP_MATH) via the new fcuda_set_tf32 ABI (cuda_backend.h), read once at init in cuda.go (tf32Enabled) and exposed for a host to flip post-init through the cross-build EnableCUDATF32() seam (tf32_cuda.go/tf32_nocuda.go, the twin pattern of EnableCUDAGraph). It is default-off, so the witnessed device-vs-cpuref cosine floors hold unchanged on the pedantic FP32-core path; TF32 keeps the f32 exponent, so only the mantissa narrows (a small, disclosed precision cost). The fak-cuda-tf32 bench engine (tools/gcp_bench.py) is the apples-to-apples prefill row: fak’s f32 device path with FAK_CUDA_TF32=1 vs the pedantic-FP32 fak-cuda row and llama.cpp. It REUSES the modelbench-cuda binary fak-cuda builds (guarded) and keeps the -require-non-reference honesty gate. Opt-in until a green Hopper run witnesses the prefill gain. Host-free coverage: tf32_enable_test.go (the seam stays callable in the default non-cuda build) + the gcp_bench_test.py engine/render tests.

The F16 row is NOT yet bench-wiring-sized (correction to the record). The F16 device GEMM (uploadF16/fcuda_matmul_f16, floor cudaFP16CosineMin = 0.997) exists at the compute layer (#484), but the modelbench Session forward only routes uploads as Q8 (s.Quant) or Q4_K (s.Q4K) — there is no F16 upload-dtype path threaded through matWeightHAL. So an fak-cuda-f16 engine needs a Session-level F16 routing change first (a new device-dtype select), not just a bench fragment — tracked as a separate not yet. TF32 is the genuinely zero-Session-change Lever-4 win, which is why it shipped first: it retunes the existing f32 GEMM and needs no new forward path.

Next checkable step. On a 1× H100: python tools/gcp_bench.py --tier a3-high-h100-1g --spot --engine llama,fak-cuda,fak-cuda-tf32. Expectation: fak-cuda-tf32 prefill tok/s rises materially over the pedantic-FP32 fak-cuda prefill row (tensor cores vs FP32 cores), with the forward still inside the Approx cosine gate. Pair it with Lever 3’s -phase-profile to confirm the prefill phase is the one that moved.

The math to 5–10×

Path Levers Compounded vs current f32 un-graphed fak-cuda
Decode Q8 (×3.9) → past parity with graph (×~1.5–2) ~6–8× → clears llama.cpp Q8
Prefill amortization fix + tensor-core/TF32 10–100× (its own large, separate headroom)

So the 5–10× is reachable and decomposed: decode via Q8 + the replay-many graph; prefill via the amortization fix + tensor cores. None of it requires new silicon — it is the same “tuning, not architecture ceiling” boundary the CPU and 4070 head-to-heads already identified, now on Hopper.

What shipped in this increment

The remaining levers are a GPU-gated not yet: the code paths exist or are scoped, but the measured Hopper number is the witness and this host has no NVIDIA GPU. The witnesses to collect, in order: Lever 1’s --engine llama,fak-cuda,fak-cuda-q8 (Q8 decode parity), Lever 4’s --engine llama,fak-cuda,fak-cuda-tf32 (TF32 prefill gain), Lever 3’s -phase-profile (which prefill phase dominates), then Lever 2 (device-resident pos/nPos for replay-many).

Reproduce / drive the next run

# apples-to-apples Q8 decode head-to-head on a 1x H100 (spot), then teardown
python tools/gcp_bench.py --tier a3-high-h100-1g --spot \
    --engine llama,fak-cuda,fak-cuda-q8

# Lever 4: the TF32 tensor-core PREFILL head-to-head (f32 weights, TF32 SGEMM math)
python tools/gcp_bench.py --tier a3-high-h100-1g --spot \
    --engine llama,fak-cuda,fak-cuda-tf32

# diagnose the prefill defect (Lever 3) on the same box, if --keep is used
modelbench-cuda -gguf <qwen2.5-3b-q8_0.gguf> -lean -backend cuda -phase-profile
# and confirm TF32 moved the prefill phase specifically:
FAK_CUDA_TF32=1 modelbench-cuda -gguf <qwen2.5-3b-q8_0.gguf> -backend cuda -phase-profile