benchmarks for prefill and decode speeds

[Avalonec]

Hi!
Have you run any dedicated throughput benchmarks for prefill and decode speeds using the hq-e8-2b KV codec? The memory density is phenomenal (~7x over bf16), but I'm curious if the E8-lattice decoding and Rice coding overhead noticeably impacts the attention kernels (tok/s) compared to standard nvfp4 or rk2v4-e8.

[cometkim]

Hey @Avalonec,

The HyperQuant implemetnation trades off the decoding speed for sure. Here some results on my side. My env right now doesn't produce ideal perf state so you need to read A/B

Prompt tokens	Comparator KV	Prefill tok/s (cmp → hq)	Decode tok/s (cmp → hq)	KV GiB (cmp → hq)
8192	int8	10066 → 9640 (−4.2%)	85.3 → 80.9 (−5.2%)	0.266 → 0.106
8192	k8v4	9984 → 9658 (−3.3%)	85.6 → 81.1 (−5.2%)	0.202 → 0.106
8192	nvfp4	9863 → 9583 (−2.8%)	85.7 → 81.1 (−5.4%)	0.145 → 0.106
262144	int8	3083 → 2247 (−27.1%)	59.2 → 37.7 (−36.3%)	8.26 → 2.29
262144	k8v4	2542 → 2234 (−12.1%)	63.5 → 37.5 (−40.9%)	6.29 → 2.29
262144	nvfp4	2319 → 2234 (−3.7%)	64.0 → 37.6 (−41.2%)	4.50 → 2.29

The decoding performance degrades in proportion to the context size.

For short-to-medium contexts hq-e8-2b costs ~5% decode for 2.5× KV density over nvfp4 (and ~2.5–5× over the others); for very long contexts prefill stays near nvfp4 KV parity but decode pays roughly 36–41%.

I don't consider this to be the fully optimized result yet. My goal is to restore peak performance in long-context scenarios to 80% of the level achieved with the original int8 KV's perf.

[Avalonec]

@cometkim Hey! Great work - the density numbers speak for themselves (2.29 GiB @ 262K is ~2× over nvfp4 and ~3.6× over int8), and what surprises me most is how well prefill held up: −3.7% vs nvfp4 at 262K is essentially parity given you're adding a full lattice codec into the KV path.
The shape of your A/B is very informative, by the way: ~−5% decode at 8K across all comparators, growing to −36…−41% at 262K, i.e. scaling with context. That's exactly the signature of the E8 decode being compute/latency-bound rather than bandwidth-bound: in compute-bound regimes (prefill; short-context decode where weight traffic dominates the step) the lattice decode hides under the MMA pipeline, and it only gets fully exposed once the KV pass dominates - which is precisely long-context decode. You're moving 3.6× fewer bytes than int8 yet running slower, so bandwidth is clearly not the limiter; the bottleneck is localized on the producer/decode side of the attention kernel. That's good news: it means the 80% goal is a kernel-microarchitecture problem, not a compression-ratio tradeoff.
A few things in a similar spot (same sm_120a part, warp-specialized NVFP4 KV attention with 12 producer / 4 consumer warps and an mbarrier pipeline):
Rebalance producer/consumer warps for the heavier codec. The split that saturates consumers for nvfp4 dequant may starve them once decode cost per byte goes up; sweeping producer count is cheap.
LUT-based lattice decode. A 256-entry codebook in const/shared memory plus PRMT/half2 unpacking removed most of per-byte decode cost.
Staging tricks. Ping-pong slots and fetching the scale/code box on every other K-tile (TMA's 16 B minimum box covers two tiles anyway) cut producer traffic noticeably; the same should apply to E8 planes.
80% of int8 decode at 262K (~47 tok/s from today's 37.6) looks very reachable if the loss is producer-side: that's ~10 tok/s to recover without touching your density. Happy to help however useful - review the attention kernel, run A/B on 5090.
Keep it up!

[cometkim]

Thanks. It seems worth a try. I can't try it right away because I'm currently using all my resources for quality measurements, but I'll revisit it soon.