fmt=4 CUDA attention crash — FIXED (#298); benchmark tracker + follow-ups (PROFILE accounting, cache warmth)

[woolcoxm]

fmt=4 CUDA attention crash — fixed; benchmark tracker + follow-ups

Tracking the resolution of the g64-model hard-reboot crash and the benchmark status. The crash is fixed (PR #298, commit 22a224e). This issue captures the benchmark data, the CUDA fix summary, and the open follow-ups so people following along have the full picture.

Supersedes/closes the crash hypothesis in #326 (the "possible VRAM-crash correlation" there is now confirmed and fixed — see below).

---

The crash, and the fix

Symptom: the g64 (fmt=4 grouped int4, gs=64) model hard-crashed the machine (black screen → audio loop → reboot) when run with the full CUDA stack (COLI_CUDA_ATTN=1 COLI_CUDA_PIPE=2 COLI_CUDA_MTP=1 + REPIN=16). The old per-row int4 model never crashed.

Root cause: the two MLA attention absorb kernels in backend_cuda.cu (attention_absorb_kernel, attention_absorb_batch_kernel) applied quantization scales per output row (wscale[row]) — the fmt=2 semantic. PR #298 added fmt=4 support to quant_matmul (per-group scales: scales[o*ng + g]) but missed these two kernels. The g64 kv_b tensor is fmt=4 (ng=8 groups/row), so its scale buffer is O*ng long; indexing it with a bare row index read with the wrong stride → GPU memory fault → Windows bugcheck 0x116 VIDEO_TDR_FAILURE (nvlddmkm.sys, STATUS_INSUFFICIENT_RESOURCES).

Fix (PR #298 22a224e): added an absorb_scale() __device__ helper mirroring quant_matmul's fmt=4 branch (wscale[row*ng + k/gs] for fmt=4, wscale[row] otherwise), applied it inside the Q- and V-projection accumulation loops of both kernels, and threaded w->gs, w->ng through all six launch sites. For fmt≠4, ng==1 → bit-identical to prior behavior. No extern "C" signature or header changes.

Validation: isolated flag tests (base → COLI_CUDA_ATTN=1 short/long prompt → COLI_CUDA_PIPE=2 long-prompt prefill, the exact crash config) all pass; user confirmed the full benchmark no longer crashes. Full investigation writeup: #298 (comment)

---

Benchmark data (post-fix, full CUDA stack)

Config: COLI_CUDA=1 CUDA_DENSE=1 COLI_CUDA_ATTN=1 COLI_CUDA_PIPE=2 COLI_CUDA_MTP=1 DIRECT=1 PIPE=1 PIPE_WORKERS=8 REPIN=16 RAM_GB=28 PIN_GB=8 CUDA_EXPERT_GB=4, NGEN=24.
Machine: Core Ultra 9 185H / 32 GB / RTX 5070 Ti 16 GB (driver 610.74, CUDA 12.8) / NVMe 2.6 GB/s / Windows.

metric	value
decode	0.18 tok/s (24 tok / 135s)
prefill	22 tok / 88.86s
MTP acceptance	47% (14/30), 2.40 tok/forward
expert hit rate	7.9% (cold cache)
experts loaded/token	964.7 (per-layer 12.86, baseline topk=8)
peak RSS	20.27 GB
VRAM resident	751 tensors, 11.27 GB
VRAM expert tier	39 experts (0.83 GB), 280 calls served, 23 REPIN swaps

No crash. This is the config that rebooted the machine before the fix.

The 0.18 tok/s is cold-cache honest (the PR's earlier 1.08 tok/s used EXPERT_BUDGET=4, since quarantined as garbage-producing). The dominant cost is expert disk I/O (7.9% hit rate → ~965 expert loads/token), not the CUDA path. Next lever is cache warmth, not GPU.

Full data: #298 (comment)

---

Open follow-ups

• PROFILE accounting regression. After rebasing cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness #298 onto dev, profile_print shows negative "other" times (other -258s prefill, -730s decode). Cause: dev's edisk_s() is a cumulative atomic counter, but profile_print's accounted sum treats it as a per-phase delta → accounted > elapsed → negative other. prof_report (serve mode) already does this correctly with a per-window delta (edisk_s() - b->edisk); profile_print needs the same treatment. Not a correctness issue (tok/s, hit rate, RSS are all valid) — only the PROFILE phase breakdown is affected. Small, self-contained fix.
• Cache warmth sweep. With more .coli_usage history (warmer pin set), hit rate should climb and expert-load count should drop. Worth a run after a few interactive turns warm the cache, to get a non-cold tok/s baseline.
• (Carried from New int4-g64 container is ~12% more resource-hungry than per-row int4; possible VRAM-crash correlation (unconfirmed) #326) int2 for cold experts — halves the ~965-expert-load/token disk I/O while keeping the hot pin set at int4. The engine has the matmul_i2 kernel + converter plumbing; needs a converter fix, exactness test, and quality measurement.

---

Prevention note (for the next quant switch)

When adding a new quant format or editing a GPU kernel: any kernel that multiplies a quantized weight by a scale must index the scale per-group for fmt=4 (scales[row*ng + k/gs]), not per-row. Use the absorb_scale() helper. Audit every wscale/scales access site, not just the matmul. The grouped MoE kernels (grouped_hidden*/grouped_down*) are still per-row-only and protected only by the hard reject at coli_cuda_expert_group line 623 — if that's ever lifted to support fmt=4 experts, those kernels (and GroupDesc, which lacks gs/ng) must be updated too. Full checklist in the PR comment linked above.

[woolcoxm]

PROFILE accounting regression — fixed

Resolved in PR #298 commit f192263. The accounted sum was using edisk_s() (cumulative cross-thread disk service) instead of m->t_ewait (felt compute-thread wait), making other go negative. Verified: PROFILE now sums to wall (decode wait 9.628s + matmul 1.505s + attn 0.529s + other 0.278s ≈ 11.9s). See #298 (comment).

[mohamedmastouri2000-boop]

Blackwell datapoint for the tracker � RTX 5080 (sm_120) / Core Ultra 9 285K / 128 GB / Windows 11 native, on the public g64 container (mastouri/GLM-5.2-colibri-int4-g64-with-int8-mtp), engine built from the fmt=4 branch.

#326 crash config � the one that used to VIDEO_TDR_FAILURE on Blackwell: COLI_CUDA=1 COLI_GPU=0 CUDA_DENSE=1 CUDA_EXPERT_GB=4

• No TDR, no crash � coherent output
• 0.66 tok/s cold decode, expert hit 81.6% (pin 16.5% + LRU 62.2%)
• 109 experts resident in VRAM (~2.3 GB budget), 470 calls served from VRAM, dense resident
• 15.7 GB VRAM in use

Ragged-attention correctness gates (both merged):

• cuda: batch ragged attention across independent streams #365 test_ragged_attention: ragged_relative_rms=0 (bit-identical to per-sequence reference)
• cuda: keep paged ragged KV resident across decode steps #372 resident paged KV: rms=0 + a real-weights multi-step decode A/B where the GPU resident-KV path and the CPU reference greedy-decoded the identical 16-token sequence (no drift)

So on sm_120 the full fmt=4 + ragged/resident-KV path is stable and correct end-to-end. Happy to add MTP-on numbers or a longer decode profile if useful.

[VincentMarquez]

DGX Spark / full-k8 speed campaign pointer

Posted a progress update on the Spark datapoint thread (#161): comment.

Headline (timed decode, full top-8, GB10 121 GB UMA): S=1 fused ~7 tok/s; best multi-S + free PLD e2e ~8.5–9.5 (mid peak ~9.5). Goal remains 30 tok/s full-k8 — not there yet; multi-S verify cost vs S=1 is still the wall. TOPK=1 ceiling ~14 is logged separately and not mixed into full-k8 rows.

No code push from that campaign into this issue — status/benchmarks only.

[JustVugg]

Progress on the follow-ups: #451 (grouped-int4 fmt=4 support in the CUDA expert-group kernels) landed in dev, opening the GPU tier to g64/E8 containers — which is the other half of making fmt=4 first-class on CUDA. #298 (the full fmt=4 engine+attention support this issue tracks) is still open and needs a rebase onto the post-#391 tree. Keeping this tracker open until #298 lands and the PROFILE-accounting / cache-warmth follow-ups are closed out.

[JustVugg]

#298 is merged into dev (2026-07-20). The dense/attention per-group scale fix landed with hardware verification: on RTX 5080/sm_120 the exact CUDA_DENSE=1 + g64 config that produced garbage now runs coherent — 952 dense tensors (12.48 GB) + 109 experts fully VRAM-resident, zero fallbacks (thanks @woolcoxm for the rebase and @mohamedmastouri2000-boop for the isolation + verification). This tracker's remaining items (PROFILE accounting, cache warmth) are now unblocked on a correct baseline — please re-run the fmt=4 benchmarks on current dev so the tracker numbers reflect the fixed path.

[mohamedmastouri2000-boop]

Heads-up for the tracker's Windows "CPU" rows: before #468 they weren't CPU — re-collect on current dev if precision matters.

While building the resource matrix in #467 I found that on Windows, COLI_CUDA=0 was silently overridden by the wrapper's bare-run auto-enable: a "CPU" run on a CUDA build with a GPU present actually got a ~12.6 GB VRAM expert tier. Measured effect on my 5080 box: 0.72 tok/s ("CPU" with the hidden expert tier) vs 0.62 true-CPU — a ~16% phantom speedup, with the confound invisible unless you notice the [GPU] auto-enabled line.

As of #468 (merged into dev today), COLI_CUDA=0 is honoured and clears stale device/sizing vars; --gpu none remains the canonical hard off-switch and always worked. So:

• Affected: Windows rows labelled CPU that were collected with a bare run or COLI_CUDA=0 on a GPU-equipped box, on any build before fix(coli): honour explicit COLI_CUDA=0 over the Windows bare-run auto-enable (+ drive-dependent DIRECT=1 docs) #468.
• Not affected: Linux/macOS rows, rows collected with --gpu none, and boxes with no CUDA build/GPU.

One more caveat when comparing CPU vs GPU rows anywhere in the tracker: CPU mode auto-enables MTP speculation (draft=3; worth ~1.5–2× on accepting prompts) while CUDA disables it by default — measured details in #467. Worth stating MTP state next to any CPU number.