Gb10 cuda graph fix#96
Open
sumergoconicio wants to merge 13 commits into
Open
Conversation
sumergoconicio
commented
Jul 20, 2026
sumergoconicio
commented
Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>
…ocked) Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>
co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>
…D, perf regression exposed Apply Opus 4.8 review of /tmp/diagnostic-brief.md (5 specific bugs identified). Clean rewrite of mmq-blackwell-q1.cu resolves: 1. misaligned int32 read on sB 2. per-q-block scale index (kc IS the q1-block index when bK=128) 3. warp column coverage (16 rows x 64 cols sub-tile, fully covered) 4. int32 truncation risk replaced with fp32 accumulator 5. dead first-half loop removed Verification (Phase 3.3): kernel runs end-to-end on real M=512 prefill. 86/86 q1_0/q2_0 unit tests pass with env=off or env=on. compute-sanitizer --tool memcheck: no Invalid __shared__ reads. Measurement (Phase 3.4): cold-cache pp512 tg128 on Bonsai-27B-Q1_0.gguf: Blackwell OFF (cuBLAS): pp512 925 +/- 45 tg32 40 +/- 0.2 Blackwell ON (int8 MMA): pp512 209 +/- 22 tg32 39 +/- 0.6 Decode (M=1) confirmed unchanged because dispatch gating requires M >= 128 to enter Blackwell path. Prefill regresses ~4.4x. Decision: route the perf gap to a follow-up Opus review (Phase 3.5) covering correctness cross-check and missing-pipelining hypotheses (cp.async, register spill, scale-in-loop, software pipelining, grid starvation, non-coalesced global reads). Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>
Opus 4.8 response to /tmp/diagnostic-brief-perf.md saved at
/tmp/diagnostic-brief-perf-response.md. Read-only; src untouched.
Top line: kernel is staging-bound on prologue, not tensor-core-bound.
Three ranked fixes:
Rank 1 (H6): Stage A coalesce + single int load (also fixes
sign-extension-or latent correctness bug for negative activations),
kill in-loop weight decode (4x mblk redundancy). Predicted 209
-> 300-380 from Stage A alone, 450-600 with pre-expansion.
Rank 2 (H2): swap kk-outer/f-inner to f-outer/kk-inner so only
one acc_i[4] is live per warp (halves regs). Predicted 550-700.
Rank 3 (H1): cp.async + double-buffered sA/sB SMEM. Predicted
800-900 (last mile to cuBLAS parity, not first thing to fix).
Opus revised pass criterion: two-step gate.
Step 1: >= 550 tok/s after minimum patch (PrismML-Eng#1+PrismML-Eng#2+PrismML-Eng#3).
Step 2: >= 850 tok/s after full Hopper-style pipeline.
Original single-number 856 criterion is not realistic from a hand
kernel without cp.async; revise as above.
Recorded decision: apply minimum patch as next concrete experiment
(Phase 3.6). Hold the .cu source commits until measured lift
confirms the kernel is past the 550 ground.
Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>
…ismML-Eng#3 no-op) Recorded measured performance from the two Opus minimum-patch edits: Patch PrismML-Eng#1 (Stage A coalesce + single int load): Cold-cache pp512 -> 470 +/- 18 tok/s (was 209 +/- 22; +2.25x) Above Opus's predicted 300-380 ceiling for that step. Patch PrismML-Eng#3 (Stage D frag-outer / kk-inner): Cold-cache pp512 -> 466 +/- 23 tok/s No measurable additional lift. Confirms Opus stacking order: H2/H1 lifts were predicated on H6 (pre-expansion) being in place. Without it, register-pressure relief just buys more warps blocked on Stage B global reads. Honest framing of current state in the ledger: ~470 tok/s == ~52% of the cuBLAS reference (899 tok/s) with the kernel running cleanly, decode (M=1) untouched, 86/86 unit tests, sanitizer clean, and 3.5 GB VRAM-resident throughout. Patch PrismML-Eng#2 (pre-expand to int8 in global VRAM) is rejected on the principle that 8x VRAM (27 GB resident) defeats the Bonsai 1-bit intelligence-density thesis. Open path: 4x mblk redundancy elimination in Stage B without VRAM blow-up -- separate Opus brief in flight. Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>
…3.5 GB VRAM
Lands the Blackwell int8 tensor-core MMQ path for Q1_0/Q2_0 weights.
Opt-in via GGML_BLACKWELL_Q1; cc-gated to DGX_SPARK (sm_121a); falls
through to MMQ on shape mismatch. Preserves the 3.5 GB on-disk /
3.5 GB VRAM-resident footprint of Bonsai 27B Q1_0 (no global Q1
pre-expansion that would blow up VRAM to 27 GB).
Measured cold-cache performance on Bonsai-27B-Q1_0.gguf:
Blackwell OFF (cuBLAS): pp512 ~870 tg32 ~39 tok/s
Blackwell ON bM=256: pp512 ~692 tg32 ~40 tok/s
(79.5% of cuBLAS)
Blackwell ON bM=128: pp512 ~560 tg32 ~39 tok/s (intermediate)
Driven by Opus feedback across two review rounds:
Step 3 segfault root-causing and clean rewrite of mmq-blackwell-q1.cu
(lane-based tile<> loaders, packed-int SMEM, float accumulators,
per-q-block scale index, dead-loop deletion).
Step 3.7 Stage B Q1 (blocked->interleaved coalesce) +
Stage B Q2 (per-b8 read collapse). Predict 530-600; observed ~560.
Step 3.8 bM=256 falsifier. Opus predicted 'DRAM-bound if pp512 ~600+';
got avg ~692 -> confirmed DRAM-bound -> rationalizes further
bM=512 + cp.async work to push past 700-800.
Hard constraints honored:
VRAM-resident at 3.5 GB packed; no pre-expansion.
Decode (M=1, tg32) unchanged at ~40 tok/s.
86/86 q1_0/q2_0 unit tests pass with env on AND env off.
compute-sanitizer --tool memcheck clean (after fixing warp-mapping
math on the bM=256 first cut: (f & 7) * 8 not f * 8 for col index).
All edits insertion-only with // GB10: sentinel markers
(mmvq.cu Phase 2.4 + ggml-cuda.cu Phase 3 dispatch + new
mmq-blackwell-q1.cu kernel).
Open path (Phase 3.9, NOT in this commit):
Opus: bM=512 + 1024 threads + cp.async double-buffer for 700-800.
Multi-rewrite, requires per-warp mapping under 1024 threads, M=256
routing, etc. Ship the current state at 692 tok/s (79.5% cuBLAS)
as the Bonsai 1-bit thesis-preserving Blackwell Blackwell path.
Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>
Added explicit Phase 3.9 entry describing the W-stationary path (bM=512 + 1024-block-threads + cp.async double-buffer) that Opus flagged as the route to 700-800 tok/s, plus the multi-day-rewrite caveats. Phase 4 section now populated with the actual shipped state on gb10-blackwell branch head (cbe9035) — including PR-readiness checklist, risk profile, and the optional Phase 3.9 follow-up. Summary of state meshed with measured numbers: 692 tok/s pp512 cold-cache, 79.5% cuBLAS, 3.5 GB VRAM-resident, 40 tok/s decode unchanged, 86/86 unit tests, sanitizer clean. Marks Phases 0+1+2.1+2.4+3.0-3.8 as SHIP-READY at this checkpoint, with Phase 3.9 deferred to a separate commit if/when pursued. Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>
…k/s gate Opus's cp.async + KG=2 minimum-patch prediction was 760-820 tok/s. Measured: 550 tok/s avg (574, 526 across two runs), a 20% regression versus the 692 tok/s bM=256 single-buffer baseline. Control experiment (KG=2 SMEM doubling without cp.async) gave 591 tok/s avg (592, 591 across two runs), confirming the regression source is the 2x SMEM footprint which cuts occupancy 2 -> 1 block/SM on Blackwell sm_121a. cp.async cannot hide SMEM-stall on Blackwell's sync mma.sync; Hopper's wgmma-async overlap pattern that motivated Opus's prediction does not apply on sm_121a. Per Opus's pp512 < 750 decision rule, documented the failure rather than shipping the regression. Reverted the .cu source to the ship point at 692 tok/s (bM=256 + single-buffer SMEM) head cbe9035. Routes forward that could push past 692: - bN=128 within current SMEM budget (Opus-predicted 830+). - Full wgmma-style async-MMA rewrite (multi-week, requires PTX tcgen05 intrinsics on sm_121a, NOT a 1-day patch). Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>
This commit closes the experiment. The Blackwell int8 tensor-core MMQ kernel for Q1_0/Q2_0 weights ships at 692 tok/s cold-cache pp512 (79.5% of cuBLAS reference 871), env-gated behind GGML_BLACKWELL_Q1 defaulting to off. Decode (M=1, tg32) unchanged at 40 tok/s. Memory footprint exactly Bonsai's 3.5 GB packed throughout. 86/86 unit tests pass with env on AND env off. compute-sanitizer clean. Two negative results are part of the shipped artifact, documented so future engineers don't re-investigate the same dead ends: - Patch PrismML-Eng#2 (pre-expand weights to int8 in global VRAM) rejected on the principle that 8x VRAM (27 GB resident) defeats the Bonsai 1-bit intelligence-density thesis. - Phase 3.9 cp.async + KG=2 attempt failed the 750-tok/s gate (550 tok/s measured, 20% regression vs the 692 baseline), isolated to SMEM-doubling occupancy hit on sm_121a sync mma. Branch head: gb10-blackwell @ e45ce1b. Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>
Refresh ledger references after Phase 4 SHIP conclusion commit: - Phase 4 'Sync / upstream' HEAD: cbe9035 -> 102e97b - Phase 4 'Ship conclusion' header: e45ce1b -> 102e97b - Commit list: add 102e97b, reorder so source ship (cbe9035) is distinguished from ledger commits (3d22967, e45ce1b, 102e97b). - PR readiness table status: 'ship-ready locally' -> 'shipped env-gated'. Source code is unchanged from cbe9035 (Phase 3 shippable commit). Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>
Table was cc-gated (not env-gated) and perf-neutral on sm_120 bench; nwarps 4->8 entries speculative. Revert DGX Spark launch config to GENERIC. Vec_dot coalesce variants and env-gated mmq-blackwell-q1.cu stay. Move GB10-PATCHES.md to docs/, strip /tmp/*.md paths; add dispatcher comment for M%256 fallthrough. Assisted-by: Claude (Droid)
Graph warmup requires two consecutive calls with identical properties to complete, but decode graphs change uid on every step which triggers a spurious 'properties changed' reset. Add a structural comparison that recognizes uid-only changes when the node count, ops, types, and shapes (of non-noop nodes) are identical. Data pointers are excluded from the comparison since cudaStreamCaptureModeRelaxed handles runtime pointer patching. Noop nodes (VIEW/PERMUTE/RESHAPE/TRANSPOSE) are skipped since their shapes vary with sequence length. This fixes the warmup-never-completes loop on hybrid-architecture models with gated-delta-net/SSM components. Assisted-by: Claude Sonnet
github-actions
Bot
added
documentation
There was a problem hiding this comment.
Pull request overview
Combines a CUDA graph warmup fix with unrelated DGX Spark Q1_0/Q2_0 CUDA optimizations.
Changes:
- Adds structural CUDA graph reuse detection.
- Adds Blackwell MMVQ and experimental MMQ kernels.
- Documents GB10 benchmarking and development history.
Reviewed changes
Copilot reviewed 5 out of 5 changed files in this pull request and generated 16 comments.
Show a summary per file
| File | Description |
|---|---|
ggml-cuda.cu |
Graph reuse and Blackwell dispatch |
mmq-blackwell-q1.cu |
New low-bit MMA kernel |
mmvq.cu |
Blackwell vector-dot routing |
vecdotq.cuh |
Aligned Q1/Q2 load variants |
GB10-PATCHES.md |
GB10 experiment ledger |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Comment on lines
+2541
to
+2542
| // GB10: Blackwell (sm_121 / DGX Spark) int8 tensor-core MMQ path for Q1_0/Q2_0. Opt-in (env GGML_BLACKWELL_Q1). | ||
| bool ggml_cuda_mul_mat_q1_blackwell(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst); |
Comment on lines
+3356
to
+3358
| if (structural_match) { | ||
| graph->uid = cgraph->uid; | ||
| return false; |
Comment on lines
+3334
to
+3336
| if (n->op != p->op || n->type != p->type || | ||
| memcmp(n->ne, p->ne, sizeof(n->ne)) != 0 || | ||
| memcmp(n->nb, p->nb, sizeof(n->nb)) != 0) { |
Comment on lines
+3345
to
+3349
| if (n->src[j]) { | ||
| if (memcmp(n->src[j]->ne, graph->node_props[i].node_src_ne[j], | ||
| sizeof(n->src[j]->ne)) != 0 || | ||
| memcmp(n->src[j]->nb, graph->node_props[i].node_src_nb[j], | ||
| sizeof(n->src[j]->nb)) != 0) { |
Comment on lines
+3328
to
+3333
| const ggml_tensor * n = cgraph->nodes[i]; | ||
| if (ggml_is_empty(n) || n->op == GGML_OP_RESHAPE || n->op == GGML_OP_TRANSPOSE || | ||
| n->op == GGML_OP_VIEW || n->op == GGML_OP_PERMUTE || n->op == GGML_OP_NONE) { | ||
| continue; | ||
| } | ||
| const ggml_tensor * p = &graph->node_props[i].node; |
Comment on lines
+6
to
+8
| // Block tile: bM=128, bN=64, bK=128 (= one block_q1_0 / block_q2_0 quant block). | ||
| // 8 warps (256 threads) per block; each warp owns 16 rows × 64 cols (= 8 m16n8 frags). | ||
| // 8 warps × 8 frags = 64 frags = bM*bN / (16*8) per block. |
| @@ -0,0 +1,487 @@ | |||
| // mmq-blackwell-q1.cu (Phase 3 — Blackwell int8 tensor-core MMQ for Q1_0/Q2_0, fixed per Opus review) | |||
| Build: GB10 (cc=1210, compute_cap=12.1, sm_121a), CUDA + CUTLASS enabled | ||
| Remotes: `prismml` / `upstream` / `github` (all SSH) | ||
|
|
||
| ## Phase 0 — Repo scaffold (DONE) |
Comment on lines
+499
to
+502
| expanding VRAM past the 3.5 GB packed footprint. Decode (M=1, | ||
| tg32) unchanged at 40 tok/s. 86/86 q1_0/q2_0 unit tests pass | ||
| with env on AND env off; `compute-sanitizer --tool memcheck` | ||
| clean. Additive-only changes with `// GB10:` sentinels. |
| - Decode (M=1, tg32) unchanged within noise (37.7–39.0). | ||
| - 86 / 86 q1_0/q2_0 unit tests pass on env off OR env on. | ||
| - `compute-sanitizer --tool memcheck` clean. | ||
| - **VRAM footprint unchanged: 3.5 GB packed weights throughout.** |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.