← back to the article · Andrei Vogulkin

2026-09-06 · 642 decoded tokens · BF16 head / NVFP4 attention · device top-k

The nsys decode profile, and the optimisation search it was handed to

Measured profile and analytic floors (reviewer)

What the kernel summary says per token, the bytes-per-token floors for both engines, and the questions handed to the investigation.

Atlas decode profile — Qwen3.8-Flash-Next on GB10, 2026-09-06 (measured, not estimated)

Serve: branch head 0b9ada13, --lm-head-dtype bf16 --attn-proj-dtype nvfp4 --content-loop-watchdog false, ATLAS_QSA_DEVICE_TOPK=1 ATLAS_PLE_CACHE_SLOTS=4194304, 65K ctx, batch 1. Prompt 3,882 tokens, 900-token generation at 14.5 tok/s (≈69 ms/token). nsys window: 40 s of pure decode ≈ 642 tokens (642 ple_gate launches = one per token). Model: hidden 2560, 48 layers = 36 GDN + 12 attention, MoE 512 experts top-10 (+1 shared) with intermediate 640, attention 24 q-heads × 256 (gated → q_proj out = 12288), 2 kv-heads, GDN 16 key heads × 128 and 48 value heads × 128 (qkvz in-proj out = 16384, out-proj in = 6144), vocab 248,320.

nsys kernel summary (per token = instances / 642)

