You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
Follow-up to the #428 desktop validation. The GPU draft round is the dominant DSpark cost — 54% of the round, ~27 ms on Qwen3-8B + dspark_qwen3_8b_block7 (RTX 4070 Ti), profiled as enqueue 67 / gpu-wait 1331 / heads 0 ms (totals over a 128-token run, ~50 rounds → ~27 ms/round, ~95% GPU-exec-bound).
Observation
The draft is already well-built: all projections (fc, q/k/v/o, gate/up/down, lm_head) use batched cublasGemmEx (CudaBackend.MatMulBatchedGemmF16W, weight read once for all B=7 positions), and argmax + the Markov chain run on-device — the [B×vocab] logits D2H is eliminated (hence heads=0 ms). So the laptop analysis's two named levers (sequential lm_head re-streaming; host Markov re-streaming) are already solved on GPU.
But 27 ms/round is ~3× the raw bandwidth floor: reading the ~4.7 GB fp16 head once at the 4070 Ti's ~500 GB/s is ~9 ms. The ~18 ms over floor is unexplained by attention (short ctx in the bench → tiny KV) or launches (enqueue only ~1.3 ms/round).
Hypothesis
The projections are M = B = 7 tall-skinny GEMMs, and cuBLAS doesn't saturate HBM bandwidth at M=7 (tile quantization / low occupancy for the small-M dimension). The block-forward loop is CudaDSparkDraftModel.cs:378-424 — 7 MatMulBatchedGemmF16W(..., B) calls per layer (q/k/v/o + gate/up/down) × NumLayers, plus the lm_head GEMM at :424.
Proposed work
Profile first (prerequisite — the 3× is inferred). Nsight per-kernel pass on the draft round to pin the over-floor cost: is it the small-M GEMMs, AttentionBatchedRagged (:409), the per-layer/per-markov View/Free churn (:388-402, :437-447), or the 7× sequential re-read of the 78 MB markov_w2 in the Markov chain (:433-448)? Confirm before committing to a kernel swap.
This targets the dominant draft cost, so it's the largest single DSpark lever. But per #428, it won't flip the verdict on the desktop 8B: plain decode there is 76.5 t/s (2× the laptop's 38.3, bandwidth-bound), so DSpark stays content-gated. This lever + #430 (verify graph-capture, lever 2) widen the win region — pushing the best case from ~1.05× toward ~1.2–1.3× on high-acceptance / code-completion content — rather than making DSpark beat plain on general prose. Prioritize accordingly (real ROI only for high-acceptance-content use cases).
Follow-up to the #428 desktop validation. The GPU draft round is the dominant DSpark cost — 54% of the round, ~27 ms on Qwen3-8B +
dspark_qwen3_8b_block7(RTX 4070 Ti), profiled asenqueue 67 / gpu-wait 1331 / heads 0 ms(totals over a 128-token run, ~50 rounds → ~27 ms/round, ~95% GPU-exec-bound).Observation
The draft is already well-built: all projections (fc, q/k/v/o, gate/up/down, lm_head) use batched
cublasGemmEx(CudaBackend.MatMulBatchedGemmF16W, weight read once for all B=7 positions), and argmax + the Markov chain run on-device — the[B×vocab]logits D2H is eliminated (henceheads=0 ms). So the laptop analysis's two named levers (sequential lm_head re-streaming; host Markov re-streaming) are already solved on GPU.But 27 ms/round is ~3× the raw bandwidth floor: reading the ~4.7 GB fp16 head once at the 4070 Ti's ~500 GB/s is ~9 ms. The ~18 ms over floor is unexplained by attention (short ctx in the bench → tiny KV) or launches (
enqueueonly ~1.3 ms/round).Hypothesis
The projections are M = B = 7 tall-skinny GEMMs, and cuBLAS doesn't saturate HBM bandwidth at M=7 (tile quantization / low occupancy for the small-M dimension). The block-forward loop is
CudaDSparkDraftModel.cs:378-424— 7MatMulBatchedGemmF16W(..., B)calls per layer (q/k/v/o + gate/up/down) ×NumLayers, plus the lm_head GEMM at :424.Proposed work
AttentionBatchedRagged(:409), the per-layer/per-markovView/Freechurn (:388-402, :437-447), or the 7× sequential re-read of the 78 MBmarkov_w2in the Markov chain (:433-448)? Confirm before committing to a kernel swap.CudaWsKernels), which amortizes weight HBM reads across the batch and saturates bandwidth at small M. Caveat: perf(cuda): weight-stationary batched-decode matmul — amortize weight HBM reads across the batch (#190 follow-up) #194's kernels target quantized weights; the draft weights are resident fp16, so this likely needs an fp16 WS-matvec variant (or an SoA/half2 repack). Argmax-stable is fine (DSpark is greedy, and the draft only proposes — the target verify is exact).Impact & caveat
This targets the dominant draft cost, so it's the largest single DSpark lever. But per #428, it won't flip the verdict on the desktop 8B: plain decode there is 76.5 t/s (2× the laptop's 38.3, bandwidth-bound), so DSpark stays content-gated. This lever + #430 (verify graph-capture, lever 2) widen the win region — pushing the best case from ~1.05× toward ~1.2–1.3× on high-acceptance / code-completion content — rather than making DSpark beat plain on general prose. Prioritize accordingly (real ROI only for high-acceptance-content use cases).
Refs: #428 (validation + profile), #430 (lever 2 — verify graph-capture), #194 (WS-matvec infra, closed), #427 (load-aware verify length).
🤖 Generated with Claude Code