Path G: decode-throughput optimization (target +50–100 %)

[ai-hpc]

Goal

Push decode tok/s from alpha.10's ~9.1 toward the bandwidth ceiling. Orin's LPDDR5 peak is 102 GB/s; reading the 2.4 GB Qwen3-4B Q4_K_M model per decode step gives a theoretical ~42.5 tok/s. We're at 21 % of bandwidth today. Path G is the incremental optimization track to close that gap.

Mirrors the structure of Path E (#33) and Path F (#45): small phased PRs, each validated, with verified-result comments.

Where we are

Decode has been stuck at 9.1 tok/s since alpha.5 (Path C shipped Q4_K uint32 weight loads on Wo / gate-up / QKV). Every release since (alpha.6 D, alpha.7+8 E, alpha.9+10 F) has been prefill or persistent-KV work — none touched the decode hot path.

Pre-Path-C nsys put 93 % of decode in 5 K-quant GEMV kernels. Path C vectorized 4 of them (the Q4_K paths). What's left, estimated:

Cost	Approx % of decode wall
W_down (Q6_K, M=2560 K=9728)	~30 %
Wv (Q6_K row in QKV triple)	~10 %
Wo / gate / up / Wq / Wk (Q4_K, already Path C-optimized)	~30 %
Attention decode (KV reads)	~15 %
RMSNorm + RoPE + final norm	~5 %
Final logits gemv (2560 × 151936)	~5 %
Embedding lookup + sample	~5 %

G1 will produce a fresh nsys profile to confirm or correct these estimates. Every G2+ phase is conditional on what G1 finds.

Phased plan

Phase	Scope	Estimated win
G1	Fresh nsys profile of alpha.10 decode. Half-day. Zero risk. Identifies hot kernels.	0 (validation)
G2	Q6_K uint16 weight loads. The Q6_K alignment problem from #29 was for uint32 (210 B not divisible by 4). uint16 works at 2-byte alignment, which 210 B satisfies. 2× fewer load instructions on the byte path — mirrors Path C's uint32 win, scaled down. Hits W_down + Wv.	+5–10 % decode
G3	Attention KV-tile reuse in shared memory. Flash decode currently reads KV from global mem per head; could stage one layer's K/V in shared and reuse across heads.	+5–10 %
G4	Fuse RMSNorm + first projection. Separate launches today for fused_rmsnorm then Wq/Wk/Wv. Merge into one launch to amortize per-layer dispatch overhead.	+3–5 %
G5	INT8 KV measurement + default-flip. --int8-kv already supported; never measured at alpha.10. Halves attention-read bandwidth. If output quality holds on a longer context, flip default.	+5–15 % for long context
G6	alpha.11 release prep — docs + cumulative win summary.	cumulative

Cumulative target: 9.1 → 12–15 tok/s (+30–60 %) at the low-to-mid end. Beats llama.cpp's tg64 = 6.33 ± 0.25 tok/s by another notch on top of alpha.8's already-+44 % lead.

What this is NOT

• Not speculative decoding. That's the planned Path H (issue Weeks 9-11: speculative decoding + persistent KV cache #6 second half) — a separate multi-week swing for 1.5–2.5× decode, after Path G plateaus or in parallel as research.
• Not a model architecture change.
• Not multi-batch decode. Single-stream serving is the GenieClaw use case.

Phase status

• G1 — nsys profile
• G2 — Q6_K uint16
• G3 — attention KV-tile reuse
• G4 — norm + projection fusion
• G5 — INT8 KV measurement / flip
• G6 — alpha.11

Risks worth naming up front

1. The profile may surprise us. alpha.5 numbers are now nine releases old; cache behavior + scheduler changes since may have shifted where time goes. G1 might say "actually attention is the bottleneck now, not Q6_K". Plan adjusts.
2. Q6_K uint16 may not bit-equal Q6_K bytes. Float-add reordering same as Path C; expected to be FP16-ULP-bounded. Validate on the reference prompt.
3. Negative-result phases. Path C had two (Path C Phase 2: CUDA Graphs decode capture #21, Path C Phase 3: split-K K-quant GEMV (target 1.5-2× decode) #22) and Path E had one (Path E phase E3b: wire MMQ Q4_K into prefill dispatcher (JLLM_MMQ_Q4K) #37). Budget 1–2 in Path G.
4. Bandwidth ceiling is real. At ~21 % util today, the theoretical ceiling at 42.5 tok/s assumes 100 % bandwidth utilization which is never achieved in practice; realistic asymptote is more like 25–30 tok/s. Targeting the 12–15 tok/s range is honest.

Inspirations

• Path C (Path C: decode-path optimization (target 2-2.5× decode tok/s) #19, closed) for the uint32 weight-load technique and the negative-result pattern.
• llama.cpp's mmvq.cuh for split-K and cp.async ideas, if G3 needs them.

[ai-hpc]

G1 — nsys profile of alpha.10 decode (2026-05-16)

Profile setup: Jetson Orin Nano Super 8 GB, alpha.10 main (commit `6013f41`), 25 W MAXN SUPER, GPU 918 MHz, Qwen3-4B-Q4_K_M, prompt = the standard 33-token chat-wrapped, `-n 200`. Decode hit EOS at 39 tokens.

```
nsys profile --output=/tmp/gen-decode-profile --trace=cuda,nvtx,osrt --stats=true ./build/jetson-llm ...
```

Top kernels by total GPU time (across one prefill + 39 decode tokens)

Kernel	Calls	Total ms	% total	% decode-only
`gemv_quant_pair_uint32_q4k_kernel` (gate/up)	1404	1716	35.9 %	43 %
`gemv_quant_add_uint32_q4k_kernel` (Wo)	2106	810	16.9 %	20 %
`gemm_mmq_q4k_kernel` (PREFILL only)	432	514	10.7 %	—
`gemv_quant_add_typed_kernel<14>` (W_down Q6_K, scalar)	702	383	8.0 %	10 %
`gemv_q6k_f32_kernel` (final logits, Q6_K)	40	324	6.8 %	8 %
`gemm_quant_batched_q6k_kernel` (PREFILL only)	72	306	6.4 %	—
`gemv_quant_triple_uint32_q4k_kernel` (QKV variant A)	702	275	5.8 %	7 %
`gemv_quant_triple_uint32_q4k_q4k_q6k_kernel` (QKV variant B)	702	271	5.7 %	7 %
`fused_rmsnorm_residual_kernel`	5800	82	1.7 %	2 %
`flash_attention_decode_kernel`	1404	74	1.5 %	2 %

Decode-only percentages renormalize after subtracting the prefill kernels (`gemm_mmq_q4k_kernel` 10.7 % + `gemm_quant_batched_q6k_kernel` 6.4 % = 17.1 % of profile is prefill).

Surprises vs the estimate in the original umbrella body

Pre-profile estimate	Actual	Implication
W_down Q6_K = 30 %	10 %	Lever real but smaller than expected
Wv Q6_K (in QKV)	absorbed into triples (~7 % each, 14 % combined)	Path C already partially handles
Q4_K (Path C-optimized) = 30 %	70 % combined (gate/up 43 + Wo 20 + QKV 14)	Path C is still the biggest bucket — and it's already optimized; remaining wins are marginal
Attention = 15 %	2 %	At ≈ 40-token context. Will grow with long ctx but not a priority for short turns.
Final logits gemv (Q6_K) = 5 %	8 % (8 ms × 40 tokens)	Bigger than expected. Same Q6_K uint16 strategy hits this AND W_down.

gate/up pair is compute-bound, not bandwidth-bound

1.22 ms per call × 28 MB weight read = ~23 GB/s, vs Orin's 102 GB/s LPDDR5 peak. So gate/up is CPU-side dequant-compute-bound, not bandwidth-bound. Tensor cores would help if we could use them (m16n8k16 needs N ≥ 8; decode is N=1), so the obvious lever doesn't apply. Without a clever "fake batching" approach the remaining wins on gate/up are register-level inner-loop tuning — high-risk, low-confidence.

Revised phased plan

Phase	Scope	Realistic estimate
G2	Q6_K uint16 weight loads. `block_q6_K` is 210 B (2-aligned but not 4-aligned), so the Path C uint32 trick from #29 doesn't apply, but uint16 does. Hits `gemv_quant_add_typed_kernel<14>` (W_down, 10 %) and `gemv_q6k_f32_kernel` (final logits, 8 %). Combined target = 18 % of decode.	+6–9 % decode wall reduction → ~9.6–9.9 tok/s
G3	Investigate gate/up dequant scalar arithmetic (43 % of decode). High risk; standalone bench first; may regress. Skipping unless G2 shows G3-friendly findings.	best case +20 %, realistic 0–5 %
G4	INT8 KV measurement at long context (current `--int8-kv` flag).	+5 % short ctx, +15 % long ctx
G5	alpha.11 release prep + docs	cumulative

Honest cumulative target — revised

alpha.10 9.1 → ~10.5–11 tok/s (+15–20 %) if G2 + G4 hit. Not the original +50 % — the profile shows decode is mostly already-optimized Q4_K. Bigger gains would need Path H — speculative decoding (1.5–2.5×), which is a multi-week separate effort tracked from the Path G epilogue.

Starting G2 next.

[ai-hpc]

G2-lite negative result — compiler was already doing it

PR #60 (closed): explicit d × sc[is+i] precompute hoist in dot_q6k_row measured −0.09 tok/s (20 ms slower across 40 decode tokens; 3 samples each, σ ≈ 1 ms — consistent regression, not noise).

alpha.10 main:    10.00 tok/s (4017 ± 1 ms / 40 tokens)
G2-lite:           9.91 tok/s (4037 ± 1 ms / 40 tokens)

nvcc was already scheduling the equivalent dependency chain in the original. The manual hoist added register pressure (a few extra resident floats) without the offsetting arithmetic-elision benefit.

Implication: the inner-loop scalar math on Q6_K is at the compiler's local optimum. Remaining Path G levers all need structural changes:

• G2 full — Q6_K uint16 with lane re-partition. Real layout change, 1–2 days careful design + validate, projected +6–9 % decode. The simple version doesn't apply because Q6_K's data layout doesn't naturally split to uint16 aligned loads without re-thinking the K-position partition.
• G3 — attention KV staging in shared mem. Only matters at long context (attention is 2 % at N=40, would grow at N=500+).
• G4 — INT8 KV measurement (already supported via --int8-kv flag). Pure measurement, ~1 hr.

Recommendation: skip G2-full unless decode tok/s becomes a hard blocker on a specific workload. The dispersion-of-effort risk vs the ~1 tok/s projected win is poor. G4 INT8-KV measurement remains the cheapest information-rich next step, but it might shift the decision toward leaving Path G at alpha.10 and investing in Streaming (#5) for the user-perceived win instead.

Latest alpha.10 decode baseline confirmed today: 10.0 tok/s (with new same-day measurement, slightly above the 9.1 documented earlier — thermal / power-state difference).