GB10 (DGX Spark): CUDA decode is at ~85-90% of measured memory bandwidth; decode-side micro-optimization is exhausted

## Summary

Measurement report, no code. After profiling CUDA decode on a DGX Spark
(GB10) for #772, I ran a series of decode-side optimizations that all came
out throughput-neutral, and the follow-up measurements show why: **plain
decode on GB10 sits at ~85–90% of the physically achievable memory
bandwidth**, so decode-side micro-optimization is essentially exhausted on
this hardware. Recording the numbers so nobody re-runs the same dead ends.

Machine: DGX Spark GB10 (128 GB unified, sm_121, CUDA 13.0, driver
580.173.02), `make cuda-spark`, model
`DeepSeek-V4-Flash-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-chat-v2-imatrix.gguf`
(81 GiB), main `84cc882`.

## Where a decode token goes

`DS4_METAL_GRAPH_TOKEN_PROFILE=1`, steady state:

```
encode=9.8-10.6 ms execute=45.0-45.7 ms read=0.04 ms total=55.0-56.0 ms
```

nsys confirms `execute` is genuine GPU work (dense Q8 matvecs, the
compressor pair projections, MoE vec kernels), not launch gaps.

## The floor

A trivial saturating-read kernel measures **231–234 GB/s** device DRAM read
bandwidth on this box (probe source below). Decode must read every dense
layer plus 6 routed experts × 61 layers per token — roughly 10–13 GB of
weight traffic — which puts the physical floor at **~45–55 ms/token**.
That is exactly the measured `execute` time: the CUDA decode path is
already at ~85–90% of what the memory system can deliver.

Cross-check: llama.cpp plain decode on the same box and a comparable quant
(UD-Q2_K_XL) measures 17.50 ± 0.10 t/s (`llama-bench` tg128) vs ds4's
17.9 — both engines sit on the same floor, which is what you'd expect if
the floor is the hardware.

## Decode-side experiments that measured flat (all bit-identity gated)

- CUDA port of the Metal M5 triple fusion (qkv norm + KV rope + FP8 +
raw store in one dispatch): bit-identical, −61 launches/token, flat.
- Extending decode island 0 through the KV store with a per-token device
descriptor (so the position-dependent kernels graph-capture): ~300
launches/token moved from eager to graph replay, bit-identical, flat.
- Shared-memory staging for the ordered-chunk pair matvec: bit-identical,
−3.6% (reverted).
- `DS4_CUDA_GREEDY_SPLIT_TOP1`, `DS4_CUDA_OUTPUT_FUSED_TOP1`,
`DS4_CUDA_GREEDY_VEC4`, `DS4_CUDA_EXACT_SCORE_SPLIT_GRAPH`,
`DS4_CUDA_END_STREAM_SYNC`: flat or slightly negative on this box.

## Implication

On GB10 the only decode-throughput lever left is reading fewer weight
bytes per emitted token, i.e. DSpark: the batched verify pays the dense
read once for up to `block` tokens. The measured cycle economics
(`DS4_DSPARK_STATS=1`): target eval ~60 ms, 2–5-row verify ~83–97 ms — so
every accepted draft is worth ~55 ms of avoided weight traffic, and
accept length/coverage is the whole game on this hardware. Consistent
with the QA table's 24.48 t/s direct-commit fixture number.

Happy to share the bench scripts, nsys traces, or run comparisons on this
box if useful.

bandwidth probe

```cuda
__global__ void readsum(const float4 *p, float *out, size_t n) {
size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
size_t stride = (size_t)gridDim.x * blockDim.x;
float acc = 0;
for (; i < n; i += stride) { float4 v = p[i]; acc += v.x + v.y + v.z + v.w; }
if (acc == 12345.678f) *out = acc;
}
// 8 GiB cudaMalloc buffer, <<<2048, 256>>>: 231.4 / 233.3 / 233.6 GB/s
```

Start with the measurement report, the GB10 configuration, and the inline bandwidth-probe source; use `make cuda-spark` and the listed model if reproducing the environment. Compare the decode timings, bandwidth figures, and `llama-bench` cross-check, then confirm that the documented conclusion and dead-end experiments are recorded without proposing further decode micro-optimizations.

Do mô hình lập chỉ mục viết ra từ nội dung của issue.