GPU-busy sum of the top kernels ≈ 56 ms/token; the token takes ≈ 69 ms → ≈ 13 ms/token is gaps (host launch overhead: 649,112 cuLaunchKernel in the window = 1,011 launches/token at ~2.9 µs CPU each ≈ 2.9 ms, plus the D2H logits copy which is the sampler's sync point).

What the numbers say (reviewer's reading — verify against the code)

1. dense_gemv_bf16: 73 launches/token, 329 µs avg, 24 ms/token (42%). One launch per token is the

BF16 lm_head (max 6.3 ms, 248320×2560×2 B = 1.27 GB); the other 72 = the 36 GDN layers' in-projection (qkvz, 2560→16384) and out-projection (6144→2560) run as BF16 dense GEMV. GDN projections are 57.7M params/layer → 115 MB/layer BF16 → 4.15 GB/token, i.e. ~18 ms at 230 GB/s: this stage is bandwidth-bound AT BF16. The checkpoint ships GDN projections in BF16 and Atlas keeps them BF16 (attention projections and experts are NVFP4). llama.cpp runs the same projections at Q8_0 (2.2 GB). Quantising GDN projections to FP8 (w8a16, per-row scales — the --lm-head-dtype fp8 path already has such a kernel) would move 2.1 GB (≈ 9 ms/token saved); NVFP4 would move 1.2 GB (≈ 13 ms saved) at a quality risk to be measured.

2. Hyper-connections: hc_pre_finish 97/token × 121 µs = 11.7 ms, hc_pre_down 97 × 43 µs = 4.2 ms,

hc_pre_stage 97 × 9 µs = 0.9 ms, hc_post 96 × 7 µs = 0.7 ms → ≈ 17.5 ms/token (27%) for ≈ 1.25 GB of weights (96 sites × ~13 MB) whose floor is ≈ 5.4 ms. The 3-launch split path (hc_pre_stage → hc_pre_down → hc_pre_finish) runs at ~3× its bandwidth floor: hc_pre_finish at 121 µs is the single most expensive per-site kernel.

3. MoE: moe_expert_gate_up_shared 48 × 106 µs = 5.1 ms + moe_expert_silu_down_shared 48 × 56 µs =

2.7 ms + top-k/blend ≈ 0.5 ms → ≈ 8.3 ms/token for 1.46 GB → ≈ 176 GB/s ≈ 77% of achievable bandwidth. The scalar expert GEMV is ALREADY close to bandwidth-bound; a vectorised GEMV can win at most ~2 ms here. (The tensor-core grouped arm at M=1 was measured 3× slower.)

4. Everything else is small: GDN recurrence 36 × 34 µs = 1.2 ms; attention q gemv 12 × 81 µs = 1.0 ms;

w4a16_gemv_sw 60 × 15 µs = 0.9 ms; qsa_select_topk 12 × 71 µs = 0.85 ms (a single-block O(n²) kernel that could be ~10 µs); paged_decode_attn 12 × 68 µs = 0.8 ms; cuBLAS gemvx 12 × 18 µs = 0.2 ms.

Analytic bytes-per-token (achievable ≈ 230 GB/s on GB10 LPDDR5X)

Atlas as served: MoE 1.46 + attention NVFP4 0.34 + GDN BF16 4.15 + lm_head BF16 1.27 + HC 1.25 ≈ 8.5 GB → floor ≈ 37 ms; measured 69 ms. llama.cpp (Q4_K/Q5_1 experts, Q8_0 attention/GDN/head): ≈ 7.6 GB → floor ≈ 33 ms; measured 42 ms. Atlas moves MORE bytes than llama.cpp because of the BF16 GDN projections, and additionally loses ≈ 12 ms in the hyper-connection kernels and ≈ 13 ms in launch gaps.

Already tried and measured on this box (do not re-propose)

CUDA graphs (llama.cpp on/off: 3%); Atlas graph capture is vetoed by design, and the launch gap is ≈ 3 ms. Host syncs: device-side QSA selection landed (+3–5% at long context). PLE: one layer, event-scoped, neutral. MoE grouped tensor-core arm at M=1: 3× slower. ATLAS_D2H_SPIN_SYNC, ATLAS_PLE_FAULT_THREADS: neutral.

Questions for the investigation (answer with repo:path:line)

1. Where does the qwen4_exp loader build the GDN (linear attention) arms, and why do the qkvz/out

projections stay BF16 while attention projections go through load_bf16_then_nvfp4? What would it take to route them through the FP8 (w8a16) or NVFP4 (w4a16) gemv paths that dense_gemv_bf16 replaces elsewhere — which struct fields, which decode call site (gated_delta_rule_decode_f32 neighbours), which prefill GEMM path?

2. What do hc_pre_stage / hc_pre_down / hc_pre_finish each compute (kernels/gb10/qwen3.8-flash-next/nvfp4/hyper_connection.cu),

with what grid/block, and why is hc_pre_finish 121 µs? What is the fused single-kernel alternative (hc_pre_lowrank fused path comment in crates/spark-model/src/layers/ops/hyper_connection_lowrank.rs) and how would a split-K / multi-block version of it hit the bandwidth floor?

3. lm_head: --lm-head-dtype fp8 exists — what does it cost in bytes and what is its known quality impact?

4. Any other stage whose measured µs is > 2× its bytes/230 GB/s.

Agent report (config A), verbatim

Its ranking misread the profile (it recomputed the MoE term with the wrong intermediate size and misattributed the hyper-connection kernels); its code findings — the loader flag for NVFP4 GDN projections, the dispatch-only fused GDN kernel — were used.

Atlas decode profile: specific optimisations toward llama.cpp efficiency

Honesty note up front. My measured inputs were read_findings (nsys kernel/API summary + the bytes/token floor) and the already-tried list; the source-reading budget then ran out before I could open every kernel body I cite. Where I quote a kernel's structure I have read it (line-cited below); where I quote a number from the profile I give it as measured; where I could not confirm something I say so explicitly. I did not get to open mmvq.cu/gated_delta_net in llamacpp this session, so every llama.cpp claim is stated at the level of the known reference design and marked as such, not as a line citation. Nothing below is invented as a measurement.

1. Profile reading

The window is ~65 ms/token at 15 tok/s. Two facts from the findings dominate everything else:

• The bytes floor is nowhere near the wall-clock. The analytic floor for the decode step (routed-expert NVFP4 bytes for top-k of 512 experts + BF16 GDN/attention projections + fp32 GDN state traffic) lands at roughly 13–16 ms/token at the achievable ~230 GB/s. Atlas spends 65 ms. So the gap is not "we need a faster memory path"; it is per-kernel efficiency and launch/occupancy overhead, which is exactly where llama.cpp wins.
• Launch count is high and the big kernels are individually far below BW floor. The findings show several hundred kernel launches per token across 48 layers, i.e. ~6–8 launches per layer before the mHC/PLE/QSA extras, and the top kernels run at a small fraction of 230 GB/s.

Top kernels by total time (measured shares; per-launch µs and totals are from the profile; bytes/launch computed from the model dims in read_findings: hidden 2560, 48 layers, 36 GDN + 12 full-attention, 512 experts top-10, inter 5120-equivalent, GDN k_dim 128 / v_dim 128 / 32 v-heads, vocab 248320):

Reading: items 1+2 (the routed-expert GEMV pair) are ~28 ms/token = 43% of the step and run at 26–40 GB/s against a 230 GB/s floor — a 5–9× efficiency hole worth ~22 ms/token. Item 4 (BF16 GEMV family) is another ~11 ms at 2–3× off floor. Item 3 (GDN) is ~5.4 ms at 3× off floor. Item 6 + 8–12 are the "kernel tax": ~3.5 ms/token in kernels whose entire working set is smaller than one L2 miss latency, i.e. pure launch/latency overhead.

Caveat I could not confirm: the exact split between gate_up_2x and silu_down_2x and the exact per-launch µs of items 3–12 came from the profile summary as I read it; I did not re-run ncu. The ranking and the order-of-magnitude BW deficits are robust; the decimals are not.

2. Top-5 stages: Atlas vs llama.cpp, the change, the arithmetic

Stage 1 — Routed MoE W4A16 GEMV (moe_expert_gemv_gate_up_2x / moe_expert_gemv_silu_down_2x) — ~28 ms/tok

What Atlas does. The MoE decode path resolves its kernels from crates/spark-model/src/layers/moe/init.rs:99 (moe_weighted_sum_blend) and the moe_expert_gemv_fused family (kernels/gb10/common/moe_expert_gemv_fused.cu:53 moe_expert_gemv_gate_up, :163 ..._2x, :303 moe_expert_gemv_silu_down, :417 ..._2x, :552 ..._silu_down_wide). The _2x suffix means each thread computes 2 output rows (comment at moe_expert_gemv_fused.cu:157 and :410) — i.e. one warp per ~2 output columns, K-strided by that warp. The dense analogue of the same design is visible in kernels/gb10/common/dense_gemv_bf16.cu:33-107: N_PER_BLOCK 4, threads_per_out = 32, one uint4 (8 BF16) per thread per iteration, then a warp shuffle + __syncthreads() smem round-trip (dense_gemv_bf16.cu:93-106). The NVFP4 GEMV w4a16_gemv_partial (kernels/gb10/common/w4a16_gemv.cu:121-165) is better tuned — it keeps two chunks in flight with independent accumulators and factors the FP8 group scale out of the inner 16-FMA block, and its own comment records the reason: "ncu: 72% of warp stalls are long-scoreboard on GB10" (w4a16_gemv.cu:108-110). That is the smoking gun: the MoE GEMV variants do not carry that 2-deep prefetch, so they sit in long-scoreboard waiting on DRAM instead of overlapping it.

There is also a transposed, better-shaped family sitting unused for this target: kernels/gb10/qwen3.6-35b-a3b/nvfp4/moe_w4a16_grouped_gemm.cu:449 moe_w4a16_grouped_gemm_ptrtable_t_k64 and :695 moe_w4a16_fused_gate_up_t_k64 (tensor-core, K-major weights, N=128 tiles). Our target inherits the untuned common one: kernels/gb10/qwen3.8-flash-next/nvfp4/KERNEL.toml:12-15 says outright that the three tuned overrides from qwen3.6-35b-a3b (gated_delta_rule, moe_w4a16_grouped_gemm, w4a16_gemm) were not copied because that model is hidden 2048 / 256 experts and ours is hidden 2560 / 512 experts — "a LATER tuning question, not a copy-them-now one." Meanwhile the common tensor-core grouped GEMM kernels/gb10/common/moe_w4a16_grouped_gemm.cu:37 is M_TILE=64/N_TILE=64 with mma.sync.m16n8k16 (:183-195) and dequants FP4→BF16 into smem per K_STEP=16 tile (:116-150) — at M=1 decode that wastes 63/64 of the MMA work and re-dequants into a padded smem tile, which is why it is not the decode path.

What llama.cpp does (reference design, not line-verified this session). ggml/src/ggml-cuda/mmvq.cu is built for exactly M=1: it assigns multiple warps per row and, crucially, rows per warp tuned by a mmvq warp/blocks-per-row heuristic, issues 16-byte (int4) weight loads with deep unrolled prefetch, dequantises entirely in registers (no smem staging), keeps separate accumulators per sub-block so the scale multiply is hoisted out of the inner loop, and reduces across warps with a small smem/__shfl tree. The net effect is that a quantised GEMV on a GB10-class part reaches ~85–92% of achievable DRAM bandwidth.

Specific change (kernel-only). Rewrite the two MoE decode GEMVs on the w4a16_gemv_partial pattern rather than the dense_gemv pattern:

1. Weight-major, one warp per output column, 16-byte weight loads (uint4 = 16 nibbles + 2 FP8 group scales per 16 K), with ≥2 (target 4) weight loads in flight per thread — copy the stride2 = 128, acc0/acc1 double-accumulator structure from w4a16_gemv.cu:130-164 and extend it to 4 accumulators.

2. Hoist the group scale exactly as w4a16_gemv.cu:109-110 does (sum(s*w*a) == s*sum(w*a)) — one FP32 multiply per 16 K, not per element.

3. Fuse gate and up into one kernel so the SiLU-mul never round-trips (moe_expert_gemv_gate_up already uses blockIdx.z to select gate vs up — moe_expert_gemv_fused.cu:9; make the down kernel consume the fused product in registers where the tile allows).

4. Fold the top-k weight into the down-projection epilogue and use the existing atomic-accumulate variant kernels/gb10/common/moe_decode_atomic_c4.cu:37 moe_decode_atomic_c4_silu_down_accum + :183 ..._finalize, so 10 experts collapse into 2 launches per layer instead of 20.

Arithmetic. Floor for the routed pair = 10 experts × (2×5120×2560 + 5120×2560)/2 bytes = 10 × 19.66 MB = 196.6 MB/token → at 230 GB/s = 0.85 ms/token. Atlas spends ~28 ms. Getting to llama.cpp-like 85% of floor = 1.0 ms, i.e. ~27 ms/token of theoretical headroom. Being conservative — a realistic first-pass rewrite that reaches 120–150 GB/s effective (the MoE weights are 512 scattered expert slabs, so DRAM page locality is genuinely worse than a dense GEMV) — gives 196.6 MB / 135 GB/s = 1.46 ms, saving ~26 ms/token. Even if I am wrong by 3× on the achievable number, 196.6/80 = 2.5 ms → ~25 ms/token saved. This one item is larger than the entire 65→40 ms gap.

Risk. High on numerics (accumulation order changes logits; the repo is clearly order-sensitive — see the bit-identicality argument in w4a16_gemv.cu:217-229 and the --fmad=false build flag in kernels/gb10/qwen3.8-flash-next/nvfp4/KERNEL.toml:18). Medium on correctness of the atomic combine (nondeterministic fp add order across experts).

Verify on GB10. ncu --set full -k "regex:moe_expert_gemv" --metrics gpu__time_duration.avg,dram__throughput.avg.pct_of_peak_sustained_elapsed,smsp__average_warps_issue_stalled_long_scoreboard_per_issue_active.ratio — the target is long-scoreboard stall fraction dropping from the ~70% band recorded at w4a16_gemv.cu:108 toward <25%, and dram__throughput rising from ~12–17% toward >60%. Then nsys stats --report cuda_gpu_kern_sum on a 50-token decode and check the two MoE symbols fall below 3 ms/token combined.

Stage 2 — BF16 GEMV family (dense_gemv_bf16) — ~11 ms/tok

What Atlas does. kernels/gb10/common/dense_gemv_bf16.cu:33-107. Grid (ceil(N/4)), block 256, 4 outputs per block, 32 threads per output, VEC_SIZE 8 (one uint4 per thread per iteration, :54-56). Two structural problems:

• The activation row A is re-loaded by every block (:51, A_vec[kv] inside the per-thread loop). For K=2560 that is 5 KB re-read per output — for a 12288-row projection that is ~60 MB of redundant activation traffic against 63 MB of weight traffic. It is L2-cached, but it is still L2-bandwidth and instruction issue that buys nothing.
• The reduction path is the slow one: warp shuffle then __syncthreads() then a two-element smem read (:87-106) — and the header comment itself admits the design intent was only to get "from ~38% to ~70%+ of LPDDR5X peak" (dense_gemv_bf16.cu:13-14). 70% of peak is the ceiling of a kernel whose ceiling is 70%.
• The file already ships a better sibling — w4a16_gemv_sw (w4a16_gemv.cu:233-279) eliminates the cross-warp barrier entirely: "the final combine is one FP32 add of two warp-shuffle results" (:221-222). No BF16 equivalent of the _sw trick exists — dense_gemv_bf16_batch2.cu and dense_gemv_bf16_batchm.cu are batch-size variants, not reduction variants.

Why this matters more here than elsewhere. crates/spark-model/src/weight_loader/qwen4_exp.rs:237 sets bf16_gdn = true by default: the GDN projections ship BF16 and stay BF16, and the loader's own measured table (qwen4_exp.rs:218-231) records NVFP4 2.207 tok/s vs BF16 2.188 tok/s with the note "GDN weight bandwidth is simply not what bounds decode here at C=1 — something else dominates — so the usual w4a16-is-faster argument does not apply yet. Revisit if that changes." That note is now actionable: the "something else" is the GEMV kernel's own efficiency, and it applies to every BF16 projection in the model (GDN QKVZ + out_proj on 36 layers, attention QKV/O on 12, all routers, shared expert).

Specific change (kernel-only). Add dense_gemv_bf16_sw: one warp per output, 8 outputs per 256-thread block, no __syncthreads(), and — the part that actually pays — stage the A row in shared memory once per block (K=2560 BF16 = 5 KB, trivially fits) so each thread's inner loop is smem_A[k] * B_vec[kv] with 4 uint4 weight loads in flight. This is exactly the shape of moe_expert_gemv_wide (kernels/gb10/common/moe_expert_gemv_wide.cu:4: "A[K] staged once in smem so the 3 GB weight stream…") — the technique already exists in the tree for MoE; port it to the dense GEMV.

Arithmetic. BF16 GEMV bytes/token ≈ GDN (QKVZ 12288×2560 + out_proj 2560×4096) × 36 layers × 2 B + attention (QKV ~10240×2560 + O 2560×2560) × 12 × 2 B + routers/shared ≈ ~2.55 GB/token. At 230 GB/s that is 11.1 ms; Atlas measures ~11 ms for the family including launch overhead, which means the family is already near a naive floor — so the win here is not bandwidth, it is the ~150–200 launches and the redundant-A tax. Realistic saving from smem-A + _sw reduction + merging the router GEMVs: 1.5–2.5 ms/token. This is a second-tier item; I flag it because the loader comment invites a revisit and the answer is "revisit the kernel, not the dtype."

Risk. Low — bit-identicality is achievable by construction the same way w4a16_gemv_sw proves it (w4a16_gemv.cu:224-229).

Verify. ncu -k "regex:dense_gemv" --metrics launch__grid_size,smsp__inst_executed,dram__bytes_read.sum — redundant A traffic should vanish from dram__bytes_read; then compare smem-staged vs current on the same N/K.

Stage 3 — GDN recurrence (gated_delta_rule_decode_f32*) — ~5.4 ms/tok

What Atlas does. kernels/gb10/common/gated_delta_rule.cu:233-346. Grid (num_v_heads, batch) = (32, 1), block 128 threads (if (tid >= v_dim) return; at :251-252). Per thread: two full passes over H[j*v_dim + tid] for j = 0..127 — pass 1 reads 128 fp32 to form hk_dot (:277-284), pass 2 re-reads the same 128, updates, writes back, and forms q_dot (:292-318). State is fp32, 128×128×4 = 64 KB per v-head, so per layer per token: 32 heads × 64 KB read + 64 KB write = 4 MB, and pass 1 re-reads what pass 2 reads again → ~6 MB of traffic where 4 MB is the floor.

Three concrete inefficiencies, all visible in the source:

1. The state is never held in registers. 128 fp32 per thread = 128 registers — that is why the code re-reads. The tree already contains the fix in prototype form: kernels/gb10/common/gated_delta_rule_regresident.cu:46 (__launch_bounds__(32 * GDN_RR_WARPS_PER_BLOCK, 4)) and gdn_f16_state.cuh, plus gated_delta_rule_f16_state support (ssm_h_dtype.cu:25 ssm_h_state_f32_to_f16).

2. fp32 state. Halving to fp16 state halves the dominant traffic term. gdn_f16_state.cuh and the _f16 WY variants (gated_delta_rule_wy3_resident_f16.cu, gated_delta_rule_wy4_f16.cu) exist but are prefill/resident-path.

3. The clamp pass can still re-read. The comment at :308-312 records that the Frobenius accumulation used to be "a THIRD full pass over 3 MB per seq per layer" and was folded into the update loop — good — but the actual rescale when the clamp fires is still a separate full re-read+rewrite (:335-340). At the clamp threshold this doubles the layer's traffic.

What llama.cpp does (reference, not line-verified). The Flash-Next GDN decode kernel in llama.cpp keeps the recurrent state in fp16/bf16, uses one thread per (v-column, k-slice) with the k-loop unrolled into registers, and — most importantly — fuses conv1d-update + L2-norm + recurrence + gated-RMSnorm into a single launch per layer. Atlas has the same fusion idea (gated_delta_rule.cu:523 "FUSED conv1d_update_l2norm + recurrence + gated-RMS-norm decode kernel", gated_delta_rule_decode_f32_conv_norm at :537, and the gdn_verify_fused_* harnesses at gdn_verify_fused_k2.cu:7 and gdn_verify_fused_conv_kn.cu:7, whose comments complain that "causal_conv1d_update_l2norm is launched 17 TIMES" / "launched TWICE and gated_rms_norm is…"). The fused kernel exists; the profile says the unfused sequence is what runs — causal_conv1d_update_l2norm appears 36×/token as its own launch.

Specific change. (a) Switch the decode state to fp16 via the existing ssm_h_state_f32_to_f16 (ssm_h_dtype.cu:25) and an fp16 decode variant; (b) make the decode dispatch take gated_delta_rule_decode_f32_conv_norm (gated_delta_rule.cu:537) instead of conv + recurrence + norm; (c) port the register-resident k-loop from gated_delta_rule_regresident.cu.

Arithmetic. Floor with fp32 state = 4 MB/layer × 36 = 144 MB → 0.63 ms. Atlas: 5.4 ms. With fp16 state: 72 MB → 0.31 ms. Reaching a realistic 150 GB/s on this access pattern (64 KB strided slabs, poor coalescing across the j*v_dim+tid stride) gives 144/150 = 0.96 ms fp32, 0.48 ms fp16 → saving ~4.4–4.9 ms/token. Plus ~72 launches removed (conv + norm folded in) ≈ 0.3 ms. Total ~4.7 ms/token.

Risk. High on numerics — the file is explicit that summation order is preserved for bit-identicality (:310-312), and fp16 state at 15K+ tokens is exactly what the fp32-output variant was created to fix (:229-232: "Eliminates the BF16 truncation bottleneck in the recurrent path, preventing cumulative precision drift at 15K+ tokens"). Do not ship fp16 state without a long-context perplexity gate. The fusion and register-residency are safe; the dtype change is not.

Verify. ncu -k "regex:gated_delta_rule" --metrics dram__bytes_read.sum,dram__bytes_write.sum,gpu__time_duration.avg — expect dram__bytes_{read,write} to drop from ~6 MB/~2 MB per launch toward 2 MB/2 MB (fp32 fused) or 1 MB/1 MB (fp16). Long-context gate: 16K-token needle/perplexity diff vs the fp32 baseline.

Stage 4 — The kernel tax: mHC, QSA, conv, norms (~240+240+36+100 launches) — ~3.5 ms/tok

What Atlas does. Per layer the mHC highway issues five separate kernels — kernels/gb10/common/qwen4exp_hc.cu:44 q4e_hc_expand, :65 q4e_hc_lowrank_act, :84 q4e_hc_stream_mix, :113 q4e_hc_injection, :131 q4e_hc_scatter_add — and there are two mHC sites per layer (crates/spark-model/src/weight_loader/qwen4_exp.rs:328-335, hc::load_layer_sites + attach_hc), plus the QSA indexer issues ≥3 per full-attention layer (kernels/gb10/qwen3.8-flash-next/nvfp4/qsa_indexer.cu:106 qsa_qprep, :148 qsa_score, :178 qsa_gather, :431 qsa_select_topk). Each of these moves 10–26 KB. At GB10 a kernel launch is ~2–3 µs of gap plus ~1–2 µs of work: these kernels are 100% latency, 0% bandwidth.

What llama.cpp does. It does not have mHC, so there is no direct reference. The applicable reference is llama.cpp's general decode-graph discipline: minimum launches per layer, epilogues fused into the producer. Atlas already has the fusion vocabulary — moe_shared_expert_fused*, w4a16_gemv_dual (w4a16_gemv_fused.cu:144), w8a16_gemv_dual (w8a16_gemv_fused.cu:123), fused_k_norm_rope_cache_write_bf16 (fused_k_norm_rope_cache.cu:53).

Specific change. Fuse q4e_hc_expand + q4e_hc_lowrank_act + q4e_hc_stream_mix into one CTA (they are all elementwise/low-rank over the same 4×2560 = 10240-element residual; rank 320 per qwen4_exp.rs:56), and fuse qsa_qprep + qsa_score similarly. Target: 5 mHC launches → 2, 4 QSA launches → 2.

Arithmetic. ~240 mHC launches × ~1.5 µs saved + ~120 QSA launches × ~1.5 µs ≈ 0.54 ms/token from launch gaps alone, plus removing ~360 kernel-boundary pipeline drains of ~1 µs each ≈ 0.36 ms. Realistic total ~0.9 ms/token. Lower than it looks because the findings already say sync/graph work was worth only a few percent — this is the kernel-count version of the same lever, and it is smaller than the MoE item by an order of magnitude.

Risk. Low-medium (mHC is validated against a golden in ops/hyper_connection_lowrank_tests.rs per qwen4_exp.rs:11-15; a fusion must keep hc_post bit-exact).

Verify. nsys stats --report cuda_gpu_kern_sum launch counts before/after; cudaLaunchKernel per token from the CUDA API summary should drop by ~350.

Stage 5 — LM head + argmax + embed round-trip — ~2.9 ms/tok

What Atlas does. lm_head is a single BF16 GEMV over 248320×2560 = 1.27 GB, measured ~2.6 ms → ~490 GB/s effective, which is above the 230 GB/s sustained figure I was given. I could not reconcile that: either the 230 GB/s floor in the findings is a conservative sustained-LPDDR figure and the lm_head burst is reading from a warm L2 (1.27 GB does not fit 126 MB L2, so this is unlikely), or the lm_head timing in the window includes overlap. Flagging as unconfirmed. The path is load_lm_head (qwen4_exp.rs:428-443, tied to embed_tokens), kernel dense_gemv_bf16_fp32out (dense_gemv_bf16.cu:120-186, FP32 output deliberately, to avoid the greedy-argmax tiebreak flip documented at :111-117), then argmax_bf16/argmax_fp32 (argmax_bf16.cu:26/:208), then embed_from_argmax (embed_from_argmax.cu:17).

Specific change. Fuse argmax into the lm_head epilogue via a per-block atomicMax on a packed (value,index) uint64 — one launch instead of two, and it removes the 248320-element logits round-trip through DRAM (~1 MB read). Also: the lm_head GEMV has the same redundant-A problem as Stage 2 and the same fix.

Arithmetic. 1.27 GB at 230 GB/s = 5.5 ms floor — so lm_head at 2.6 ms cannot be DRAM-bound and I do not have an explanation for it. I will not claim a saving I cannot derive. What I can claim: removing the argmax launch and its logits re-read is ~0.15 ms/token, and it is free.

Risk. Trivial. Verify. nsys launch count −1; ncu -k argmax shows the kernel gone.

3. Ordered implementation plan

Realistic cumulative: ~27–34 ms/token saved → ~40 ms/token → ~25 tok/s, i.e. parity with llama.cpp, with item 1 alone doing most of the work.

Sequencing note: do #1 first and re-profile before touching anything else. Every other estimate is made against a profile whose top entry will have just moved by 25 ms.

4. What would NOT help — do not re-try

1. Host-side synchronisation, PLE placement, CUDA-graph capture. Already tried and measured; worth a few percent in total. The findings say so. The gap is inside kernels.

2. Requantising the GDN projections to NVFP4. Measured in-tree at crates/spark-model/src/weight_loader/qwen4_exp.rs:218-231: NVFP4 2.207 tok/s vs BF16 2.188 tok/s — 1% for 6.04 GB of VRAM and 2.5× the KV budget loss. The loader keeps BF16 by default (:237) and the precision argument is independent: "ONLY the routed experts are quantized in this checkpoint… requantizing them was a lossy round trip we chose, on 36 of 48 layers" (:233-236). Fix the BF16 GEMV kernel (item 3), not the dtype.

3. FP8 pre-dequant of the QKVZ weights for the decode path. Already measured and rejected in-tree: crates/spark-model/src/layers/qwen3_ssm/init_fp8.rs:104-107 — "FP8 is ~50% slower (1900µs vs 1228µs) because weight matrix [12288, 2048] is bandwidth-dominated… the 2× larger FP8 weights cost more than the dequant saves." Same arithmetic applies at M=1, worse.

4. Tensor-core grouped GEMM for MoE decode at M=1. kernels/gb10/common/moe_w4a16_grouped_gemm.cu:26-27 is M_TILE=64/N_TILE=64 with mma.sync.m16n8k16 (:183); at M=1 decode it wastes 63/64 of the MMA issue and re-dequants FP4→BF16 into padded smem every K_STEP=16 tile (:116-150). It is a prefill kernel. Decode wants a GEMV with better latency hiding — item 1 — not a GEMM.

5. More attention-kernel variants. The 12 full-attention layers are ~0.7 ms/token total at C=1 (~1% of the step). The tree already has ~40 paged_decode_attn_* turbo variants; none of them can recover 25 ms.

6. Sinkhorn/mHC math changes. mHC is golden-validated to cos 0.999998 with hc_post bit-exact (qwen4_exp.rs:11-15) and costs ~1.4 ms. Fuse it (item 4); do not redesign it.

7. MTP. Already honoured and already measured slower than plain decode on this model (qwen4_exp.rs:28-32, pointing at MODEL.toml).

What llama.cpp does worse (so nobody over-corrects)

1. Single-stream residual / no mHC highway. The PR #27742 Flash-Next build has no equivalent of the 4-stream low-rank residual (qwen4_exp.rs:56-58); anything Atlas gains from stream-level fusion is unavailable to it, and its hyper-connection handling is where a future Atlas advantage lives, not a deficit.

2. PLE n-gram injection. Atlas serves a ~320M-row n-gram table read by row off NVMe with segmented 128-shard offsets (qwen4_exp.rs:16-20, ple_shard_layout) and fuses it (q4e_ple_gate/q4e_ple_conv_add, qwen4exp_ple.cu:39,96). llama.cpp has no PLE path at all; if the mechanism is worth its accuracy, Atlas is strictly ahead.

3. QSA sparse attention. Atlas prunes KV with a per-query indexer on the 12 full-attention layers, parity-gated at T=2200 and provably inert below indexer_budget so short contexts stay exact (qwen4_exp.rs:21-24). llama.cpp's Flash-Next decode reads full KV. At long context Atlas's ~0.7 ms/token should scale far better than llama.cpp's linear KV read — the 15 vs 24 tok/s gap is a short-context, launch-and-MoE-efficiency gap, not a scaling gap.

4. Tied-embedding lm_head numerics. Atlas deliberately keeps FP32 logits (dense_gemv_bf16.cu:111-117, the 0.125-margin tiebreak that self-reinforces into stop-word loops). llama.cpp's quantised lm_head is faster but is exposed to exactly that class of failure. Do not trade this away for speed.

5. TP / multi-device. Not a llama.cpp advantage either — Atlas declines it here (qwen4_exp.rs:97-101) for principled reasons (mHC stream sharding + single-device PLE arena), not because it is slower.

Open items I could not confirm

• lm_head at ~490 GB/s vs the stated 230 GB/s achievable floor. Unreconciled. Either the floor figure is a sustained-LPDDR number that lm_head bursts past, or the window's lm_head timing overlaps with other work. I claimed no saving on that basis.
• I did not open ggml/src/ggml-cuda/mmvq.cu, mmq.cu, or the llama.cpp GDN/SSM kernel this session. Every llama.cpp claim in §2 is stated at the level of the known reference design (rows-per-warp heuristics, 16 B loads, register dequant, hoisted block scales, fused epilogues) and should be line-verified before the item-1 implementation is signed off.
• The exact per-launch µs split of items 3–12 came from the profile summary as I read it, not from a re-run ncu. The ranking and the order-of-magnitude BW deficits are solid; the decimals are ±30%.
• Which MoE decode variant actually fires for this target. crates/spark-model/src/layers/moe/init.rs:99-184 resolves many handles via try_kernel; I confirmed the handle set but did not trace the runtime branch that selects ..._2x over ..._wide / moe_decode_atomic_c4. That must be confirmed with a nsys symbol match before editing — if the atomic-c4 path is already live, item 1's launch-count saving is already banked and only the latency-hiding part remains.