Conversation
|
Cross-reference: #215, opened the same day, attacks the same batch-1 wall from the same diagnosis: the AoS q8_1 activation loads cost about 36 L1 wavefronts per warp instruction in the small-K geometry, and 88 of 128 threads idle at K=5120. The two activation layouts differ (#215 groups 32 K blocks with word w contiguous, this PR stores each column as 8 planes of 16-byte pieces plus a plane of scales), both change What this PR has that #215 does not:
What #215 has that this PR does not: exact int16 activation sums with the bias subtracted once per block (one fewer op per word in the trit decode, worth adopting here too) and the GDN gather fusion. Proposed split, if it helps review: #215's exact sums and GDN fusion, this PR's multi-column path, invariance switch and tests, rebased onto whichever layout the maintainers prefer. Extending #215's word-grouped layout to N columns is a per-column stride, so that rebase is mechanical; happy to do it on request. RTX 3060 numbers for the #215 branch will follow once the card is free (#215 was measured on a 4070). |
bri-prism
left a comment
bri-prism
left a comment
There was a problem hiding this comment.
Agent review: posted by the maintainer's coding agent at their request.
No additional correctness finding in this source pass. Please reconcile the activation layout and MMVQ changes with #215 before integration; these are competing implementations rather than independent patches. Keeping the gather fusion separate would make that comparison easier.
The description correctly records whole-model batch invariance only for 1-4 tokens. Please keep the switch's documentation consistent with that limitation rather than promising 1-8. CUDA compilation, numerical tests and throughput were not rerun in this review.
Reviewed commit: 588318660c5076795ad43487c3410f1a44970bf8.
bri-prism
left a comment
bri-prism
left a comment
There was a problem hiding this comment.
Agent review follow-up, posted at the maintainer’s request: reproduced a build blocker at head 5883186.
[P1] Keep architecture-dependent restrict qualifiers off the new kernel signature — ggml/src/ggml-cuda/mmvq-ptq1_0.cuh:275–277.
On RTX 5090, CUDA 12.8.93 with GNU host compiler 13.3.0, building llama-bench and test-backend-ops for CMAKE_CUDA_ARCHITECTURES=120 (normalized by CMake to 120a) succeeds at base 9a9394a and fails at this PR head. Both quantize.cu and mmvq.cu fail while compiling the generated host stubs for mul_mat_vec_ptq1_0_pt:
error: template-id '__wrapper__device_stub_mul_mat_vec_ptq1_0_pt<1, 4, true, true>' ... does not match any template declaration
The generated specialization takes const void*& / float*&, while its candidate declaration takes const void* __restrict__& / float* __restrict__&. The new formal parameters use GGML_CUDA_RESTRICT; that macro expands differently in the host pass and the Hopper-or-newer device pass when PDL is enabled. This makes the new kernel's signature inconsistent across those passes.
Please follow the existing kernels in mmvq.cu: accept unqualified vx_ptr, vy_ptr, and dst_ptr parameters, then introduce GGML_CUDA_RESTRICT local aliases inside the kernel body. Removing the formal parameter qualifiers is another minimal candidate. These fixes are suggested from the diagnostic and existing pattern; a patched rebuild has not been run. This is a compile-time failure, so setting the runtime GGML_CUDA_PDL environment variable does not resolve it.
No head performance or correctness results can be reported on this configuration until the build succeeds.
|
Agent benchmark follow-up, posted at the maintainer's request. Pinned head
Selected CPU-reference backend checks passed on the compared arms. The CUDA 12.8 build blocker on Hopper/Blackwell is documented in the earlier review. No patched source was substituted. These timings do not establish batch invariance or speculative acceptance. The percentages describe these paired runs; small changes should not be interpreted as established improvements. No long-context, multi-slot serving, or end-to-end logit-parity claim is made. Exploratory small-batch throughput changes (one pair, three repetitions per invocation):
|
…per-row small-K GEMV (+14% TG)
Two changes to the batch-1 PTQ1_0 mat-vec path, found by tracing a Bonsai 2 27B decode graph
on an RTX 4070 (sm_89) with CUPTI. Same weights, same bits, output byte-identical.
1. Exact integer sums + warp-transposed (SoA) q8_1 activation layout.
quantize_q8_1<exact_isum> stores the int16 sum of each 32-value block in ds.y (instead of
the float input sum) and writes the activations grouped by 32 K-blocks with word w of the
group contiguous. The ternary vec-dot then accumulates raw trits {0,1,2} with dp4a and
corrects once per block, removing a __vsub4 per 4 weights (no native SIMD byte subtract on
sm_8x), and a warp-wide load of "word w" in the small-K geometry hits 1 L1 wavefront instead
of ~36 (lanes were reading block_q8_1 structs 144 B apart). That LSU pressure, not GDDR,
capped every PTQ1_0 GEMV at ~370 GB/s. Bytes per column are unchanged when K is padded to
32*128. Only PTQ1_0 opts in (ggml_cuda_q8_1_exact_isum); other types keep the AoS layout.
2. Warp-per-row geometry for small-K PTQ1_0 GEMVs.
The generic small_k loop strides K-blocks across the whole 128-thread block, so for K=5120
(40 blocks) only threads 0..39 ever load anything and two of four warps just hold SM slots
at __syncthreads. Each warp now owns one row and its 32 lanes stride that row's K-blocks;
shuffle reduction, no shared memory, no barrier. qkv 391 -> 445 GB/s, lm_head 414 -> 460.
The recurrent-state gather fusion that was part of the first revision of this PR is now its
own PR so that this activation layout can be compared with PrismML-Eng#218 on its own.
RTX 4070 12 GB, Ternary-Bonsai-2-27B-PTQ1_0, llama-bench -fa 1 -ctk/-ctv q4_0, stock clocks,
three alternating rounds of r=3:
prism @ 9a9394a 51.0 t/s tg128
+ 1 + 2 (this PR) 56.4 (+10.5%)
+ GDN gather fusion (follow-up) 58.5 (+14.7%)
test-backend-ops CUDA0 vs CPU: MUL_MAT ptq1_0 45/45 supported cases, MUL_MAT_ID ptq1_0 75/75.
400 greedy tokens identical with the change on and off.
… (+3.8% TG on Ada) build_rs gathers each layer's live recurrent state via GET_ROWS(cache, s_copy) into a temp (3 MB per layer for Bonsai 2 27B) that only the GDN kernel reads. Traced in-graph with CUPTI on an RTX 4070 the gather kernel costs ~8 us/layer, and its dirty output sits in L2 until the FFN weight stream evicts it, which made the gate/up GEMVs in GDN layers ~15% slower than the identical GEMVs in attention layers. When the GET_ROWS output has no consumer other than (optionally through a RESHAPE) the GDN node's state input, the graph evaluator skips the GET_ROWS node and records the gather (cache base, ids, row stride) for that GDN node. The kernel then indexes the cache row s_ids[seq] directly; no temp, no kernel, no writeback. Single-sequence only: with several sequences a gathered row may alias a row another sequence writes in the same op. GGML_CUDA_GDN_GATHER_FUSION=0 disables it. The registrations live in the evaluating ggml_backend_cuda_context (ggml_cuda_gdn_gather_context, beside the concurrent-stream state), keyed by node pointer and reset at the start of every graph evaluation, so two contexts or host threads decoding at once never see each other's skips. This replaces the process-wide static map in the first revision of PrismML-Eng#215. Independent of the activation layout change; split out of PrismML-Eng#215 so that layout can be compared with PrismML-Eng#218 on its own. RTX 4070 12 GB, Ternary-Bonsai-2-27B-PTQ1_0, llama-bench -fa 1 -ctk/-ctv q4_0, stock clocks, three alternating rounds of r=3: tg128 56.4 -> 58.5 t/s on top of PrismML-Eng#215 (+3.8%). pp512 unchanged. test-backend-ops CUDA0 vs CPU: GATED_DELTA_NET 39/39 supported cases, GET_ROWS unchanged. 400 greedy tokens identical with GGML_CUDA_GDN_GATHER_FUSION=0 and 1.
…per-row small-K GEMV (+14% TG)
Two changes to the batch-1 PTQ1_0 mat-vec path, found by tracing a Bonsai 2 27B decode graph
on an RTX 4070 (sm_89) with CUPTI. Same weights, same bits, output byte-identical.
1. Exact integer sums + warp-transposed (SoA) q8_1 activation layout.
quantize_q8_1<exact_isum> stores the int16 sum of each 32-value block in ds.y (instead of
the float input sum) and writes the activations grouped by 32 K-blocks with word w of the
group contiguous. The ternary vec-dot then accumulates raw trits {0,1,2} with dp4a and
corrects once per block, removing a __vsub4 per 4 weights (no native SIMD byte subtract on
sm_8x), and a warp-wide load of "word w" in the small-K geometry hits 1 L1 wavefront instead
of ~36 (lanes were reading block_q8_1 structs 144 B apart). That LSU pressure, not GDDR,
capped every PTQ1_0 GEMV at ~370 GB/s. Bytes per column are unchanged when K is padded to
32*128. Only PTQ1_0 opts in (ggml_cuda_q8_1_exact_isum); other types keep the AoS layout.
2. Warp-per-row geometry for small-K PTQ1_0 GEMVs.
The generic small_k loop strides K-blocks across the whole 128-thread block, so for K=5120
(40 blocks) only threads 0..39 ever load anything and two of four warps just hold SM slots
at __syncthreads. Each warp now owns one row and its 32 lanes stride that row's K-blocks;
shuffle reduction, no shared memory, no barrier. qkv 391 -> 445 GB/s, lm_head 414 -> 460.
The recurrent-state gather fusion that was part of the first revision of this PR is now its
own PR so that this activation layout can be compared with PrismML-Eng#218 on its own.
RTX 4070 12 GB, Ternary-Bonsai-2-27B-PTQ1_0, llama-bench -fa 1 -ctk/-ctv q4_0, stock clocks,
three alternating rounds of r=3:
prism @ 9a9394a 51.0 t/s tg128
+ 1 + 2 (this PR) 56.4 (+10.5%)
+ GDN gather fusion (follow-up) 58.5 (+14.7%)
test-backend-ops CUDA0 vs CPU: MUL_MAT ptq1_0 45/45 supported cases, MUL_MAT_ID ptq1_0 75/75.
400 greedy tokens identical with the change on and off.
|
Ada numbers for the comparison, since the maintainer runs only had #218 on the 4090. RTX 4070 12 GB,
Same picture the 3090/4090 pair suggested: your kernel's win is at 2-8 columns and on Ampere, the #215 layout's win is at one column on Ada. Neither covers the other's case yet. The pp8 gap the other way is worth a look on your side: with I am fine with the split you proposed. Concretely, the pieces that need to end up on one layout are: the exact int16 per-block sums (bias subtracted once per block, which is also what makes your 3060 numbers for #215 whenever the card is free would close the loop. Raw JSON for the table: |
|
Follow-up: #221 is the integration branch with #217, #218, #214, #215, #216 and #220 together plus the two pieces that make the combination fit on 12 GB: a hybrid dispatch that uses the #215 layout at one column and the #218 kernel at 2-4 columns (one Result on an RTX 4070 12 GB: Bonsai 2 27B with the MTP head at the full 262,144 window resident, 86 tok/s greedy three-prompt mean (64 without the draft head, 51 stock), byte-identical output under |
|
Confirming @bri-prism's P1 on The failure is not architecture-specific. On Windows with MSVC 14.44 building Root cause: template<int A> void f(const void * __restrict &);
template<> void f<1>(const void *&); // cl.exe: error C2912Fix, mirroring the pattern already used in this file for --- a/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh
+++ b/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh
@@ -281,10 +281,17 @@
static __global__ void mul_mat_vec_ptq1_0_pt(
- const void * GGML_CUDA_RESTRICT vx, const void * GGML_CUDA_RESTRICT vy, const ggml_cuda_mm_fusion_args_device fusion,
- float * GGML_CUDA_RESTRICT dst,
+ const void * vx_, const void * vy_, const ggml_cuda_mm_fusion_args_device fusion,
+ float * dst_,
const int ncols_x, const int nrows_x, const int stride_row_x, const int stride_col_y, const int stride_col_dst,
const int rows_per_cta, const uint3 bpr_fd, const uint3 rpc_fd) {
+ const void * GGML_CUDA_RESTRICT vx = vx_;
+ const void * GGML_CUDA_RESTRICT vy = vy_;
+ float * GGML_CUDA_RESTRICT dst = dst_;
extern __shared__ float partials[];With that hunk the tree builds clean for Two notes that may matter for the Hopper/Blackwell builds that currently fail:
Measured after the fix on an RTX 5070 Ti Laptop (sm_120), paired runs against the stock fork binary, same PTQ1_0 file, |
|
The C2912 / The 5070 Ti timings attached to that comment are not this PR. #218 does not touch the MMQ tile loader. Standalone, #218's own win, as already measured: 2–8 column PTQ1_0 and Ampere decode (3090 tg128 +27%; 4070 pp2 65→92, pp4 73→149). Ada batch-1 is the #215 layout. Please keep those columns on this thread so review does not think this patch doubles prefill. |
|
On the one real code finding (C2912 / It is not in So: apply the local-alias pattern already used by |
|
Pushed
Regression check on an RTX 3060 (sm_86), same PTQ1_0 file, The reconciliation with #215 stays as proposed above. #221 already carries this branch's commits and the hybrid dispatch, so |
C2912 / stub mismatch on sm_90 and sm_120 when GGML_CUDA_RESTRICT is on the formal parameters of mul_mat_vec_ptq1_0_pt (same hunk as 2578fdf on PrismML-Eng#218). Ampere (cc 800-889) takes the planar 1-col kernel; Ada and newer keep SoA. 3060: PrismML-Eng#218 was +5.9% tg128 vs PrismML-Eng#215 at one column.
| if (!ptq1_0_pt_enabled() || nchannels_dst != 1 || nsamples_dst != 1 || ncols_x % QK_PTQ1_0 != 0 || | ||
| ncols_dst < 1 || ncols_dst > PTQ1_0_PT_MAX_COLS || 2 * (ncols_x / QK_PTQ1_0) * ncols_dst > PTQ1_0_PT_SMEM_FLOATS) { | ||
| return false; | ||
| } |
| // GGML_CUDA_BATCH_INVARIANT=1: prefer kernels whose per-column arithmetic does not depend on the | ||
| // number of columns in the batch, so that a token decoded alone and a token verified inside a | ||
| // speculative batch see the same logits bit for bit. Whole-model invariance is established for | ||
| // batches of 1 to 4 columns; 5 to 8 columns agree with each other but can differ from 1 to 4. | ||
| // Costs some throughput at 2 to 8 columns. |
|
Agent review follow-up: the exact-tree delta to current head |
…ag changes The comment in common.cuh read as a whole-model promise. The flag only changes the F16 and BF16 mat-vec paths it selects, the PTQ1_0 mat-vec and flash attention up to 8 queries, for batches of 1 to 4 columns; other weight types and attention shapes outside those kernels can still pick batch-dependent kernels. Say that, and keep the 1 to 4 wording.
…MMQ tile path With the branch-free PTQ1_0 MMQ tile loader (PrismML-Eng#214) in the tree, the MMQ path is the faster one from 5 columns on. On an RTX 3060 with the K = 5120 projections of Bonsai 2 27B, llama-bench -p 8 runs a batch of 8 in 64.4 ms through the new MMQ tiles against 112.7 ms through the 8-column mat-vec (five-arm a/b, r=3, github.com/sudoingX/bonsai2-small-gpu kernel/ab/ab_215_214_rtx3060.md). The mat-vec keeps its lead at 2 to 4 columns (1.20x / 1.49x / 1.77x of a single pass against 1.46x / 2.21x / 2.75x for the tile path), which is the range speculative verification uses. ggml_cuda_should_use_mmvq now routes PTQ1_0 to the mat-vec for ne11 <= 4 (PTQ1_0_PT_MAX_COLS, was 8) and the 5 to 8 column instantiations of the dedicated kernel are gone. GGML_CUDA_BATCH_INVARIANT keeps its 1 to 4 column guarantee; the comment no longer mentions 5 to 8, which the flag never covered once those batches take the tile path.
3443dde to
285542d
Compare
|
Rebased onto the current The mat-vec column limit stops at 4. sm_86 after the rebase, RTX 3060 12GB, CUDA 12.4, original Ternary-Bonsai-2-27B-PTQ1_0.gguf, head 285542d:
Both earlier findings are closed. @bri-prism reviewed 3443dde on Sep 21: the guard and the launcher share Could you re-review when you have a moment? The changes-requested is from Sep 19 and refers to the restrict failure fixed in 2578fdf. |
bri-prism
left a comment
bri-prism
left a comment
There was a problem hiding this comment.
Re-reviewed at 285542d. Every point I raised is addressed, the standing build condition is met, and the new column cap is measured rather than asserted. Clearing my change request.
P1, the restrict qualifiers. Resolved as suggested: the formal parameters of mul_mat_vec_ptq1_0_pt are unqualified with local GGML_CUDA_RESTRICT aliases inside, and the comment at mmvq-ptq1_0.cuh:295 records why, which is the part that stops someone reinstating it later. Independently build-verified by @AlexGabbia on sm_120 with CUDA 13.0 and MSVC, 405/405 targets and zero errors. That also satisfies the fresh Hopper/Blackwell build check I asked to retain before merge, so consider that condition discharged.
The batch-invariance documentation. Narrowed from 1 to 8 down to 1 to 4 in both common.cuh and fattn.cu, with the 5 to 8 behaviour stated explicitly rather than left implied, and the code predicate unchanged. That was the ask.
The shared-memory budgeting fix. I reviewed 844dd06 separately and had no finding. The entry guard and the launcher now compute the same number through ptq1_0_pt_smem_bytes(), including rows per CTA and the gate multiplier, with oversized launches falling back to the generic path.
The new column cap is the right shape and I checked it holds in code. PTQ1_0_PT_MAX_COLS is 4 at mmvq-ptq1_0.cuh:35, the kernel entry guard rejects outside 1 to 4 at :448, and ggml_cuda_should_use_mmvq routes on ne11 <= PTQ1_0_PT_MAX_COLS at mmvq.cu:304. Host routing and device guard read the same constant rather than two copies of the same number, which is what keeps a future edit from splitting them. This is the failure mode where a dispatch gate and the kernel it dispatches to disagree about their own limits, and it is avoided here by construction.
The justification is measured, on an RTX 3060 with the K = 5120 projections, five arms at r = 3, 64.4 ms through the MMQ tiles against 112.7 ms through the 8-column mat-vec at a batch of 8, with the mat-vec keeping its lead at 2 to 4 columns where speculative verification actually lives. Handing the range you lose to #214 rather than defending it is the right call, and publishing the arms is what makes the cap reviewable instead of a magic number.
One thing still outstanding, and it is sequencing rather than code. The activation layout still has to converge with #215. I approved #215 earlier today on its own contents, so both are now clear on their merits while carrying different layouts, and #221 is where the hybrid dispatch reconciles them. I have not reviewed #221. Whoever merges these should decide the layout there rather than letting merge order decide it, because merge order will otherwise pick a winner silently.
Method note: source reading at 285542d plus the build and benchmark reports on this thread. I did not rebuild or rerun anything myself, and the 3060 numbers are the author's.
… (+3.8% TG on Ada) build_rs gathers each layer's live recurrent state via GET_ROWS(cache, s_copy) into a temp (3 MB per layer for Bonsai 2 27B) that only the GDN kernel reads. Traced in-graph with CUPTI on an RTX 4070 the gather kernel costs ~8 us/layer, and its dirty output sits in L2 until the FFN weight stream evicts it, which made the gate/up GEMVs in GDN layers ~15% slower than the identical GEMVs in attention layers. When the GET_ROWS output has no consumer other than (optionally through a RESHAPE) the GDN node's state input, the graph evaluator skips the GET_ROWS node and records the gather (cache base, ids, row stride) for that GDN node. The kernel then indexes the cache row s_ids[seq] directly; no temp, no kernel, no writeback. Single-sequence only: with several sequences a gathered row may alias a row another sequence writes in the same op. GGML_CUDA_GDN_GATHER_FUSION=0 disables it. The registrations live in the evaluating ggml_backend_cuda_context (ggml_cuda_gdn_gather_context, beside the concurrent-stream state), keyed by node pointer and reset at the start of every graph evaluation, so two contexts or host threads decoding at once never see each other's skips. This replaces the process-wide static map in the first revision of PrismML-Eng#215. Independent of the activation layout change; split out of PrismML-Eng#215 so that layout can be compared with PrismML-Eng#218 on its own. RTX 4070 12 GB, Ternary-Bonsai-2-27B-PTQ1_0, llama-bench -fa 1 -ctk/-ctv q4_0, stock clocks, three alternating rounds of r=3: tg128 56.4 -> 58.5 t/s on top of PrismML-Eng#215 (+3.8%). pp512 unchanged. test-backend-ops CUDA0 vs CPU: GATED_DELTA_NET 39/39 supported cases, GET_ROWS unchanged. 400 greedy tokens identical with GGML_CUDA_GDN_GATHER_FUSION=0 and 1.
…per-row small-K GEMV (+14% TG) Merged onto the planar-transposed activation layout and dedicated 2-8 column kernel (PrismML-Eng#218, sudoingX). One decision, ggml_cuda_q8_1_layout_for(type, ncols_dst, ids), is evaluated by both the quantizer call and the kernel switch: PTQ1_0, 1 column, no ids -> SOA_ISUM (this change): exact int16 sums, warp-per-row small-K geometry PTQ1_0, 2-8 columns / MoE -> PT (PrismML-Eng#218): shared weight decode, full lane utilization anything else -> plain block_q8_1 GGML_CUDA_BATCH_INVARIANT=1 forces the one-column case onto PT as well so a token decoded alone and a token verified inside a speculative batch still see identical arithmetic.
C2912 / stub mismatch on sm_90 and sm_120 when GGML_CUDA_RESTRICT is on the formal parameters of mul_mat_vec_ptq1_0_pt (same hunk as 2578fdf on PrismML-Eng#218). Ampere (cc 800-889) takes the planar 1-col kernel; Ada and newer keep SoA. 3060: PrismML-Eng#218 was +5.9% tg128 vs PrismML-Eng#215 at one column.
… (+3.8% TG on Ada) build_rs gathers each layer's live recurrent state via GET_ROWS(cache, s_copy) into a temp (3 MB per layer for Bonsai 2 27B) that only the GDN kernel reads. Traced in-graph with CUPTI on an RTX 4070 the gather kernel costs ~8 us/layer, and its dirty output sits in L2 until the FFN weight stream evicts it, which made the gate/up GEMVs in GDN layers ~15% slower than the identical GEMVs in attention layers. When the GET_ROWS output has no consumer other than (optionally through a RESHAPE) the GDN node's state input, the graph evaluator skips the GET_ROWS node and records the gather (cache base, ids, row stride) for that GDN node. The kernel then indexes the cache row s_ids[seq] directly; no temp, no kernel, no writeback. Single-sequence only: with several sequences a gathered row may alias a row another sequence writes in the same op. GGML_CUDA_GDN_GATHER_FUSION=0 disables it. The registrations live in the evaluating ggml_backend_cuda_context (ggml_cuda_gdn_gather_context, beside the concurrent-stream state), keyed by node pointer and reset at the start of every graph evaluation, so two contexts or host threads decoding at once never see each other's skips. This replaces the process-wide static map in the first revision of PrismML-Eng#215. Independent of the activation layout change; split out of PrismML-Eng#215 so that layout can be compared with PrismML-Eng#218 on its own. RTX 4070 12 GB, Ternary-Bonsai-2-27B-PTQ1_0, llama-bench -fa 1 -ctk/-ctv q4_0, stock clocks, three alternating rounds of r=3: tg128 56.4 -> 58.5 t/s on top of PrismML-Eng#215 (+3.8%). pp512 unchanged. test-backend-ops CUDA0 vs CPU: GATED_DELTA_NET 39/39 supported cases, GET_ROWS unchanged. 400 greedy tokens identical with GGML_CUDA_GDN_GATHER_FUSION=0 and 1.
…per-row small-K GEMV (+14% TG) Merged onto the planar-transposed activation layout and dedicated 2-8 column kernel (PrismML-Eng#218, sudoingX). One decision, ggml_cuda_q8_1_layout_for(type, ncols_dst, ids), is evaluated by both the quantizer call and the kernel switch: PTQ1_0, 1 column, no ids -> SOA_ISUM (this change): exact int16 sums, warp-per-row small-K geometry PTQ1_0, 2-8 columns / MoE -> PT (PrismML-Eng#218): shared weight decode, full lane utilization anything else -> plain block_q8_1 GGML_CUDA_BATCH_INVARIANT=1 forces the one-column case onto PT as well so a token decoded alone and a token verified inside a speculative batch still see identical arithmetic.
C2912 / stub mismatch on sm_90 and sm_120 when GGML_CUDA_RESTRICT is on the formal parameters of mul_mat_vec_ptq1_0_pt (same hunk as 2578fdf on PrismML-Eng#218). Ampere (cc 800-889) takes the planar 1-col kernel; Ada and newer keep SoA. 3060: PrismML-Eng#218 was +5.9% tg128 vs PrismML-Eng#215 at one column.
|
Kernel-level numbers from an RTX 4090 (sm_89, CUDA 12.8), comparing this PR's head (285542d) with #215's head (ad61b84) and
The dedicated mat-vec does what it says for plain MUL_MAT. As the description notes, MoE calls stay on the generic kernel. On the first expert shape that path is about 14% slower than on |
…per-row small-K GEMV (+14% TG)
Two changes to the batch-1 PTQ1_0 mat-vec path, found by tracing a Bonsai 2 27B decode graph
on an RTX 4070 (sm_89) with CUPTI. Same weights, same bits, output byte-identical.
1. Exact integer sums + warp-transposed (SoA) q8_1 activation layout.
quantize_q8_1<exact_isum> stores the int16 sum of each 32-value block in ds.y (instead of
the float input sum) and writes the activations grouped by 32 K-blocks with word w of the
group contiguous. The ternary vec-dot then accumulates raw trits {0,1,2} with dp4a and
corrects once per block, removing a __vsub4 per 4 weights (no native SIMD byte subtract on
sm_8x), and a warp-wide load of "word w" in the small-K geometry hits 1 L1 wavefront instead
of ~36 (lanes were reading block_q8_1 structs 144 B apart). That LSU pressure, not GDDR,
capped every PTQ1_0 GEMV at ~370 GB/s. Bytes per column are unchanged when K is padded to
32*128. Only PTQ1_0 opts in (ggml_cuda_q8_1_exact_isum); other types keep the AoS layout.
2. Warp-per-row geometry for small-K PTQ1_0 GEMVs.
The generic small_k loop strides K-blocks across the whole 128-thread block, so for K=5120
(40 blocks) only threads 0..39 ever load anything and two of four warps just hold SM slots
at __syncthreads. Each warp now owns one row and its 32 lanes stride that row's K-blocks;
shuffle reduction, no shared memory, no barrier. qkv 391 -> 445 GB/s, lm_head 414 -> 460.
The recurrent-state gather fusion that was part of the first revision of this PR is now its
own PR so that this activation layout can be compared with PrismML-Eng#218 on its own.
RTX 4070 12 GB, Ternary-Bonsai-2-27B-PTQ1_0, llama-bench -fa 1 -ctk/-ctv q4_0, stock clocks,
three alternating rounds of r=3:
prism @ 9a9394a 51.0 t/s tg128
+ 1 + 2 (this PR) 56.4 (+10.5%)
+ GDN gather fusion (follow-up) 58.5 (+14.7%)
test-backend-ops CUDA0 vs CPU: MUL_MAT ptq1_0 45/45 supported cases, MUL_MAT_ID ptq1_0 75/75.
400 greedy tokens identical with the change on and off.
… (+3.8% TG on Ada) build_rs gathers each layer's live recurrent state via GET_ROWS(cache, s_copy) into a temp (3 MB per layer for Bonsai 2 27B) that only the GDN kernel reads. Traced in-graph with CUPTI on an RTX 4070 the gather kernel costs ~8 us/layer, and its dirty output sits in L2 until the FFN weight stream evicts it, which made the gate/up GEMVs in GDN layers ~15% slower than the identical GEMVs in attention layers. When the GET_ROWS output has no consumer other than (optionally through a RESHAPE) the GDN node's state input, the graph evaluator skips the GET_ROWS node and records the gather (cache base, ids, row stride) for that GDN node. The kernel then indexes the cache row s_ids[seq] directly; no temp, no kernel, no writeback. Single-sequence only: with several sequences a gathered row may alias a row another sequence writes in the same op. GGML_CUDA_GDN_GATHER_FUSION=0 disables it. The registrations live in the evaluating ggml_backend_cuda_context (ggml_cuda_gdn_gather_context, beside the concurrent-stream state), keyed by node pointer and reset at the start of every graph evaluation, so two contexts or host threads decoding at once never see each other's skips. This replaces the process-wide static map in the first revision of PrismML-Eng#215. Independent of the activation layout change; split out of PrismML-Eng#215 so that layout can be compared with PrismML-Eng#218 on its own. RTX 4070 12 GB, Ternary-Bonsai-2-27B-PTQ1_0, llama-bench -fa 1 -ctk/-ctv q4_0, stock clocks, three alternating rounds of r=3: tg128 56.4 -> 58.5 t/s on top of PrismML-Eng#215 (+3.8%). pp512 unchanged. test-backend-ops CUDA0 vs CPU: GATED_DELTA_NET 39/39 supported cases, GET_ROWS unchanged. 400 greedy tokens identical with GGML_CUDA_GDN_GATHER_FUSION=0 and 1.
…per-row small-K GEMV (+14% TG) Merged onto the planar-transposed activation layout and dedicated 2-8 column kernel (PrismML-Eng#218, sudoingX). One decision, ggml_cuda_q8_1_layout_for(type, ncols_dst, ids), is evaluated by both the quantizer call and the kernel switch: PTQ1_0, 1 column, no ids -> SOA_ISUM (this change): exact int16 sums, warp-per-row small-K geometry PTQ1_0, 2-8 columns / MoE -> PT (PrismML-Eng#218): shared weight decode, full lane utilization anything else -> plain block_q8_1 GGML_CUDA_BATCH_INVARIANT=1 forces the one-column case onto PT as well so a token decoded alone and a token verified inside a speculative batch still see identical arithmetic.
C2912 / stub mismatch on sm_90 and sm_120 when GGML_CUDA_RESTRICT is on the formal parameters of mul_mat_vec_ptq1_0_pt (same hunk as 2578fdf on PrismML-Eng#218). Ampere (cc 800-889) takes the planar 1-col kernel; Ada and newer keep SoA. 3060: PrismML-Eng#218 was +5.9% tg128 vs PrismML-Eng#215 at one column.
|
Heads-up from testing this branch on an RTX 5080 (sm_120): at The fix is one line, at no speed cost: sudoingX#1. Isolation table and the full 5080 numbers: sudoingX/bonsai2-small-gpu#3. |
|
Same hole is in #221 (this kernel is the hybrid 2?4 column path). Picked up the one-line |
|
Thanks @LamplighterPaul, seriously appreciate you catching this and doing the isolation work. My 4070 would never have exposed it. The fix is in our #221 combo branch and patch 0023 in the setup bundle. It uses the existing PDL dependency wait before the kernel reads the quantized activations. Your PR on Sudo's fork is still the route into #218. Validation is now green:
For anyone on the Linux bundle, from the repo directory: git pull --ff-only
bash build/build_linux.shThen restart with the rebuilt server. Existing downloaded binaries don't pick this up just from pulling the patches. Your 5080 isolation remains the runtime evidence; I haven't independently rerun it on Hopper/Blackwell and I'm not claiming a new performance result. If anything still breaks on the fixed combo branch, send me the commit, model, launch command and failing output and I'll dig into our side. Thanks again for helping get this right. |
…ernary GEMV The multi-row decode-once kernel ports the approach of PrismML-Eng/llama.cpp#218 (CUDA: fast, batch-invariant small-batch PTQ1_0 mat-vec) by @sudoingX, by way of its Metal port on jasontitus/llama.cpp. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SHJDMbEnsswDXJU3LAQQZ7
|
@professorpalmer on it, i can easily re-run it shouldn't be a problem. |
…ernary GEMV The multi-row decode-once kernel ports the approach of PrismML-Eng/llama.cpp#218 (CUDA: fast, batch-invariant small-batch PTQ1_0 mat-vec) by @sudoingX, by way of its Metal port on jasontitus/llama.cpp. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01SHJDMbEnsswDXJU3LAQQZ7
…on (Bonsai 1 + 2) (#530) * perf: 2-bit ternary GEMV for 1..8 rows, per GPU generation, on Bonsai 1 too Ports the multi-column mat-vec idea from the Bonsai Metal work on the llama.cpp fork: each packed 2-bit word decodes once and serves every activation row (plain decode, batched decode, MTP verify), so 2..8-row steps stop paying one full decode per row. - qmv2: one `msv_qmv2_rows` kernel over M = 1..8, bf16 or f16 activations (x read as half2 for f16, float2 for bf16), R rows per simdgroup and G simdgroups per threadgroup as template args. Needs biases == -scales. - `planFor` picks the geometry per GPU generation from kernel sweeps on the Bonsai 27B shapes: g13 (M1 Ultra) R4G2; g16 (M4 Pro) R4G2 odd M, R2G8 even; g17 (M5 Max) R4G2 at M=1 f16, R4G8/R2G8 above, stock at M=5 and bf16 M=1. Unmeasured generations and phones keep the previous dispatch exactly. - Non-Hadamard ternary packs (Prism Bonsai 1) get the kernel: a load-time scan arms `ternary_2bit` only when every matmul weight's biases are exactly -scales (an embedding only gathered is exempt; the first offending tensor is logged otherwise). - On measured generations a generic-bias weight of a Hadamard pack keeps the bias-aware `qmv` at M=1. Tests: parity vs f32 truth at every M, geometry, dtype and bias layout; per-generation routing incl. legacy byte-for-byte; the ternary scan; the qmatmul wiring (red when the ternary branch is removed). Rejected after measurement: a fused gate/up/SwiGLU GEMV (bit-exact, ~1% at M=1 on M1 Ultra, slower at M>=2). The GDN state gather/copy-back removed on the fork does not exist here. * perf: qmv2 plans send narrow outputs and M5's non-MLP single rows to stock Kernel sweeps on the remaining Bonsai shapes: below 2048 output rows (GDN a/b at 48, attention k/v at 1024) too few threadgroups stream K and stock is up to 2x faster on M1 Ultra and M5 Max at every width; on M5 Max the single-row kernel only wins at the MLP width (lm_head 1.06x, 17408 rows) and loses 3-12% below it. Unmeasured generations keep the legacy dispatch. * docs: Bonsai ternary kernel CHANGELOG line and the per-chip tuning gotcha * docs: credit @sudoingX's CUDA small-batch kernels behind the Bonsai ternary GEMV The multi-row decode-once kernel ports the approach of PrismML-Eng/llama.cpp#218 (CUDA: fast, batch-invariant small-batch PTQ1_0 mat-vec) by @sudoingX, by way of its Metal port on jasontitus/llama.cpp. * docs: leave CHANGELOG to the v26.9.6 release PR (#518) The entry moves to the PR description so it lands in #518's Changes list instead of a second v26.9.6 header that conflicts with #518 and #531. * test: tell the ternary route from stock by bias, not by rounding luck The routing test assumed the kernel and stock round differently on a random fixture. On M4 Max they came out bit-identical, so the vacuity guard failed. Biases of +scale make them disagree by construction: the ternary kernel never reads biases.
|
Ran #221 (09b6cce) on the RTX 5080 (sm_120, CUDA 13.3). The PDL fix holds and output is coherent across text, image, code, prose and bash. Decode is +8% over our fixed #218 build (tg128 104.8 vs 96.9), MTP n-max 2 gets 169 / 137 / 173 tok/s on code / prose / bash at 131K, and VRAM is ~470 MiB lower. Nice work. Two small notes, in case they're useful:
|
|
Thanks for running The +8% tg128 and the 131K MTP numbers are the hybrid stack (#215 one-column + this PR at 2–4 + in-place FA + GDN). Good to have them on a Blackwell part. Both notes are real and belong on #221 / #220, not here.
No patch from this box; we'll pick those two up on the combo branch. |
… Ada: #218 + #215 hybrid PTQ1_0 dispatch, in-place q4_0/q8_0 K/V flash attention (integration checkpoint) (#221) * Add: planar-transposed activation layout for the PTQ1_0 mat-vec path The mmvq path gave every thread one 128-weight PTQ1_0 block and read the activations as 32 scattered 4-byte loads per column out of 36-byte block_q8_1 structs, so every extra column cost a full pass: 47 / 104 / 154 / 201 us for 1 / 2 / 3 / 4 columns on a 4096 x 14336 matrix (RTX 3060, test-backend-ops perf). Quantize PTQ1_0 activations into a planar-transposed layout instead: per column, 8 planes of 16-byte quant pieces and one plane of the 4 (d, s) half2 scales, 16 bytes per 128-element block each, with the same bytes per row and the same column stride as block_q8_1. A thread then reads its block's activations as 9 aligned 16-byte loads with adjacent lanes on adjacent pieces: 46 / 55 / 65 / 75 us. The trit decode is unchanged; the byte-wise minus one is (q + 0x7F7F7F7F) ^ 0x80808080 instead of __vsub4. nwarps is pinned to 4 for every column count so the K partition, and with it the fp32 summation order, does not depend on the batch. mmvq now handles PTQ1_0 up to 8 columns (MMQ took 395 us at 8). HIP keeps the old layout. * Add: dedicated PTQ1_0 mat-vec kernel with full lane utilization With 128-element blocks the generic kernel's 128 threads cover 128 K blocks per iteration, so a K = 5120 projection (40 blocks) keeps 31% of the lanes busy and K = 17408 (136 blocks) 53%. This kernel handles plain 2D MUL_MAT (no batch dims, no expert ids, K a multiple of 128, up to 8 columns): it flattens (row group, K block) pairs into one index space with rows_per_cta chosen on the host to fill whole 128-thread iterations, gives each thread 4 rows per K block (2 at 5 to 8 columns) for latency hiding and activation reuse, stores one fp32 partial per (row, column, K block) in dynamic shared memory and has one warp per (row, column) sum them in a fixed order (lane-strided sequential, then a butterfly). The per-block partial is d * sum_k d8_k * sumi_k with exact integer sumi_k, written as __fmaf_rn and __fmul_rn so every instantiation rounds alike, and the order depends only on the weight shape: a column's result is bit-identical for every column count 1 to 8. Fusion (bias, gated GLU epilogue) is supported for one column as in the generic kernel, which remains the path for batched and MoE calls. Bonsai 2 shapes, RTX 3060 12GB, us per call, 1 column: attn_qkv 5120 x 10240 72.4 to 41.8, ffn_up 5120 x 17408 120.0 to 66.7, attn_gate 5120 x 6144 44.8 to 27.7. Whole model llama-bench tg32 26.1 to 39.8 tok/s, pp2 34.1 to 61.1, pp3 33.3 to 72.0. * Test: add Bonsai 2 projection shapes to the mul_mat perf cases * Add: GGML_CUDA_BATCH_INVARIANT for batch-invariant small-batch kernels Off by default, nothing changes. When set, kernels are picked so that the per-column arithmetic of a token does not depend on how many columns (1 to 8) share the launch: F16 and BF16 matrices with at most 512 rows take the mul_mat_vec_f kernel for 1 to 8 columns on Ampere and newer (its K partition depends on K only), flash attention with up to 8 queries takes the vector kernel on Turing and newer instead of switching to the tensor-core kernel at 2 queries, and the KV split (parallel_blocks) is sized as for one query tile so the partial-softmax combine order does not depend on the batch. Together with the batch-invariant PTQ1_0 mat-vec this makes the logits of a token bit-identical whether it is decoded alone or verified inside a speculative batch of 2, 3 or 4 tokens (measured on Ternary Bonsai 2 27B: greedy output with draft-mtp equals greedy output without it on three prompts at n-max 1, 2 and 3). Batches of 5 to 8 agree with each other but not with 1 to 4. Costs throughput at depth (the vector attention kernel at 3 queries over 42K keys). * Fix: use the mat-vec kernel for bf16 matrices under 64 rows at 2 to 8 columns The qwen35 gated-delta-net gate projections are bf16 [5120 x 48]. At 2 to 8 columns on Ampere they took a 10.5 us path where mul_mat_vec_f takes 4.4 to 7.3 us (test-backend-ops perf, RTX 3060); at one column both already used mul_mat_vec_f (3.3 us). 96 calls per token on the 27B model. * cuda: fold the recurrent-state gather into the GATED_DELTA_NET kernel (+3.8% TG on Ada) build_rs gathers each layer's live recurrent state via GET_ROWS(cache, s_copy) into a temp (3 MB per layer for Bonsai 2 27B) that only the GDN kernel reads. Traced in-graph with CUPTI on an RTX 4070 the gather kernel costs ~8 us/layer, and its dirty output sits in L2 until the FFN weight stream evicts it, which made the gate/up GEMVs in GDN layers ~15% slower than the identical GEMVs in attention layers. When the GET_ROWS output has no consumer other than (optionally through a RESHAPE) the GDN node's state input, the graph evaluator skips the GET_ROWS node and records the gather (cache base, ids, row stride) for that GDN node. The kernel then indexes the cache row s_ids[seq] directly; no temp, no kernel, no writeback. Single-sequence only: with several sequences a gathered row may alias a row another sequence writes in the same op. GGML_CUDA_GDN_GATHER_FUSION=0 disables it. The registrations live in the evaluating ggml_backend_cuda_context (ggml_cuda_gdn_gather_context, beside the concurrent-stream state), keyed by node pointer and reset at the start of every graph evaluation, so two contexts or host threads decoding at once never see each other's skips. This replaces the process-wide static map in the first revision of #215. Independent of the activation layout change; split out of #215 so that layout can be compared with #218 on its own. RTX 4070 12 GB, Ternary-Bonsai-2-27B-PTQ1_0, llama-bench -fa 1 -ctk/-ctv q4_0, stock clocks, three alternating rounds of r=3: tg128 56.4 -> 58.5 t/s on top of #215 (+3.8%). pp512 unchanged. test-backend-ops CUDA0 vs CPU: GATED_DELTA_NET 39/39 supported cases, GET_ROWS unchanged. 400 greedy tokens identical with GGML_CUDA_GDN_GATHER_FUSION=0 and 1. * cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-per-row small-K GEMV (+14% TG) Merged onto the planar-transposed activation layout and dedicated 2-8 column kernel (#218, sudoingX). One decision, ggml_cuda_q8_1_layout_for(type, ncols_dst, ids), is evaluated by both the quantizer call and the kernel switch: PTQ1_0, 1 column, no ids -> SOA_ISUM (this change): exact int16 sums, warp-per-row small-K geometry PTQ1_0, 2-8 columns / MoE -> PT (#218): shared weight decode, full lane utilization anything else -> plain block_q8_1 GGML_CUDA_BATCH_INVARIANT=1 forces the one-column case onto PT as well so a token decoded alone and a token verified inside a speculative batch still see identical arithmetic. * cuda: PTQ1_0 hybrid dispatch: template y_soa, MMQ from 5 columns The SoA/planar choice is a kernel template parameter (only instantiated true for PTQ1_0 at one column) instead of a runtime branch inside the K loop; the runtime branch cost 5% TG on Ada. ggml_cuda_should_use_mmvq keeps PTQ1_0 on the mat-vec path up to 4 columns: with the branch-free MMQ tile loader the tile path overtakes the planar mat-vec kernel at n=5 (RTX 4070, Bonsai 2 27B: 163 vs 137 t/s at n=5, 258 vs 148 at n=8). * cuda: flash attention MMA reads q4_0/q8_0 K/V in place (no F16 scratch copy) The tensor-core flash attention path converted the whole quantized K/V cache to F16 in a scratch buffer on every micro-batch: 1 GB per context at 262k with q4_0 caches, 2 GB with a draft context, and a full re-conversion per ubatch during prefill. With Bonsai 2 27B + MTP that pushed the 4070 12 GB over the edge and WDDM spilled to system RAM (41 tok/s). flash_attn_ext_f16 now takes type_KV as a template parameter. For q4_0 and q8_0 the tile loader dequantizes blocks straight from the cache into the shared-memory half2 tiles (aligned 32-bit word reads, scale from the block header), the cp.async multi-stage pipeline is disabled for those types (cp.async cannot dequantize), and K/V strides are byte strides. The allocator no longer reserves the F16 copy when the native kernel handles the call. Dispatch is limited to D=128/256 heads with matching K/V types and word-aligned strides; anything else keeps the old conversion path. test-backend-ops: 48 new FLASH_ATTN_EXT cases for q4_0/q8_0 at D=128/256, GQA, batch 1-8, KV lengths incl. non-multiple of the tile, permuted layouts. All pass against the CPU reference; the full suite is 2984/2984. Bonsai 2 27B PTQ1_0 + MTP, RTX 4070 12 GB, 262144 context, ub 128: resident, 222 MiB shared (no spill), 86 tok/s mean over 3 prompts. * cuda: PTQ1_0 multi-column mat-vec: raw digits with exact activation sums, per-pair epilogue (+13% 3-col, +24% 4-col) Two changes to mul_mat_vec_ptq1_0_pt, found with Nsight Compute on an RTX 4070 (sm_89). 1. The trit decode returns the raw base-3 digits {0,1,2} and the dot product uses a mixed-sign dp4a (u8 x s8) against the activations, subtracting the exact int16 sum of the 32 quantized activations once per sub-block: sum((q-1)*a) = sum(q*a) - sum(a). The PT quantizer now stores that sum bit-cast in the ds.y half slot (the SoA layout already did). Two ALU ops fewer per four weights; the result is the same integer, so output is bit-identical. 2. The epilogue summed each (row, column) pair with one warp: a runtime integer division per pair, a lane-strided loop and a 5-step shuffle butterfly. Under the profiler that was 23% of issued instructions and ~50% of warp stall samples (FADD on SHFL scoreboard, MUFU.RCP/I2F/F2I from the division) while the CTA held its slot and moved no bytes. Now one thread per pair sums the K blocks sequentially from shared memory with four interleaved accumulators, over a padded (odd) row stride so the reads are bank-conflict-free. The order still depends only on the weight shape, so the batch-invariant guarantee holds. Profile of the 3-column gate|up GEMV (K=5120, N=17408) before -> after, ncu locked clocks: 2135 -> 1493 instructions per warp, 56.3 -> 43.5 us, DRAM 62% -> 83% of peak. llama-bench, +1500 MHz memory: pp3 150 -> 169 tok/s, pp4 172 -> 212. pp1 unchanged (SoA path). test-backend-ops -b CUDA0 -o MUL_MAT: 1283/1283. * qwen35: MTP graph publishes only the output rows of h_nextn when embeddings_nextn_masked llama_context::decode copies the first n_outputs rows of t_h_nextn when embeddings_nextn_masked is set, but the QWEN35 MTP builder assigned t_h_nextn before the inp_out_ids gather, so the tensor had n_tokens rows. Any MTP batch with more tokens than outputs (e.g. catch-up rows decoded together with a draft row) returned row 0's hidden state for the draft row and the next draft step ran on the wrong input. Never triggered before because every MTP decode had n_tokens == n_outputs or no outputs. Gather before publishing in the masked case, exactly as the trunk graph does at its last layer; the unmasked case is unchanged. * speculative: draft-mtp decodes the catch-up rows together with the first draft row process() used to run the draft over every row of the target's verify batch immediately, purely to fill the draft's memory, and partly on rows the target then rejected. That is one graph launch and one host round trip per step spent on nothing the next draft needs yet. Single-head MTP with its own memory now keeps those rows (tokens, positions, shifted h) and decodes only the accepted prefix at the start of draft(), in the same llama_decode as the first draft row. Positions are contiguous and the draft row attends to the catch-up rows in-batch, so the draft is unchanged; accept() trims the stash to the accepted prefix, and rows are decoded on their own when a batch does not lead straight into a draft (prompt fed in tiny ubatches, drafting stopped, rollback replay overwrites them). Shared-memory and chained-head drafters keep the eager path, as does LLAMA_MTP_EAGER_CATCHUP=1 for A/B. RTX 4070, Bonsai-2-27B PTQ1_0 + MTP head, 262k context, q4_0 K/V, greedy, --backend-sampling: 23.0 -> 22.4 ms/step, three-prompt mean 101.0 -> 103.3 tok/s with identical draft acceptance (0.86 / 0.46 / 0.65) and byte-identical output vs the draft head off. Draft-side launches per step 4 -> 3. * cuda: Hadamard transform quantizes its own output when every consumer is a PTQ1_0 mat-vec A Hadamard-rotated activation ([MUL signs, RESHAPE,] MUL_MAT with GGML_HINT_SRC0_IS_HADAMARD) whose every use is src1 of a PTQ1_0 mat-vec on the mmvq path never needs to exist in F32: each consumer immediately quantizes it to q8_1. The transform kernel now quantizes as it goes and writes the q8_1 rows, in the layout the consumers expect (AoS / PT / SoA with exact integer sums and the same padding rule as ggml_cuda_mul_mat_vec_q), into its own output buffer; the mat-vecs find the rows through a per-context registry keyed by tensor identity and skip their quantize launch. On Bonsai 2 this removes ~390 launches per decode step. Safety rules the pass enforces: - every later node that reads the transform output (directly or via a whole-tensor reshape view) must be a PTQ1_0 MUL_MAT that ggml_cuda_mul_mat routes to mmvq, and the graph use count must match, so nothing ever reads the buffer as F32 - the registry is keyed by ggml_tensor pointer, never by data pointer: ggml-alloc recycles a dead output's block for later tensors in the same graph - when the output buffer aliases the input (ggml-alloc does this for the output-projection input on every Bonsai layer: x is a reshape view of a cont, there is no in-place reuse through a view parent, so the cont is released at the MUL and mm takes its block two nodes later) the rows go to a pool block held until the next graph evaluation. Writing into the aliased buffer raced the kernel's own input reads and produced sporadic garbage that, on a hybrid model, entered the recurrent state and poisoned the rest of the sequence (surfaced as INT_MAX from the backend argmax over an all-NaN logits row, about once per ~170 requests) - GGML_CUDA_FWHT_FUSION=0 disables the pass fwht.cuh: the block butterfly is a shared inline so fwht.cu and the fused quantizer stay one implementation. Validated: 340 consecutive requests (probe.py x 20 at 262144 ctx, q4_0 KV, draft-mtp n_max=2, backend sampling) with zero errors; output byte-identical to the unfused path. * sampling: reject out-of-vocab token ids from the backend sampler and the draft A backend-sampled id outside [0, n_vocab) (the CUDA argmax/top-k pad sentinel INT_MAX when a logits row is all NaN) used to travel as-is into the server, where the tokenizer's vector::at() threw "invalid vector subscript" and killed the request. Log the origin (idx, n_vocab, first candidate) and fall back to the CPU chain instead; the draft-mtp loop drops a draft whose top candidate id is out of vocab rather than feeding it to the verify decode. * speculative: draft-mtp drops stale deferred rows when a new task starts on the slot begin() used to decode whatever catch-up rows the previous task left pending. That is right only when the new prompt reuses them (same tokens at the same positions, continuing what ctx_dft holds). After a cancelled task or a different conversation landing on the slot, the server has already trimmed ctx_dft past those rows and decoding them writes stale positions into the recurrent draft state. Check first, flush when they belong, discard otherwise. * cuda: release the FWHT-q8 held pool blocks LIFO and before the pools are destroyed The registry keeps pool blocks alive across a graph when the fused Hadamard quantize would otherwise write over its own input. Its implicit destructor released them oldest first, and it is declared before the pools, so at context teardown the frees hit a destroyed VMM pool: GGML_ASSERT(ptr == pool_addr + pool_used) at the end of llama-bench. Release newest first in a destructor, and reset() from the context destructor ahead of the pools. * server: --spec-draft-depth-max stops drafting once the sequence is deep Deep in the context a decode step is bound by reading the KV cache. The draft passes and the multi-column verify add to that read without shortening it, so past some depth speculation costs more than the accepted tokens save: with Bonsai 2 27B on an RTX 4070 the MTP draft is +85% at zero depth, even near 24-32k tokens and -27% at 64k. Past the cutoff the slot decodes one token per step and the process() hook is skipped for batches that sit entirely beyond it (by position, so the early ubatches of a long prompt still reach the draft context and remain reusable as a prefix). Default 0 keeps the old behaviour. 0 / 32k / 64k tokens: 101 / 48 / 37 tok/s with a 24k cutoff against 101 / 45 / 27 without. * cuda: restrict off PTQ1_0 kernel params; Ampere uses the PT 1-col path. C2912 / stub mismatch on sm_90 and sm_120 when GGML_CUDA_RESTRICT is on the formal parameters of mul_mat_vec_ptq1_0_pt (same hunk as 2578fdf on #218). Ampere (cc 800-889) takes the planar 1-col kernel; Ada and newer keep SoA. 3060: #218 was +5.9% tg128 vs #215 at one column. * cuda: address Copilot review on native q4/q8 FA, MTP catch-up, and BOM. Disable drafting when deferred catch-up decode fails. Gate in-place quantized K/V FA to NVIDIA and require ggml_cuda_is_aligned on K and V. Strip UTF-8 BOMs. The FWHT fusion comment no longer opens with the GDN gather notes. * cuda: share the PTQ1_0 mat-vec smem budget between guard and launch The combo kernel pads each partials row by one float. Use the same byte count in the entry guard as the launch, including that pad, so oversized gated shapes fall back instead of requesting more than smpb. * spec: flush MTP catch-up when it plus anchors would overflow n_batch process() can stash n_batch rows; draft() then added one more anchor. Flush first when the combined count would exceed llama_n_batch. * cuda: restore the scalar PTQ1_0 vec-dot on MUSA Same hole as #215: the generic entry returned zero once MUSA was kept off the SoA path. MUSA now takes the HIP scalar implementation. * spec: keep MTP catch-up across a failed draft decode catchup_failed was set on a flush failure and never cleared, so one bad decode disabled that slot for every later task. begin() now resets it with the rest of the per-task state. draft() also cleared pending while building the batch, so a failed llama_decode dropped those rows from both the stash and ctx_dft. pending stays set until that first decode succeeds, and a failure marks the slot so it is not drafted again until the next begin(). flush_deferred waits for success the same way. The layout_for comment now says it is the column-count helper. ggml_cuda_q8_1_layout_host is the decision both sides call. * cuda: define the PTQ1_0 layout helper after ggml_cuda_info The host layout decision reads the current device's compute capability. It was declared above ggml_cuda_info and ggml_cuda_get_device, so a clean sm_89 build failed in every CUDA translation unit. The helper now sits below those declarations. * cuda: land remaining #221 review notes: pass q8 layout down, skip HIP/MUSA q4/q8 FA instances, exercise ALiBi/softcap. * cuda: wait on the PDL dependency in the PTQ1_0 planar mat-vec (Hopper/Blackwell). * cuda: restore warp-reduce epilogue under BATCH_INVARIANT so MTP-on matches MTP-off. The four-accumulator path (5300cd1) is faster but a 5080 run flipped a late near-tie. GDN now waits on PDL before reading s_ids. --------- Co-authored-by: sudoingX <200180104+sudoingX@users.noreply.github.com> Co-authored-by: Cary Palmer <24235924+professorpalmer@users.noreply.github.com>
Branch:
pr-ptq1-mmv, five commits on top ofprism(9a9394a89), 8 files, 596 insertions, 21 deletions.Measured on one RTX 3060 12GB (sm_86, CUDA 12.4, driver 550.144.03) with Ternary Bonsai 2 27B.
What
Add: planar-transposed activation layout for the PTQ1_0 mat-vec path(quantize.cu,mmvq.cu,new
mmvq-ptq1_0.cuh). PTQ1_0 activations are quantized into a layout where the 128 quants athread needs are 8 aligned 16-byte pieces plus one 16-byte piece of scales, adjacent lanes on
adjacent pieces; same bytes per row and the same column stride as
block_q8_1, so nothing in thelaunchers changes. Trit decode unchanged; the byte-wise
-1is(q + 0x7F7F7F7F) ^ 0x80808080.nwarps pinned to 4 for all column counts; mmvq handles PTQ1_0 for 1 to 4 columns, 5 and above take the
MMQ tile path now that cuda: branch-free PTQ1_0 MMQ tile loader and full Ampere tile table (2x prefill) #214 is in the tree (see the last commit). HIP unchanged.
Add: dedicated PTQ1_0 mat-vec kernel with full lane utilization(mmvq-ptq1_0.cuh,mul_mat_vec_ptq1_0_pt). For plain 2D MUL_MAT: (row group, K block) work items flattened soevery thread has work, 4 rows per thread per K block, one fp32 partial per (row, column, K block)
in dynamic shared memory, one warp per (row, column) reducing in a fixed order. The per-block
partial is
d * sum_k d8_k * sumi_kwith exact integersumi_k(__fmaf_rn,__fmul_rn), theorder depends only on the weight shape: a column's result is bit-identical for every column count
1 to 4. Fusion for one column as in the generic kernel; batched and MoE calls keep the generic
kernel (which also reads the new layout).
Test: add Bonsai 2 projection shapes to the mul_mat perf cases(tests/test-backend-ops.cpp):PTQ1_0 and Q4_0 at the six Bonsai 2 projection shapes for 1, 2, 3, 4, 8 columns (8 now exercises the
MMQ tile path), plus the bf16 [5120 x 48] gate projection.
Add: GGML_CUDA_BATCH_INVARIANT for batch-invariant small-batch kernels(common.cuh,mmvf.cu,fattn.cu,fattn-common.cuh). Off by default. When set: small F16/BF16 matrices takemul_mat_vec_ffor 1 to 8 columns, flash attention with up to 8 queries takes the vector kernel,and its KV split is sized as for one query tile.
Fix: use the mat-vec kernel for bf16 matrices under 64 rows at 2 to 8 columns(mmvf.cu).Fix: keep GGML_CUDA_RESTRICT off the PTQ1_0 mat-vec kernel signature(mmvq-ptq1_0.cuh): thehost stub for sm_90 and sm_120 rejects restrict-qualified formal parameters; aliases inside the body.
Fix: budget the PTQ1_0 mat-vec shared memory the launch will request(mmvq-ptq1_0.cuh,tests/test-backend-ops.cpp): the entry guard and the launcher shareptq1_0_pt_smem_bytes(), foureval cases at the 48 KiB boundary with and without gate fusion.
Fix: cap the PTQ1_0 mat-vec at 4 columns and send 5 and above to the MMQ tile path(mmvq.cu,mmvq-ptq1_0.cuh,common.cuh): with cuda: branch-free PTQ1_0 MMQ tile loader and full Ampere tile table (2x prefill) #214's branch-free tile loader inprism, MMQ does a batch of 8in 64.4 ms where the 8-column mat-vec took 112.7 ms (RTX 3060, K = 5120 shapes, llama-bench pp8, r=3;
github.com/sudoingX/bonsai2-small-gpu/blob/main/kernel/ab/ab_215_214_rtx3060.md), so
ggml_cuda_should_use_mmvqroutes PTQ1_0 to the mat-vec for 1 to 4 columns only.Two
Docs:commits scope the batch-invariance statement to 1 to 4 columns and to the paths the flag covers.Why
The fork's PTQ1_0 mat-vec charged 1.5 single-token passes for a 2-token batch and 2.3 for a
3-token batch, so speculative decoding (
--spec-type draft-mtp) could not pay for its drafts on aternary model, and its logits changed with the batch size, so greedy output with the draft head
differed from greedy output without it. Two causes, both in the same kernel: the activations were
read as 32 scattered 4-byte loads per column out of 36-byte structs, and on the K = 5120 projections
(three quarters of the 27B's weights) 88 of every 128 threads sat idle because one thread owned a
whole 128-element block.
test-backend-ops perfon a 4096 x 14336 PTQ1_0 matrix: 47 / 104 / 154 /201 us for 1 / 2 / 3 / 4 columns; Q4_0 on the same shape 101 / 102 / 115 / 151.
Before / after
llama-bench,Ternary-Bonsai-2-27B-PTQ1_0with the MTP head appended, 4096 context,-fa 1 -ctk q4_0 -ctv q4_0, 8 repetitions:GGML_CUDA_BATCH_INVARIANT=1measures the same within noise (tg32 39.70).Kernel level,
test-backend-ops perf, us per call (before = the fork's kernel on the same layoutchange, so this isolates the dedicated kernel; the layout change alone took the 4096 x 14336 case
from 47 / 104 / 154 / 201 to 46 / 55 / 65 / 75 us):
Speculative decoding end to end (
llama-server, 131072 context, one slot, q4_0 K/V, thinking off,client-side tok/s over streamed tokens, medians of 3 runs on 3 prompts):
draft-mtpn-max 1,GGML_CUDA_BATCH_INVARIANT=1GGML_CUDA_BATCH_INVARIANT=1At 41.8K tokens of context: 22.2 tok/s flag off, 26.9 with n-max 1 (invariant mode), 28.3 with
n-max 2 on the default kernels.
Correctness
test-backend-ops test -o MUL_MAT -p type_a=ptq1_0: 47 of 47 against the CPU reference, including1 to 9 columns and batched shapes (generic kernel).
-o MUL_MAT_ID: 77 of 77. The flag-off greedytext changes versus the prebuilt build at near ties (the K summation order changed), as expected for
any kernel change; the batched-prefill check that exposed the prebuilt fork's own decode and prefill
paths disagreeing by up to 0.1 nats now shows bit-identical logprobs for 1 to 4 tokens.
How to reproduce
The invariance check (same prefix, last N tokens in one batch, compare the next-token logprobs) is
tools/batch_numerics.pyin github.com/sudoingX/bonsai2-small-gpu against a runningllama-server.Known limits
1.55x of pp1). Removing the activation loads from the kernel (perf only) makes 1 to 8 columns cost
the same, so the loads are the cost; broadcasting them across 8 lanes (warp tiles), more rows per
thread, tighter launch bounds and two
mma.sync.m16n8k16.s8variants (correct, bit-identical byconstruction) did not beat the dp4a kernel at 2 to 4 columns on this card. Details and diffs of
the attempts are in the repo's
KERNEL_REPORT.mdandresults/kernel/experiments/.GGML_CUDA_BATCH_INVARIANT=1makes batches of 1 to 4 bit-identical. Batches of 5 and above take theMMQ tile path (since the column cap at 4) and are not covered by the flag; before the cap, 5 to 8
agreed with each other but differed from 1 to 4 by up to 0.03 nats on the tested steps (top-1
unchanged), and the kernel that switched between 4 and 5 columns was not identified. Prompt processing
is not covered either.
n-max 2 loses at depth (21.1 vs 22.2 tok/s flag off) where the default kernels gain (28.3).
it was not benchmarked (no ternary MoE file at hand).
ptq1_0_pt_enabled()is false there); the new code iscompiled out with
!defined(GGML_USE_HIP).