cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-per-row small-K GEMV (+10.5% TG, +14.7% with #220) - #215
Conversation
|
Same wall, reached on an RTX 3060 (sm_86) with Details and a proposed split are in a comment on #218. Short version: your exact int16 sums and the GDN fusion are things #218 does not have, the multi-column path and the invariance switch are things this PR does not have, and the two can be rebased onto one layout, yours if the maintainers prefer it. Independent of that, the wide MMQ tiles and the branch-free loader in #214 are welcome on this card too: PTQ1_0 pp512 on the 3060 sits at 268 t/s. Will post 3060 numbers for this branch when the card is free. |
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.
Request changes: P1 - Own the gather registry per CUDA context/evaluation.
ggml/src/ggml-cuda/gated_delta_net.cu:8-18 introduces one mutable static unordered_map shared by every CUDA context and host thread. Each graph evaluation clears it. Two independently decoded contexts can interleave A's registration/GET_ROWS skip, B's clear, then A's GDN lookup. A then reads the gathered tensor whose producer it skipped. Concurrent find/insert/clear also creates a C++ data race, including when the contexts use different GPUs.
Please store the registrations in the owning backend/evaluation and pass that state to the consumer. Locking individual map operations alone would not prevent another graph clearing A's registrations.
Confirmed by source tracing; GPU concurrency was not reproduced here. Separately, coordinate the activation layout and MMVQ changes with #218 before integration.
Reviewed commit: b7660800c2460d8b87c4375222829181012727a9.
… (+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.
b766080 to
1cd8ffc
Compare
professorpalmer
changed the title
|
Revision 2 pushed ( P1, gather registry: the GDN gather fusion is no longer in this PR. It is now #220 on its own, with the registrations owned by the evaluating Coordination with #218: I built
On Ada the two layouts split cleanly: this one is ahead at one column, #218's dedicated multi-column kernel is far ahead at 2-4 columns, and this PR barely moves there. That matches your 4090 result for #218 (tg128 -0.1%) next to its 3090 result (+26.9%). I have no Ampere card; @sudoingX has offered 3060 numbers for this branch, which would complete the picture. I agree with the split @sudoingX proposed: exact int16 sums and warp-per-row from here, multi-column kernel, invariance switch and test shapes from there, on one layout. Either layout can host the other's pieces. If the maintainers state a preference I will do the rebase in that direction; if you would rather see one combined PR from the two of us, also fine. Full body updated above with the table and raw run JSON. |
|
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 |
|
RTX 3060 12GB numbers for this branch, as promised. Same PTQ1_0 file,
Three reads:
All arms sat on the card's 150 W cap (medians 142 to 149 W, SM clock 1665 to 1837 MHz). Verbatim tables, the pp1 to pp8 ms ladder and the power samples: https://github.com/sudoingX/bonsai2-small-gpu/blob/main/kernel/ab/ab_215_214_rtx3060.md |
|
This is the 3060 table we needed. Thank you. Ampere ranking is the opposite of Ada, agreed: #218 wins at one column here (+5.9% over #215), Ada was the other way. #221's hybrid is currently "SoA at 1 column, PT at 2–4" because that was the 4070 win. On sm_86 it should prefer the #218 kernel at one column too — per-arch, not only per column count. I'll flip that in #221 unless you want it in this PR. 40.4 tg128 on a 150 W 3060 is not a weak result. Stock on that card is 26.2; #218 is +54%. The 4070 headline (67.6 kernel / 97 served) is a different measurement on a different part: 504 GB/s, ~2.7 GHz, 200 W, and the served number includes the MTP head + deferred catch-up, which none of these five arms have. Same stack on this 3060 should look like ~40 llama-bench and then whatever the draft accepts on top, not 97. #214 doubling pp512 (268 → 515) and beating the 8-column mat-vec at pp8 matches the 4070. Dropping #218's column limit from 8 to 4 once #214 is in is the right cut. If you get a free window, |
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.
|
Landed on the integration branch ( |
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
MUSA currently enables behavior that the documented scope says is disabled, and several comments violate repository conventions.
Get a fresh assessment by requesting another Copilot review.
Review effort: Balanced
Findings: 1
| static constexpr __host__ __device__ bool ggml_cuda_q8_1_exact_isum(ggml_type type_src0) { | ||
| #if defined(GGML_USE_HIP) | ||
| GGML_UNUSED(type_src0); | ||
| return false; | ||
| #else | ||
| return type_src0 == GGML_TYPE_PTQ1_0; |
There was a problem hiding this comment.
You're right: MUSA took the non-HIP branch, so exact-isum and the SoA vec-dot would have gone live. 1b14d26 gates ggml_cuda_q8_1_exact_isum and those vec-dot paths on !HIP && !MUSA. We do not have a MUSA box to re-validate; the old quantizer/vec-dot is what that backend compiled before this PR.
| // For these src0 types the row-quantizer stores the exact integer sum of the q8 values (int16 | ||
| // bits) in block_q8_1::ds.y instead of the float input sum, so the ternary vec-dot can accumulate | ||
| // raw digits {0,1,2} with dp4a and subtract the bias once per 32-block instead of per 4 weights. | ||
| // The quantizer also writes the warp-transposed layout below for these types. |
| // Warp-per-row geometry for small K. | ||
| // The generic small_k loop strides K-blocks across the whole block (kbx = tid; kbx += 128), so | ||
| // for K=5120 (40 blocks) only threads 0..39 ever load anything: warp 0 is full, warp 1 has | ||
| // 8 lanes, warps 2..3 only wait at __syncthreads while holding SM warp slots. Measured | ||
| // in-graph with CUPTI on an RTX 4070: 350-390 GB/s for every K=5120 GEMV vs 460 GB/s for the | ||
| // K=17408 geometry where all four warps stream. Here warp w owns row row0+w and its 32 lanes | ||
| // stride that row's K-blocks: every warp issues loads, the reduction is a shuffle, and | ||
| // there is no shared memory or block barrier. Activations are re-read per warp but they | ||
| // are 20 KB and L1/L2 resident. |
There was a problem hiding this comment.
Trimmed to the launch invariant in 1b14d26. Numbers stay in the PR body.
| // exact_isum: store the integer sum of the quantized values (bit-cast into the ds.y half slot) | ||
| // instead of the float sum of the inputs. Ternary vec-dots (PTQ1_0) use it to fold the | ||
| // digit bias {0,1,2} -> {-1,0,+1} into one subtraction per 32-block instead of a SIMD byte | ||
| // subtract per 4 weights, with results bit-identical to the biased path. |
| // Lane-coalesced activation words: word ww of this K-block is at yw[ww*32]; 32 lanes of a | ||
| // warp own 32 consecutive K-blocks, so every load below is 32 consecutive words. |
|
Agent benchmark follow-up, posted at the maintainer's request. Historical measurement: tested head
Selected CPU-reference backend checks passed on the compared arms. The current head 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):
|
|
Thanks for pinning the SHA caveat. I agree with it, and I want it on the record in the same words. What you measured. What the 4070 can still corroborate, because we already split the two pieces on this card (stock clocks, same flags, three alternating rounds;
pp512 on that box is 600 → 1256 only because #214 is in the same binary; it is not a #215 result (your campaign already shows pp512 flat, which is the expected MMVQ-only shape). I am not quoting our pp8 here for the same reason: once #214 is present, batch 8 takes the MMQ tile path. Shape vs your four cards. Ada is the small-percent / high-absolute lane: 4090 +6.9% (fused historical) sits next to 4070 +10.5% (layout only) and +14.7% (layout + fusion). Ampere is the large-percent lane: 3090 +34.1% fused historical, and @sudoingX's 3060 at current-style Your exploratory pp2/pp4 (4090 +4.9% / +15.2%) is the same direction as 4070 +11% / +26%. #218 still owns the 2-4 column win; this PR's job at those widths is not to compete with that kernel. Remaining review items on |
bri-prism
left a comment
bri-prism
left a comment
There was a problem hiding this comment.
Re-reviewed at head 1b14d26. The gather registry point is resolved, but the MUSA scoping commit introduced a new problem that I think blocks merge.
P1 is resolved. The file-static gather registry is gone from this PR entirely, and the fusion now lives in #220 against a per-context registry, which is the ownership model I asked for. I have not audited #220, so that is resolved for this PR rather than resolved outright.
New: PTQ1_0 mat-vec returns zero on MUSA.
vec_dot_ptq1_0_q8_1 in ggml/src/ggml-cuda/vecdotq.cuh keeps its real implementation under #if defined(GGML_USE_HIP), and the #else branch now returns 0.0f after GGML_UNUSED on all four arguments. MUSA takes that #else.
At the merge base 9a9394a the same #else branch called vec_dot_ptq1_0_q8_1_multi<1>(vbq, bq8_1, kbx, iqs, 0, &result) and returned a real value, so MUSA had a working path. Two changes in this PR closed it:
vecdotq.cuh:807now gates_multion#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA), so MUSA no longer has_multiat all.- The PTQ1_0 fast paths in
mmvq.cuat :670, :807 and :992 carry the same gate, so on MUSA they all compile out and PTQ1_0 falls through to the generic path.
The generic path resolves PTQ1_0 through mmvq.cu:16, case GGML_TYPE_PTQ1_0: return vec_dot_ptq1_0_q8_1, which is the stub. So every PTQ1_0 mat-vec on MUSA produces zeros. The GGML_UNUSED calls silence the warnings that would otherwise flag an unused-argument stub, so this builds clean and fails silently at runtime.
Suggested fix, one line at the top of that function:
#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA)
That routes MUSA back to the scalar implementation, which matches what the accompanying comment in common.cuh already claims is happening. The comment says HIP and MUSA keep the old quantizer and vec-dot. That is accurate for the quantizer, since ggml_cuda_q8_1_exact_isum returns false under MUSA, but not for the vec-dot.
Two smaller things, neither blocking:
quantize.cu:92 computes float sum = xi; sum = warp_reduce_sum<QK8_1>(sum); inside the exact_isum branch, which then computes and uses its own integer isum. The float reduction looks dead there.
I traced the SoA addressing and did not find a problem with it. ne10_padded is padded to GGML_CUDA_PTQ1_K_PAD only when exact_isum is true, col_base matches the vec-dot's stride, the group-word arithmetic is consistent, and tail padding is handled. The single quantize_row_q8_1_cuda call site means the MoE path sees the same layout. No issue found.
Method note: this is source tracing against 1b14d26, not a build or a run. I have no MUSA hardware either, so if you read the chain differently I would rather be corrected than have you patch around a misreading.
On sequencing, I still think the activation layout needs to converge with #218 before either lands. I can see #221 exists as the integration branch; I have not reviewed it.
Same hole as PrismML-Eng#215: the generic entry returned zero once MUSA was kept off the SoA path. MUSA now takes the HIP scalar implementation.
|
You read the chain correctly. The MUSA gate compiled out Pushed The float Same stub was on the integration branch. |
bri-prism
left a comment
bri-prism
left a comment
There was a problem hiding this comment.
Verified at ad61b84. The MUSA regression is fixed and I traced the whole path rather than taking the commit message for it.
vec_dot_ptq1_0_q8_1 now opens with #if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA), so MUSA reaches the scalar implementation again instead of the stub.
I also checked the claim the new comment makes, that the generic entry is unused on NVIDIA, because that is what makes the remaining return 0.0f safe rather than merely unreached by luck. It holds. In mmvq.cu the if constexpr (type == GGML_TYPE_PTQ1_0) block routes to vec_dot_ptq1_0_q8_1_multi, and the generic vec_dot_q_cuda call sits in its else. The whole construct is inside #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA), so NVIDIA always takes the _multi branch and HIP and MUSA compile down to the generic block alone, which now lands on the scalar path. The same shape holds at the small-K path and in the MoE variant.
Clearing my change request. Nine minutes from report to correct fix was a good turnaround.
Two things carried forward, neither blocking this approval.
The activation layout still needs to converge with #218 before either lands. I can see #221 exists as the integration branch and that this PR is deliberately standalone for review, which is a reasonable split. I have not reviewed #221, so I am approving this PR on its own contents and not the integration.
Separately, PR #211 rewrites the HIP branch of this same function. Neither branch contains the other and the two versions differ by roughly fifty lines in one function, so they will conflict. The resolution is now less dangerous than it was, since the MUSA stub is gone from your side, but whoever merges second should confirm the #if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) condition survives intact. It is the kind of line that a conflict resolution quietly reverts.
Method note: source reading at ad61b84. I did not build, and I have no MUSA or AMD hardware.
… (+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.
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.
Same hole as PrismML-Eng#215: the generic entry returned zero once MUSA was kept off the SoA path. MUSA now takes the HIP scalar implementation.
… (+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.
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.
Same hole as PrismML-Eng#215: the generic entry returned zero once MUSA was kept off the SoA path. MUSA now takes the HIP scalar implementation.
|
Kernel-level numbers from an RTX 4090 (sm_89, CUDA 12.8), comparing this PR's head (ad61b84) with #218's head (285542d) and
With the weights larger than L2, both PRs close the small-K gap about equally. The difference is on the MoE expert shapes: with this PR, PTQ1_0 runs at PQ2_0 speed, while on The branch conflicts with |
…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.
MUSA took the non-HIP branch and would have used the new quantizer and vec-dot. Also trim the hard-wrapped benchmark comments.
The MUSA gate compiled out the SoA fast path and left the generic entry returning zero. MUSA now takes the same scalar implementation as HIP. The exact-isum quantizer no longer reduces an unused float sum.
ad61b84 to
68ea9b4
Compare
|
Rebased onto current The only conflict was
Happy to have the 4090 |
… (+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.
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.
Same hole as PrismML-Eng#215: the generic entry returned zero once MUSA was kept off the SoA path. MUSA now takes the HIP scalar implementation.
… 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>
|
Landed as part of #221 ( What is in @bri-prism, on the line you asked to survive the #211 merge: it did. |
Summary
Two changes to the batch-1 PTQ1_0 mat-vec path, motivated by CUPTI traces of a Bonsai 2 27B decode graph on an RTX 4070 (sm_89, 504 GB/s). Same weights, same bits, byte-identical output. Independent of #214 (prefill).
Revision 2: the GDN gather fusion that was part 3 of the first revision is now its own PR (#220) with the per-context registry requested in review, so this PR is only the activation layout and the small-K GEMV geometry and can be compared with #218 on its own. Nothing in this PR touches
gated_delta_net.*orggml-cuda.cu.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 inds.y(instead of the float input sum) and writes the activations grouped by 32 K-blocks with word w of the group stored contiguously (ggml_cuda_ptq1_q8_word). Two effects:__vsub4(no native SIMD byte subtract on sm_8x) from every 4-weight round. Throughput was flat from this alone but board power dropped 13 W at the cap, which the memory system then used.Only PTQ1_0 opts in via
ggml_cuda_q8_1_exact_isum(); every other type keeps the existing layout and the existingvec_dot_q_cudapath. HIP/MUSA compile with the feature off.2. Warp-per-row geometry for small-K PTQ1_0 GEMVs (
mmvq.cu)The generic
small_kloop strides K-blocks across the whole 128-thread block, so for K=5120 (40 blocks) only threads 0..39 ever load anything: warp 0 is full, warp 1 has 8 lanes, warps 2-3 wait at__syncthreadsholding SM slots. Measured in-graph: 350-390 GB/s for every K=5120 GEMV vs 460 GB/s for the K=17408 geometry where all four warps stream. Each warp now owns one row and its 32 lanes stride that row's K-blocks: shuffle reduction, no shared memory, no barrier. The 20 KB activation vector is re-read per warp but is L1/L2 resident. qkv 391 -> 445 GB/s, lm_head 414 -> 460.Measurements
RTX 4070 12 GB,
Ternary-Bonsai-2-27B-PTQ1_0.gguf, CUDA 13.3,llama-bench -ngl 99 -fa 1 -ctk q4_0 -ctv q4_0 -r 3, stock clocks (SM ~2760-2790, GDDR6X 10251 MHz, all arms at the 200 W cap), three alternating rounds per arm, means. Base is the Prism release binary at 9a9394a. The #218 column issudoingX/llama.cpp@pr-ptq1-mmvbuilt from source with the same toolchain, for the comparison the review asked for.So on Ada the two layouts split cleanly: this one is ahead at one column (batch-1 decode, which is the case that matters without a draft head), #218's dedicated multi-column kernel is far ahead at 2-4 columns (speculative verification), and this PR barely moves there because its small-K path still charges one pass of activation loads per column. That is consistent with the maintainer's 4090 numbers for #218 (-0.1% tg128) and its 3090 numbers (+26.9%): the win from that layout is on Ampere, the win from this one shows on Ada. I have not been able to run this branch on an Ampere card; sudoingX has offered 3060 numbers.
I agree with the split proposed on #218: the exact int16 sums and the warp-per-row geometry from here, the multi-column kernel, invariance switch and test shapes from there, on one layout. Either layout can carry the other's pieces; this one already stores exact per-block integer sums, which #218 would need to add for the bias-once-per-block decode. Happy to do the rebase in whichever direction the maintainers prefer.
Raw runs:
artifacts/h2h_215_vs_218_4070_stockclocks.jsonin https://github.com/professorpalmer/bonsai-ada-surgery.Correctness
test-backend-ops -b CUDA0: MUL_MATtype_a=ptq1_045/45 supported cases, MUL_MAT_IDtype_a=ptq1_075/75 (the MoEmul_mat_vec_q_moepath is updated for the new layout too).llama-perplexity -c 2048 -b 2048: 7.6740 vs 7.6742 (float noise).Scope
CUDA only. The SoA layout and warp-per-row path are inside
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA);ggml_cuda_q8_1_exact_isumreturns false on HIP so the quantizer and vec-dot keep their existing behaviour there. Not tested on Hopper/Blackwell/GB10; the changes use no arch-specific intrinsics beyond what the file already uses.Profiler, traces and the notes on what did not help (LUT unpack, L2 prefetch/persistence, nwarps, forced MMQ, PDL off, write-through state stores): https://github.com/professorpalmer/bonsai-ada-surgery