diff --git a/README.md b/README.md
index 643eb24..a612657 100644
--- a/README.md
+++ b/README.md
@@ -23,38 +23,46 @@ matching chat template).
| Model | Repo | Size | Backend | Prefill t/s | Decode t/s | Notes |
|---|---|---:|---|---:|---:|---|
-| SmolLM2 1.7B Instruct | [HuggingFaceTB](https://huggingface.co/HuggingFaceTB/SmolLM2-1.7B-Instruct-GGUF) | 1 GB | CPU | 16.6 | 38.9 | AVX2 fused dequant-matvec |
-| SmolLM2 1.7B Instruct | (same) | 1 GB | Vulkan `-g -1` | 42.0 | **139.7** | GLSL `subgroupAdd` reduce |
-| SmolLM2 1.7B Instruct | (same) | 1 GB | **CUDA** `-g -1` | **181.1** | **158.1** | NVRTC `__dp4a` + Q8_1 |
-| Qwen3 8B | [Qwen](https://huggingface.co/Qwen/Qwen3-8B-GGUF) | 5 GB | CPU | 2.9 | 11.7 | dense forward pass, no KV compression |
-| Qwen3 8B | (same) | 5 GB | CPU `--tq` | 3.2 | **11.9** | 3-bit KV → up to 40 960 ctx. FastScan K + V kernels (issue #34) deliver ~20× over the per-block AVX2 path on the attention K+V hot loop; visible end-to-end at long context: **10.2 t/s @ 3K, 9.4 t/s @ 6K** (vs ~5 t/s baseline at 6K). Batched prefill is TQ-aware (issue #34 follow-up) so prompt processing matches the no-`--tq` rate |
-| Qwen3 8B | (same) | 5 GB | Vulkan `-g -1` | 23.0 | 45.8 | 11.4K auto-ctx |
-| Qwen3 8B | (same) | 5 GB | Vulkan `-g -1 --tq` | 21.7 | 45.5 | 3-bit KV → 40 960 ctx |
-| Qwen3 8B | (same) | 5 GB | **CUDA** `-g -1` | **65.9** | **58.6** | ~2.8× Vulkan prefill |
-| Qwen3 8B | (same) | 5 GB | **CUDA** `-g -1 --no-thinking` | **66.0** | **58.2** | Same per-token rate; reasoning suppressed in chat template, so all decoded tokens are visible answer |
-| Qwen3 8B | (same) | 5 GB | **CUDA** `-g -1 --tq` | **65.9** | **58.4** | 3-bit KV → 40 960 ctx; 17 t/s @ 8K, 10 t/s @ 16K |
-| Qwen3 8B | (same) | 5 GB | **CUDA** `-g -1 --tq --no-thinking` | **66.1** | **58.1** | Same per-token rate as `--tq` alone; reasoning suppressed |
-| OLMoE 1B-7B Instruct (MoE) | [allenai](https://huggingface.co/allenai/OLMoE-1B-7B-0924-Instruct-GGUF) | 4 GB | CPU | 21.6 | 55.7 | 64 experts / 8 active; per-channel QK-norm; `norm_topk_prob=false` |
-| OLMoE 1B-7B Instruct (MoE) | (same) | 4 GB | Vulkan `-g -1` | 18.9 | **121.2** | 16 layers all on VRAM; greedy on this prompt is unstable across backends — use `--temp 0.6 --top-p 0.95` for usable output |
-| OLMoE 1B-7B Instruct (MoE) | (same) | 4 GB | **CUDA** `-g -1` | **117.4** | **111.7** | Same; greedy varies, sampling coherent |
-| Qwen3-Coder 30B-A3B (MoE) | [Qwen](https://huggingface.co/Qwen/Qwen3-Coder-30B-A3B-Instruct-GGUF) | 17 GB | CPU | 13.3 | 21.1 | 128 experts / 8 active |
-| Qwen3-Coder 30B-A3B (MoE) | (same) | 17 GB | CPU `--tq` | 13.7 | 21.0 | 3-bit KV. FastScan K + V kernels (issue #34) keep attention cost bounded as context grows: **15.5 t/s decode @ 3.2K ctx** (27 % slowdown for ~27× context growth); without FastScan the per-block K+V path would drop this to ~13 t/s |
-| Qwen3-Coder 30B-A3B (MoE) | (same) | 17 GB | Vulkan `-g -1` (hybrid) | 1.1 | 5.3 | 29 GPU + 19 CPU layers, SLRU expert slot cache. Next-layer predictive prefetch (PR #77 / issue #50) on by default; no-op until the cache is under pressure — disable with `--no-moe-predict-prefetch` |
-| Qwen3-Coder 30B-A3B (MoE) | (same) | 17 GB | **CUDA** `-g -1` (hybrid) | **12.3** | **25.0** | 29 GPU + 19 CPU layers; routed experts stream through `CudaExpertSlotManager` SLRU (2220 / 3712 slots) instead of the prior eager whole-layer upload (PR #77 / issue #72). Decode lifts 22.2 → 25.0 (+13 %) and prefill 10.6 → 12.3 (+16 %) on top of #77's baseline once the warm-pin + frequency-biased eviction (#74) and async upload stream (#78) hide miss latency in the steady-state working set. Set `SHARPI_EXPERT_STATS=path` to inspect per-layer hit rates |
+| SmolLM2 1.7B Instruct | [HuggingFaceTB](https://huggingface.co/HuggingFaceTB/SmolLM2-1.7B-Instruct-GGUF) | 1 GB | CPU | 40.4 | 38.9 | AVX2 fused dequant-matvec |
+| SmolLM2 1.7B Instruct | (same) | 1 GB | Vulkan `-g -1` | 123.2 | **139.7** | GLSL `subgroupAdd` reduce |
+| SmolLM2 1.7B Instruct | (same) | 1 GB | **CUDA** `-g -1` | **163.1** | **158.1** | NVRTC `__dp4a` + Q8_1 |
+| Qwen3 8B | [Qwen](https://huggingface.co/Qwen/Qwen3-8B-GGUF) | 5 GB | CPU | 9.9 | 11.7 | dense forward pass, no KV compression |
+| Qwen3 8B | (same) | 5 GB | CPU `--tq` | 9.5 | **11.9** | 3-bit KV → up to 40 960 ctx. FastScan K + V kernels (issue #34) deliver ~20× over the per-block AVX2 path on the attention K+V hot loop; visible end-to-end at long context: **10.2 t/s @ 3K, 9.4 t/s @ 6K** (vs ~5 t/s baseline at 6K). Batched prefill is TQ-aware (issue #34 follow-up) so prompt processing matches the no-`--tq` rate |
+| Qwen3 8B | (same) | 5 GB | Vulkan `-g -1` | 45.4 | 45.8 | 11.4K auto-ctx |
+| Qwen3 8B | (same) | 5 GB | Vulkan `-g -1 --tq` | 40.7 | 45.5 | 3-bit KV → 40 960 ctx |
+| Qwen3 8B | (same) | 5 GB | **CUDA** `-g -1` | **61.7** | **58.6** | ~2.8× Vulkan prefill |
+| Qwen3 8B | (same) | 5 GB | **CUDA** `-g -1 --no-thinking` | **61.8** | **58.2** | Same per-token rate; reasoning suppressed in chat template, so all decoded tokens are visible answer |
+| Qwen3 8B | (same) | 5 GB | **CUDA** `-g -1 --tq` | **57.4** | **58.4** | 3-bit KV → 40 960 ctx; 17 t/s @ 8K, 10 t/s @ 16K |
+| Qwen3 8B | (same) | 5 GB | **CUDA** `-g -1 --tq --no-thinking` | **57.5** | **58.1** | Same per-token rate as `--tq` alone; reasoning suppressed |
+| OLMoE 1B-7B Instruct (MoE) | [allenai](https://huggingface.co/allenai/OLMoE-1B-7B-0924-Instruct-GGUF) | 4 GB | CPU | 51.6 | 55.7 | 64 experts / 8 active; per-channel QK-norm; `norm_topk_prob=false` |
+| OLMoE 1B-7B Instruct (MoE) | (same) | 4 GB | Vulkan `-g -1` | 112.3 | **121.2** | 16 layers all on VRAM; greedy on this prompt is unstable across backends — use `--temp 0.6 --top-p 0.95` for usable output |
+| OLMoE 1B-7B Instruct (MoE) | (same) | 4 GB | **CUDA** `-g -1` | **117.6** | **111.7** | Same; greedy varies, sampling coherent |
+| Qwen3-Coder 30B-A3B (MoE) | [Qwen](https://huggingface.co/Qwen/Qwen3-Coder-30B-A3B-Instruct-GGUF) | 17 GB | CPU | 19.4 | 21.1 | 128 experts / 8 active |
+| Qwen3-Coder 30B-A3B (MoE) | (same) | 17 GB | CPU `--tq` | 18.8 | 21.0 | 3-bit KV. FastScan K + V kernels (issue #34) keep attention cost bounded as context grows: **15.5 t/s decode @ 3.2K ctx** (27 % slowdown for ~27× context growth); without FastScan the per-block K+V path would drop this to ~13 t/s |
+| Qwen3-Coder 30B-A3B (MoE) | (same) | 17 GB | Vulkan `-g -1` (hybrid) | 1.1 | 5.3 | 29 GPU + 19 CPU layers, SLRU expert slot cache. Next-layer predictive prefetch (PR #77 / issue #50) on by default; no-op until the cache is under pressure — disable with `--no-moe-predict-prefetch`. (Prefill not re-measured at ~1K like the other rows — the Vulkan-hybrid MoE path errored on the longer prompt; value is the original short-ctx run) |
+| Qwen3-Coder 30B-A3B (MoE) | (same) | 17 GB | **CUDA** `-g -1` (hybrid) | **24.6** | **25.0** | 29 GPU + 19 CPU layers; routed experts stream through `CudaExpertSlotManager` SLRU (2220 / 3712 slots) instead of the prior eager whole-layer upload (PR #77 / issue #72). Decode lifts 22.2 → 25.0 (+13 %) and prefill 10.6 → 12.3 (+16 %) on top of #77's baseline once the warm-pin + frequency-biased eviction (#74) and async upload stream (#78) hide miss latency in the steady-state working set. Set `SHARPI_EXPERT_STATS=path` to inspect per-layer hit rates |
| Llama-4 Scout 17B-16E (MoE) | [meta-llama](https://huggingface.co/meta-llama/Llama-4-Scout-17B-16E-Instruct) | 61 GB | CPU | 2.1 | 4.3 | 48 layers, 17B active params; split GGUF (Q4_K_M) |
| Llama-4 Scout 17B-16E (MoE) | (same) | 61 GB | CUDA `-g -1` (hybrid) | 1.2 | 2.6 | 7 GPU + 41 CPU layers — model still dwarfs the 12 GB card so CPU-only wins, but per-expert SLRU streaming (PR #77 / issue #72) lifts decode 2.1 → 2.6 (+24 %) and prefill 0.9 → 1.2 (+33 %) over the prior eager whole-layer upload |
-| Qwen3.6-35B-A3B (GDN+MoE) | [unsloth](https://huggingface.co/unsloth/Qwen3.6-35B-A3B-GGUF) | 22 GB | CPU | 6.7 | 8.5 | hybrid GDN/attn, 256 experts / 8 active |
-| Qwen3.6-35B-A3B (GDN+MoE) | (same) | 22 GB | **CUDA** `-g -1` (hybrid) | **14.7** | **23.2** | 10 attn + 30 GDN on GPU; MoE auto-routed to CPU, batched-expert dispatch (8 experts × 3 ops into 2 Parallel.For sweeps), shared expert kept on GPU and overlapped with the CPU routed loop |
-| Qwen3.6-27B-MTP (GDN) | [unsloth](https://huggingface.co/unsloth/Qwen3.6-27B-MTP-GGUF) | 16 GB | CPU `--no-thinking` | 2.8 | **3.8** | dense 27B, hybrid GDN/attn, native MTP head; auto-engages MTP self-spec (issue #25) at greedy + `--no-thinking`. 95% draft acceptance (38/40); batched N=2 verify (#30) + fused Q6_K·Q8_K 2-input dot (#42) lift decode from 2.7 (sequential N=1) to 3.8 — 1.4× over MTP-off baseline |
-| Qwen3.6-27B-MTP (GDN) | (same) | 16 GB | **CUDA** `-g -1 --no-thinking` (hybrid) | **5.7** | **10.7** | 20/64 dense FFN layers on GPU (3.3 GB) + GDN + attn KV resident; 44/64 FFN layers on CPU mmap. 95% draft acceptance; batched verify lifts decode from 6.2 to 10.7 (1.73× over MTP-off baseline). The CPU FFN majority batches via `CpuDenseFfn2` and the on-GPU FFN layers now batch via `MatMulN2` (issue #43 — one weight read per row, two outputs). Direct-pinned `Download/UploadInto` (#48) and async `_lastHidden` overlap (#49) shave per-layer host stall on the MoE-MTP/dense-FFN-MTP hot path |
-| Qwen3.6-27B-MTP (GDN) | (same) | 19 GB | CPU `--no-thinking` `Q5_K_M` | 2.5 | **3.5** | Q5_K_M variant, ~10% slower than Q4_K_M as expected from weight bandwidth. 100% draft acceptance (40/40) on this prompt; batched verify lifts decode from 2.4 to 3.5 (1.46×) |
-| Qwen3.6-27B-MTP (GDN) | (same) | 19 GB | **CUDA** `-g -1 --no-thinking` `Q5_K_M` (hybrid) | 4.4 | **7.9** | 13/64 FFN layers on GPU (2.4 GB) + GDN + attn KV resident; 51/64 FFN on CPU mmap. Uses `llm_embed_lookup_q5k` direct-read kernel (issue #39) and the Q5_K `MatMulN2` kernel (issue #43) for the on-GPU layers. 98% draft acceptance; batched verify lifts decode from 4.3 to 7.9 (1.84×). Direct-pinned `Download/UploadInto` (#48) + async `_lastHidden` overlap (#49) land 0.5-0.8 t/s on top of #43's batched MatMul win — measurable here because Q5_K MatMul is more PCIe-bound than Q4_K |
-| Qwen3.6-35B-A3B-MTP (GDN+MoE) | [unsloth](https://huggingface.co/unsloth/Qwen3.6-35B-A3B-MTP-GGUF) | 22 GB | CPU `--no-thinking` | 6.9 | 8.0 | hybrid GDN/attn + 256-expert MoE + native MTP head (issue #44). 100% draft acceptance. Issue #45 enables `BatchForward2` for MoE MTP — attn/GDN/lm_head batch across t1/t2 but routed experts run sequentially per token (top-K differs), so the win is small (lm_head save + frame overhead). At parity with the 8.2 MTP-off baseline within CPU jitter |
-| Qwen3.6-35B-A3B-MTP (GDN+MoE) | (same) | 22 GB | **CUDA** `-g -1 --no-thinking` (hybrid) | **13.0** | **22.9** | Requires `SHARPI_CPU_MOE=1`: 30 GDN + 10 attn + shared expert on GPU, MoE routed experts + MTP MoE FFN mmap'd on CPU. 100% draft acceptance. Issue #45 lifts decode from 22.9 (sequential MTP at 0.99× MTP-off baseline) to at-or-above MTP-off — modest because routed-expert weight reads can't share between tokens; the bandwidth-bound CPU MoE FFN runs sequentially per token. Issues #47 (async UploadViaStaging) + #49 (overlap `_lastHidden` D2H with lm_head MatMul) shave further µs/layer |
-| Carnice (Qwen3.6-35B-A3B-MTP finetune) | [mudler](https://huggingface.co/mudler/Carnice-Qwen3.6-MoE-35B-A3B-APEX-MTP-GGUF) | 17 GB | **CUDA** `-g -1 --no-thinking` (hybrid) | **20.6** | **25.0** | APEX-MTP-I-Compact mixed-precision quant: Q3_K + Q8_0 routed experts (rest Q4_K/Q5_K). Hermes-style agentic/tool-calling finetune of 35B-A3B-MTP; same qwen35moe arch + MTP head; 77% draft acceptance on the `bench-carnice.ps1` 60-token prompt (the `bench-all` "capital of France" default elicits a 1-token EOS on Carnice's terser agentic tuning, so it gets its own bench cell). Numbers are apples-to-apples warm-cache; the prior 5.3 t/s decode in this row was a cold-cache artefact (mmap'd expert pages). Issues #101 / #99 / #107 add `DotQ3K_Q8KS_Avx2` + `DotQ8_0_Q8KS_Avx2` (dual-acc-chain int-domain dots over a **Q8_KS** per-32-element-scale input — 8 FP scales per 256-elem super-block instead of Q8_K's one per 256), **auto-enabled at model load** when the GGUF tensor index shows matching routed-expert dtypes (APEX mixed-precision tier) — this lifts warm decode from 23.9 (forced off) to 25.0 t/s (+4.6%). Force off with `SHARPI_Q3K_Q8K=0` / `SHARPI_Q8_0_Q8K=0` to fall back to the FP `DotQ3K` / `DotQ8_0` path; output is not bit-identical (cumulative trunk drift on Qwen3.6-class models, see `feedback_q4k_q8k_no_parity_win`) but Q8_KS cuts the parity envelope ~4× vs the plain-Q8_K probe (#103 validation: ±13 pp MTP-accept → ±3 pp, every prompt argmax-stable through ≥32 tokens), and all 294 Tests.ForwardPass tests stay green with the kernels engaged |
-| Gemma 4 E4B-it Q8 | [unsloth](https://huggingface.co/unsloth/gemma-4-E4B-it-GGUF) | 8 GB | CPU | 4.3 | 5.0 | dense 42-layer gemma4 (E4B): per-layer head_dim (256 SWA / 512 global), dual-RoPE (10K SWA / 1M global with `rope_freqs.weight` baked in), KV-share tail (shared_kv_layers=18; last 18 layers alias layer 22 or 23), 5 SWA : 1 global attention pattern, post-attn + post-ffw RmsNorms, `layer_output_scale` per-layer scalar, final-logit softcap=30, GeluTanh FFN, PLE-256 injection (per_layer_token_embd ~4.2 GB stays mmap-resident, per_layer_model_proj BF16→F32 dequant once). `SimdKernels.RmsNormWide` (AVX-512, 16 floats/iter) + K/V `MatVecDual` fuse model-gated to per-layer-head_dim path; other models stay on AVX2 because Qwen3.6-27B-MTP byte-parity is sensitive to the ~ULP reduction-order shift |
-| Gemma 4 E4B-it Q8 | (same) | 8 GB | **CUDA** `-g -1 -c 512` | **51.3** | **44.1** | All 42 layers in 12 GB at small context. Attention scale fix (Gemma 4 wants 1.0, kernels apply 1/sqrt(head_dim) so Q is pre-scaled by sqrt(layerHd) — would otherwise produce a degenerate 2-cycle repeat after first decode). KV-share alias dispatch + SWA vs full attention split per layer; `rope_freqs.weight` applied via `RoPEWithFactors` on global layers, plain `RoPE` on SWA. PLE projections (~215 MB across 42 layers) upload at construction; per-token row dequant + projection runs on GPU |
-| Gemma 4 E4B-it Q8 | (same) | 8 GB | **CUDA** `-g 22 -c 2048` (hybrid) | 7.1 | 6.8 | 22 GPU + 20 CPU layers @ 2K context. `-g <= 22` constraint: own-KV source layers (22, 23) must live on CPU so the shared-KV tail (24..41, all on CPU) can read them without cross-tier KV reads. CPU half dominates decode (dense FFN reads ~28 MB Q8 per layer × 20 layers = ~560 MB/token, bandwidth-bound on Zen 4 DDR4-3200); GPU half overlaps the embed/output norm. SHARPI_CUDA_PROFILE=1 dumps per-phase breakdown (embed / ple / qkv / rope+qknorm / kv+attn / o-proj+res / ffn / final) |
+| Qwen3.6-35B-A3B (GDN+MoE) | [unsloth](https://huggingface.co/unsloth/Qwen3.6-35B-A3B-GGUF) | 22 GB | CPU | 8.4 | 8.5 | hybrid GDN/attn, 256 experts / 8 active |
+| Qwen3.6-35B-A3B (GDN+MoE) | (same) | 22 GB | **CUDA** `-g -1` (hybrid) | **63.7** | **23.2** | 10 attn + 30 GDN on GPU; MoE auto-routed to CPU, batched-expert dispatch (8 experts × 3 ops into 2 Parallel.For sweeps), shared expert kept on GPU and overlapped with the CPU routed loop. Issue #114-B collapses the per-position GDN recurrence into one fused sequential-scan kernel and the KV-append + SDPA into batched-query launches (`llm_full_seq_attention`), bit-identical to the per-token path; the win grows with context (A/B at ~4K ctx: +17%, 55.3 → 64.5 t/s, over the #111 GEMM-batched-projection trunk). Decode is the near-zero-ctx generation rate |
+| Qwen3.6-27B-MTP (GDN) | [unsloth](https://huggingface.co/unsloth/Qwen3.6-27B-MTP-GGUF) | 16 GB | CPU `--no-thinking` | 3.2 | **3.8** | dense 27B, hybrid GDN/attn, native MTP head; auto-engages MTP self-spec (issue #25) at greedy + `--no-thinking`. 95% draft acceptance (38/40); batched N=2 verify (#30) + fused Q6_K·Q8_K 2-input dot (#42) lift decode from 2.7 (sequential N=1) to 3.8 — 1.4× over MTP-off baseline |
+| Qwen3.6-27B-MTP (GDN) | (same) | 16 GB | **CUDA** `-g -1 --no-thinking` (hybrid) | **6.4** | **10.7** | 20/64 dense FFN layers on GPU (3.3 GB) + GDN + attn KV resident; 44/64 FFN layers on CPU mmap. 95% draft acceptance; batched verify lifts decode from 6.2 to 10.7 (1.73× over MTP-off baseline). The CPU FFN majority batches via `CpuDenseFfn2` and the on-GPU FFN layers now batch via `MatMulN2` (issue #43 — one weight read per row, two outputs). Direct-pinned `Download/UploadInto` (#48) and async `_lastHidden` overlap (#49) shave per-layer host stall on the MoE-MTP/dense-FFN-MTP hot path |
+| Qwen3.6-27B-MTP (GDN) | (same) | 19 GB | CPU `--no-thinking` `Q5_K_M` | 2.8 | **3.5** | Q5_K_M variant, ~10% slower than Q4_K_M as expected from weight bandwidth. 100% draft acceptance (40/40) on this prompt; batched verify lifts decode from 2.4 to 3.5 (1.46×) |
+| Qwen3.6-27B-MTP (GDN) | (same) | 19 GB | **CUDA** `-g -1 --no-thinking` `Q5_K_M` (hybrid) | 4.2 | **7.9** | 13/64 FFN layers on GPU (2.4 GB) + GDN + attn KV resident; 51/64 FFN on CPU mmap. Uses `llm_embed_lookup_q5k` direct-read kernel (issue #39) and the Q5_K `MatMulN2` kernel (issue #43) for the on-GPU layers. 98% draft acceptance; batched verify lifts decode from 4.3 to 7.9 (1.84×). Direct-pinned `Download/UploadInto` (#48) + async `_lastHidden` overlap (#49) land 0.5-0.8 t/s on top of #43's batched MatMul win — measurable here because Q5_K MatMul is more PCIe-bound than Q4_K |
+| Qwen3.6-35B-A3B-MTP (GDN+MoE) | [unsloth](https://huggingface.co/unsloth/Qwen3.6-35B-A3B-MTP-GGUF) | 22 GB | CPU `--no-thinking` | 8.5 | 8.0 | hybrid GDN/attn + 256-expert MoE + native MTP head (issue #44). 100% draft acceptance. Issue #45 enables `BatchForward2` for MoE MTP — attn/GDN/lm_head batch across t1/t2 but routed experts run sequentially per token (top-K differs), so the win is small (lm_head save + frame overhead). At parity with the 8.2 MTP-off baseline within CPU jitter |
+| Qwen3.6-35B-A3B-MTP (GDN+MoE) | (same) | 22 GB | **CUDA** `-g -1 --no-thinking` (hybrid) | **65.0** | **22.9** | Requires `SHARPI_CPU_MOE=1`: 30 GDN + 10 attn + shared expert on GPU, MoE routed experts + MTP MoE FFN mmap'd on CPU. 100% draft acceptance. Issue #45 lifts decode from 22.9 (sequential MTP at 0.99× MTP-off baseline) to at-or-above MTP-off — modest because routed-expert weight reads can't share between tokens; the bandwidth-bound CPU MoE FFN runs sequentially per token. Issues #47 (async UploadViaStaging) + #49 (overlap `_lastHidden` D2H with lm_head MatMul) shave further µs/layer. Issue #114-B's fused GDN sequential-scan + batched-query SDPA (`llm_full_seq_attention`) is bit-identical to the per-token path and grows with context (A/B at ~4K ctx: +17%, 55.4 → 64.7 t/s, over the #111 per-position trunk). Decode is the near-zero-ctx MTP generation rate |
+| Carnice (Qwen3.6-35B-A3B-MTP finetune) | [mudler](https://huggingface.co/mudler/Carnice-Qwen3.6-MoE-35B-A3B-APEX-MTP-GGUF) | 17 GB | **CUDA** `-g -1 --no-thinking` (hybrid) | **43.6** | **25.0** | Issue #114-B collapses the per-position GDN recurrence (fused sequential-scan kernel) + KV-append/SDPA (`llm_full_seq_attention`) into one launch each, bit-identical to the per-token path (per-kernel bit-exactness + `BatchedTrunk_BitwiseMatchesSequentialTrunk` oracles, full Tests.ForwardPass suite green); the win grows with context (A/B at ~4K ctx: +31%, 92.1 → 120.8 t/s, over the #111 GEMM-batched-projection trunk). Decode shown is the near-zero-ctx MTP generation rate. APEX-MTP-I-Compact mixed-precision quant: Q3_K + Q8_0 routed experts (rest Q4_K/Q5_K). Hermes-style agentic/tool-calling finetune of 35B-A3B-MTP; same qwen35moe arch + MTP head; 77% draft acceptance on the `bench-carnice.ps1` 60-token prompt (the `bench-all` "capital of France" default elicits a 1-token EOS on Carnice's terser agentic tuning, so it gets its own bench cell). Numbers are apples-to-apples warm-cache; the prior 5.3 t/s decode in this row was a cold-cache artefact (mmap'd expert pages). Issues #101 / #99 / #107 add `DotQ3K_Q8KS_Avx2` + `DotQ8_0_Q8KS_Avx2` (dual-acc-chain int-domain dots over a **Q8_KS** per-32-element-scale input — 8 FP scales per 256-elem super-block instead of Q8_K's one per 256), **auto-enabled at model load** when the GGUF tensor index shows matching routed-expert dtypes (APEX mixed-precision tier) — this lifts warm decode from 23.9 (forced off) to 25.0 t/s (+4.6%). Force off with `SHARPI_Q3K_Q8K=0` / `SHARPI_Q8_0_Q8K=0` to fall back to the FP `DotQ3K` / `DotQ8_0` path; output is not bit-identical (cumulative trunk drift on Qwen3.6-class models, see `feedback_q4k_q8k_no_parity_win`) but Q8_KS cuts the parity envelope ~4× vs the plain-Q8_K probe (#103 validation: ±13 pp MTP-accept → ±3 pp, every prompt argmax-stable through ≥32 tokens), and all 294 Tests.ForwardPass tests stay green with the kernels engaged |
+| Gemma 4 E4B-it Q8 | [unsloth](https://huggingface.co/unsloth/gemma-4-E4B-it-GGUF) | 8 GB | CPU | 4.9 | 5.0 | dense 42-layer gemma4 (E4B): per-layer head_dim (256 SWA / 512 global), dual-RoPE (10K SWA / 1M global with `rope_freqs.weight` baked in), KV-share tail (shared_kv_layers=18; last 18 layers alias layer 22 or 23), 5 SWA : 1 global attention pattern, post-attn + post-ffw RmsNorms, `layer_output_scale` per-layer scalar, final-logit softcap=30, GeluTanh FFN, PLE-256 injection (per_layer_token_embd ~4.2 GB stays mmap-resident, per_layer_model_proj BF16→F32 dequant once). `SimdKernels.RmsNormWide` (AVX-512, 16 floats/iter) + K/V `MatVecDual` fuse model-gated to per-layer-head_dim path; other models stay on AVX2 because Qwen3.6-27B-MTP byte-parity is sensitive to the ~ULP reduction-order shift |
+| Gemma 4 E4B-it Q8 | (same) | 8 GB | **CUDA** `-g -1 -c 2048` | **47.7** | **44.1** | All 42 layers fit in 12 GB (still fits at `-c 2048`, used here for the ~1K-token prefill cell; decode shown is the small-ctx rate). Attention scale fix (Gemma 4 wants 1.0, kernels apply 1/sqrt(head_dim) so Q is pre-scaled by sqrt(layerHd) — would otherwise produce a degenerate 2-cycle repeat after first decode). KV-share alias dispatch + SWA vs full attention split per layer; `rope_freqs.weight` applied via `RoPEWithFactors` on global layers, plain `RoPE` on SWA. PLE projections (~215 MB across 42 layers) upload at construction; per-token row dequant + projection runs on GPU |
+| Gemma 4 E4B-it Q8 | (same) | 8 GB | **CUDA** `-g 22 -c 2048` (hybrid) | 6.6 | 6.8 | 22 GPU + 20 CPU layers @ 2K context. `-g <= 22` constraint: own-KV source layers (22, 23) must live on CPU so the shared-KV tail (24..41, all on CPU) can read them without cross-tier KV reads. CPU half dominates decode (dense FFN reads ~28 MB Q8 per layer × 20 layers = ~560 MB/token, bandwidth-bound on Zen 4 DDR4-3200); GPU half overlaps the embed/output norm. SHARPI_CUDA_PROFILE=1 dumps per-phase breakdown (embed / ple / qkv / rope+qknorm / kv+attn / o-proj+res / ffn / final) |
+
+_Methodology: **Prefill t/s** is the warm-cache rate at a **~1K-token prompt** (≈1000–1120 tokens
+per row — steady-state prompt processing), re-measured across every on-disk row so the column is
+comparable rather than the old ~10–60-token launch-overhead cells. **Decode t/s** is the
+near-zero-ctx generation rate (the conventional generation-speed metric; the per-issue before/after
+figures in the notes are historical at-the-time measurements). Llama-4 Scout is not on the bench
+machine (kept at its prior numbers); Qwen3-Coder Vulkan-hybrid errored on the ~1K prompt so its
+prefill is the original short-ctx value._
`--backend auto` (default) picks CUDA when available, sizing the GPU/CPU split from
VRAM via TierPlanner; falls through to Vulkan only when CUDA isn't present.
diff --git a/scripts/bench-allrows-1k.ps1 b/scripts/bench-allrows-1k.ps1
new file mode 100644
index 0000000..ac196e9
--- /dev/null
+++ b/scripts/bench-allrows-1k.ps1
@@ -0,0 +1,90 @@
+# Re-benchmark every on-disk README row at a uniform ~1K-token "normal" context,
+# warm-cache, current code (#114-B batched trunk on by default). Goal: make the
+# README "Prefill t/s" column consistent (warm @~1K ctx) instead of the old
+# ~10-token launch-overhead cells. Decode is captured for reference but the README
+# decode column / notes stay as the near-zero-ctx generation rate.
+$ErrorActionPreference = "Continue"
+
+$C = "C:\p\sharpi\models"
+$E = "E:\models"
+
+# ~1K-token prompt: a few varied technical sections + a task (not pure repetition).
+$para = @(
+"Modern large language model inference is dominated by memory bandwidth rather than raw compute.",
+"Each decode step streams the full weight matrix for every layer, so quantization formats like Q4_K and Q5_K trade a small accuracy loss for a large reduction in bytes moved per token.",
+"Mixture-of-experts models complicate this: only a handful of the hundreds of experts fire per token, and which fire varies token to token, defeating simple weight caching.",
+"Gated DeltaNet layers replace quadratic attention with a linear recurrent state update, bounding per-token cost as context grows, at the price of a strictly sequential scan.",
+"Hybrid placement keeps the attention and recurrent trunk on the accelerator while streaming routed-expert weights from host memory, overlapping the two so neither stalls."
+) -join " "
+$sb = [System.Text.StringBuilder]::new()
+[void]$sb.Append("Read the following engineering notes and then write a concise technical summary.`n`n")
+for ($i = 1; $i -le 6; $i++) { [void]$sb.Append("Section $i. $para`n`n") }
+[void]$sb.Append("Summarize the main performance trade-offs across the sections above.")
+$prompt = $sb.ToString()
+
+# Job table: Tag, Model, Args, Env (CPU_MOE), Timeout. Order groups by model file so
+# the first run of each model warms the OS page cache for the ones after it.
+$jobs = @(
+ @{ Tag="smol-cpu"; M="$C\SmolLM2-1.7B-Instruct-Q4_K_M.gguf"; A=@(); T=300 }
+ @{ Tag="smol-vulkan"; M="$C\SmolLM2-1.7B-Instruct-Q4_K_M.gguf"; A=@("-g","-1","--backend","vulkan"); T=300 }
+ @{ Tag="smol-cuda"; M="$C\SmolLM2-1.7B-Instruct-Q4_K_M.gguf"; A=@("-g","-1","--backend","cuda"); T=300 }
+
+ @{ Tag="qwen3-cpu"; M="$C\Qwen3-8B-Q4_K_M.gguf"; A=@(); T=400 }
+ @{ Tag="qwen3-cpu-tq"; M="$C\Qwen3-8B-Q4_K_M.gguf"; A=@("--tq"); T=400 }
+ @{ Tag="qwen3-vulkan"; M="$C\Qwen3-8B-Q4_K_M.gguf"; A=@("-g","-1","--backend","vulkan"); T=400 }
+ @{ Tag="qwen3-vulkan-tq"; M="$C\Qwen3-8B-Q4_K_M.gguf"; A=@("-g","-1","--backend","vulkan","--tq"); T=400 }
+ @{ Tag="qwen3-cuda"; M="$C\Qwen3-8B-Q4_K_M.gguf"; A=@("-g","-1","--backend","cuda"); T=400 }
+ @{ Tag="qwen3-cuda-nt"; M="$C\Qwen3-8B-Q4_K_M.gguf"; A=@("-g","-1","--backend","cuda","--no-thinking"); T=400 }
+ @{ Tag="qwen3-cuda-tq"; M="$C\Qwen3-8B-Q4_K_M.gguf"; A=@("-g","-1","--backend","cuda","--tq"); T=400 }
+ @{ Tag="qwen3-cuda-tq-nt";M="$C\Qwen3-8B-Q4_K_M.gguf"; A=@("-g","-1","--backend","cuda","--tq","--no-thinking"); T=400 }
+
+ @{ Tag="olmoe-cpu"; M="$C\OLMoE-1B-7B-0924-Instruct-Q4_K_M.gguf"; A=@(); T=400 }
+ @{ Tag="olmoe-vulkan"; M="$C\OLMoE-1B-7B-0924-Instruct-Q4_K_M.gguf"; A=@("-g","-1","--backend","vulkan"); T=400 }
+ @{ Tag="olmoe-cuda"; M="$C\OLMoE-1B-7B-0924-Instruct-Q4_K_M.gguf"; A=@("-g","-1","--backend","cuda"); T=400 }
+
+ @{ Tag="coder-cpu"; M="$C\Qwen3-Coder-30B-A3B-Instruct-Q4_K_M.gguf"; A=@(); T=600 }
+ @{ Tag="coder-cpu-tq"; M="$C\Qwen3-Coder-30B-A3B-Instruct-Q4_K_M.gguf"; A=@("--tq"); T=600 }
+ @{ Tag="coder-vulkan"; M="$C\Qwen3-Coder-30B-A3B-Instruct-Q4_K_M.gguf"; A=@("-g","-1","--backend","vulkan"); T=1800 }
+ @{ Tag="coder-cuda"; M="$C\Qwen3-Coder-30B-A3B-Instruct-Q4_K_M.gguf"; A=@("-g","-1","--backend","cuda"); T=600 }
+
+ @{ Tag="qwen36-35b-cpu"; M="$E\Qwen3.6-35B-A3B-UD-Q4_K_M.gguf"; A=@(); T=900 }
+ @{ Tag="qwen36-35b-cuda"; M="$E\Qwen3.6-35B-A3B-UD-Q4_K_M.gguf"; A=@("-g","-1","--backend","cuda"); T=900 }
+
+ @{ Tag="27b-mtp-q4-cpu"; M="$E\Qwen3.6-27B-MTP-Q4_K_M.gguf"; A=@("--no-thinking"); T=900 }
+ @{ Tag="27b-mtp-q4-cuda"; M="$E\Qwen3.6-27B-MTP-Q4_K_M.gguf"; A=@("-g","-1","--backend","cuda","--no-thinking"); T=900 }
+ @{ Tag="27b-mtp-q5-cpu"; M="$E\Qwen3.6-27B-MTP-Q5_K_M.gguf"; A=@("--no-thinking"); T=900 }
+ @{ Tag="27b-mtp-q5-cuda"; M="$E\Qwen3.6-27B-MTP-Q5_K_M.gguf"; A=@("-g","-1","--backend","cuda","--no-thinking"); T=900 }
+
+ @{ Tag="35b-mtp-cpu"; M="$E\Qwen3.6-35B-A3B-MTP-UD-Q4_K_M.gguf"; A=@("--no-thinking"); CpuMoe=$true; T=900 }
+ @{ Tag="35b-mtp-cuda"; M="$E\Qwen3.6-35B-A3B-MTP-UD-Q4_K_M.gguf"; A=@("-g","-1","--backend","cuda","--no-thinking"); CpuMoe=$true; T=900 }
+
+ @{ Tag="carnice-cuda"; M="$E\Carnice-Qwen3.6-MoE-35B-A3B-APEX-MTP-I-Compact.gguf"; A=@("-g","-1","--backend","cuda","--no-thinking"); CpuMoe=$true; T=900 }
+
+ @{ Tag="gemma4-cpu"; M="$E\gemma-4-E4B-it-Q8_0.gguf"; A=@(); T=900 }
+ @{ Tag="gemma4-cuda"; M="$E\gemma-4-E4B-it-Q8_0.gguf"; A=@("-g","-1","--backend","cuda","--ctx-size","2048"); T=900 }
+ @{ Tag="gemma4-cuda-hyb"; M="$E\gemma-4-E4B-it-Q8_0.gguf"; A=@("-g","22","--backend","cuda","--ctx-size","2048"); T=900 }
+)
+
+$warmed = @{}
+$rows = @()
+foreach ($j in $jobs) {
+ if (-not (Test-Path $j.M)) { Write-Host "[skip] $($j.Tag): $($j.M) missing" -ForegroundColor Yellow; continue }
+ if ($j.CpuMoe) { $env:SHARPI_CPU_MOE = "1" } else { Remove-Item env:SHARPI_CPU_MOE -ErrorAction SilentlyContinue }
+
+ # Warm the OS page cache for this model file once (short prompt, discarded).
+ if (-not $warmed.ContainsKey($j.M)) {
+ Write-Host "--- warming $($j.Tag) model ---" -ForegroundColor DarkGray
+ $null = .\scripts\bench-textgen.ps1 -Model $j.M -Tag "$($j.Tag)-warm1k" -NTokens 8 -Prompt "Hello, world." -TimeoutSec $j.T -ExtraArgs $j.A
+ $warmed[$j.M] = $true
+ }
+
+ $r = .\scripts\bench-textgen.ps1 -Model $j.M -Tag "$($j.Tag)-1k" -NTokens 60 -Prompt $prompt -TimeoutSec $j.T -ExtraArgs $j.A
+ Remove-Item env:SHARPI_CPU_MOE -ErrorAction SilentlyContinue
+ $rows += [PSCustomObject]@{ Tag=$j.Tag; PrefTok=$r.PrefillTok; PrefillTps=$r.PrefillTps; DecodeTps=$r.DecodeTps; Mtp=$r.MtpAccept; Wall=$r.WallSec; TO=$r.TimedOut }
+ Write-Host (" {0,-18} pref={1,7} t/s dec={2,6} t/s ({3} tok, {4}s{5})" -f $j.Tag,$r.PrefillTps,$r.DecodeTps,$r.PrefillTok,$r.WallSec,($(if($r.TimedOut){" TIMEOUT"}else{""}))) -ForegroundColor Green
+}
+Write-Host ""
+Write-Host "=== All-rows @~1K ctx (warm) ===" -ForegroundColor Cyan
+$rows | Format-Table -AutoSize
+$rows | Export-Csv -NoTypeInformation -Path "tools\bench\allrows-1k.csv"
+Write-Host "CSV: tools\bench\allrows-1k.csv"
diff --git a/src/SharpInference.Cuda/CudaBackend.cs b/src/SharpInference.Cuda/CudaBackend.cs
index 730fc03..34c9d73 100644
--- a/src/SharpInference.Cuda/CudaBackend.cs
+++ b/src/SharpInference.Cuda/CudaBackend.cs
@@ -175,6 +175,17 @@ public sealed unsafe class CudaBackend : IComputeBackend, IImageOpsBackend, IDis
private nint _gdnTileHeadsKernel;
private nint _gdnRecurrenceDecodeKernel;
+ // Issue #114-B: batched GDN trunk + batched-query SDPA kernels.
+ private nint _gdnConv1dDecodeBatchedKernel;
+ private nint _gdnConv1dStateUpdateBatchedKernel;
+ private nint _gdnL2NormPerHeadBatchedKernel;
+ private nint _gdnTileHeadsBatchedKernel;
+ private nint _gdnRecurrenceScanKernel;
+ private nint _kvAppendBatchedKernel;
+ private nint _kvAppendBatchedBf16Kernel;
+ private nint _fullSeqAttentionKernel;
+ private nint _fullSeqAttentionBf16Kernel;
+
// Persistent Q8_1 scratch for the Q4_K matvec input. Grows on demand and
// never shrinks. Sized in 36-byte sub-blocks (one block_q8_1 per 32 elements).
private nint _q81Buf;
@@ -2242,6 +2253,114 @@ public void AttentionBf16(Tensor q, Tensor kCache, Tensor vCache, Tensor output,
if (r != 0) throw new InvalidOperationException($"cuLaunchKernel(attention_bf16) failed: {r}");
}
+ // ── Issue #114-B: batched prompt-prefill SDPA ──────────────────────────
+
+ ///
+ /// Batched : writes the K/V vectors for
+ /// tokens into the cache at consecutive positions startPos .. startPos+nTok-1
+ /// in a single launch. / are
+ /// [nTok × kvDim] token-major. Bit-identical to nTok sequential
+ /// calls.
+ ///
+ public void KvAppendBatched(Tensor kAll, Tensor vAll, Tensor kCache, Tensor vCache,
+ int kvDim, int startPos, int maxSeqLen, int nTok)
+ {
+ EnsureImageKernels();
+ if (!_imageKernelsAvailable)
+ throw new NotSupportedException("NVRTC kernels are not available.");
+
+ nint kPtr = GetDevPtr(kAll), vPtr = GetDevPtr(vAll);
+ nint kcP = GetDevPtr(kCache), vcP = GetDevPtr(vCache);
+ int pKD = kvDim, pSP = startPos, pMSL = maxSeqLen, pN = nTok;
+ nint* args = stackalloc nint[8]
+ {
+ (nint)(&kPtr), (nint)(&vPtr), (nint)(&kcP), (nint)(&vcP),
+ (nint)(&pKD), (nint)(&pSP), (nint)(&pMSL), (nint)(&pN)
+ };
+ uint grid = (uint)((kvDim + 255) / 256);
+ int r = NvrtcInterop.LaunchKernel(_kvAppendBatchedKernel, grid, (uint)nTok, 1, 256, 1, 1, 0, _stream, args, null);
+ if (r != 0) throw new InvalidOperationException($"cuLaunchKernel(kv_append_batched) failed: {r}");
+ }
+
+ /// Bf16-store variant of (default KV dtype).
+ public void KvAppendBatchedBf16(Tensor kAll, Tensor vAll, Tensor kCache, Tensor vCache,
+ int kvDim, int startPos, int maxSeqLen, int nTok)
+ {
+ EnsureImageKernels();
+ if (!_imageKernelsAvailable)
+ throw new NotSupportedException("NVRTC kernels are not available.");
+
+ nint kPtr = GetDevPtr(kAll), vPtr = GetDevPtr(vAll);
+ nint kcP = GetDevPtr(kCache), vcP = GetDevPtr(vCache);
+ int pKD = kvDim, pSP = startPos, pMSL = maxSeqLen, pN = nTok;
+ nint* args = stackalloc nint[8]
+ {
+ (nint)(&kPtr), (nint)(&vPtr), (nint)(&kcP), (nint)(&vcP),
+ (nint)(&pKD), (nint)(&pSP), (nint)(&pMSL), (nint)(&pN)
+ };
+ uint grid = (uint)((kvDim + 255) / 256);
+ int r = NvrtcInterop.LaunchKernel(_kvAppendBatchedBf16Kernel, grid, (uint)nTok, 1, 256, 1, 1, 0, _stream, args, null);
+ if (r != 0) throw new InvalidOperationException($"cuLaunchKernel(kv_append_batched_bf16) failed: {r}");
+ }
+
+ ///
+ /// Batched-query scaled dot-product attention (issue #114-B). All
+ /// prompt queries attend over their causal prefix in a
+ /// single launch (grid = numHeads × nTok), instead of nTok sequential
+ /// launches. Query i (row of ,
+ /// [nTok × numHeads·headDim]) attends over cache positions
+ /// [0, startPos+i+1); output written to in the
+ /// same layout. Bit-identical to the per-token path.
+ ///
+ /// Constraint: uses the shared-scores fast path only, so the caller
+ /// MUST guarantee startPos + nTok ≤ 4096 (every block's seqLen stays
+ /// ≤ MAX_STORED_SCORES). Beyond that, fall back to the per-token loop.
+ ///
+ public void AttentionBatched(Tensor qAll, Tensor kCache, Tensor vCache, Tensor outAll,
+ int numHeads, int numKvHeads, int headDim,
+ int startPos, int maxSeqLen, int nTok)
+ {
+ EnsureImageKernels();
+ if (!_imageKernelsAvailable)
+ throw new NotSupportedException("NVRTC kernels are not available.");
+ if (startPos + nTok > 4096)
+ throw new ArgumentException(
+ $"AttentionBatched requires startPos+nTok ≤ 4096 (shared-scores path); got {startPos}+{nTok}.");
+
+ nint qP = GetDevPtr(qAll), kP = GetDevPtr(kCache), vP = GetDevPtr(vCache), oP = GetDevPtr(outAll);
+ int pNH = numHeads, pNKV = numKvHeads, pHD = headDim, pSP = startPos, pMSL = maxSeqLen, pN = nTok;
+ nint* args = stackalloc nint[10]
+ {
+ (nint)(&qP), (nint)(&kP), (nint)(&vP), (nint)(&oP),
+ (nint)(&pNH), (nint)(&pNKV), (nint)(&pHD), (nint)(&pSP), (nint)(&pMSL), (nint)(&pN)
+ };
+ int r = NvrtcInterop.LaunchKernel(_fullSeqAttentionKernel, (uint)numHeads, (uint)nTok, 1, 256, 1, 1, 0, _stream, args, null);
+ if (r != 0) throw new InvalidOperationException($"cuLaunchKernel(full_seq_attention) failed: {r}");
+ }
+
+ /// Bf16-read variant of (default KV dtype).
+ public void AttentionBatchedBf16(Tensor qAll, Tensor kCache, Tensor vCache, Tensor outAll,
+ int numHeads, int numKvHeads, int headDim,
+ int startPos, int maxSeqLen, int nTok)
+ {
+ EnsureImageKernels();
+ if (!_imageKernelsAvailable)
+ throw new NotSupportedException("NVRTC kernels are not available.");
+ if (startPos + nTok > 4096)
+ throw new ArgumentException(
+ $"AttentionBatchedBf16 requires startPos+nTok ≤ 4096 (shared-scores path); got {startPos}+{nTok}.");
+
+ nint qP = GetDevPtr(qAll), kP = GetDevPtr(kCache), vP = GetDevPtr(vCache), oP = GetDevPtr(outAll);
+ int pNH = numHeads, pNKV = numKvHeads, pHD = headDim, pSP = startPos, pMSL = maxSeqLen, pN = nTok;
+ nint* args = stackalloc nint[10]
+ {
+ (nint)(&qP), (nint)(&kP), (nint)(&vP), (nint)(&oP),
+ (nint)(&pNH), (nint)(&pNKV), (nint)(&pHD), (nint)(&pSP), (nint)(&pMSL), (nint)(&pN)
+ };
+ int r = NvrtcInterop.LaunchKernel(_fullSeqAttentionBf16Kernel, (uint)numHeads, (uint)nTok, 1, 256, 1, 1, 0, _stream, args, null);
+ if (r != 0) throw new InvalidOperationException($"cuLaunchKernel(full_seq_attention_bf16) failed: {r}");
+ }
+
// ================================================================
// SnapKV (issue #58): prefill KV eviction support
// ================================================================
@@ -2754,6 +2873,143 @@ public void GdnRecurrenceDecode(
if (r != 0) throw new InvalidOperationException($"cuLaunchKernel(gdn_recurrence_decode) failed: {r}");
}
+ // ── Issue #114-B: batched GDN trunk over a chunk of N prompt tokens ────────
+
+ ///
+ /// Batched GDN depthwise conv1d over tokens (read-only
+ /// state). / are [nTok × channels];
+ /// is the carried [(K-1) × channels]. Bit-identical to
+ /// nTok sequential calls. State is advanced separately by
+ /// (so concurrent token blocks read one snapshot).
+ ///
+ public void GdnConv1dDecodeBatched(Tensor x, Tensor state, Tensor weight, Tensor output,
+ int channels, int kernelSize, int nTok)
+ {
+ EnsureImageKernels();
+ if (!_imageKernelsAvailable)
+ throw new NotSupportedException("NVRTC kernels are not available.");
+
+ nint xP = GetDevPtr(x), sP = GetDevPtr(state), wP = GetDevPtr(weight), oP = GetDevPtr(output);
+ int pC = channels, pK = kernelSize, pN = nTok;
+ nint* args = stackalloc nint[7]
+ {
+ (nint)(&xP), (nint)(&sP), (nint)(&wP), (nint)(&oP),
+ (nint)(&pC), (nint)(&pK), (nint)(&pN)
+ };
+ uint grid = (uint)((channels + 255) / 256);
+ int r = NvrtcInterop.LaunchKernel(_gdnConv1dDecodeBatchedKernel, grid, (uint)nTok, 1, 256, 1, 1, 0, _stream, args, null);
+ if (r != 0) throw new InvalidOperationException($"cuLaunchKernel(gdn_conv1d_decode_batched) failed: {r}");
+ }
+
+ /// Advance the conv1d state past a chunk of tokens
+ /// (matches the sequential state evolution). See .
+ public void GdnConv1dStateUpdateBatched(Tensor x, Tensor state, int channels, int kernelSize, int nTok)
+ {
+ EnsureImageKernels();
+ if (!_imageKernelsAvailable)
+ throw new NotSupportedException("NVRTC kernels are not available.");
+
+ nint xP = GetDevPtr(x), sP = GetDevPtr(state);
+ int pC = channels, pK = kernelSize, pN = nTok;
+ nint* args = stackalloc nint[5] { (nint)(&xP), (nint)(&sP), (nint)(&pC), (nint)(&pK), (nint)(&pN) };
+ uint grid = (uint)((channels + 255) / 256);
+ int r = NvrtcInterop.LaunchKernel(_gdnConv1dStateUpdateBatchedKernel, grid, 1, 1, 256, 1, 1, 0, _stream, args, null);
+ if (r != 0) throw new InvalidOperationException($"cuLaunchKernel(gdn_conv1d_state_update_batched) failed: {r}");
+ }
+
+ ///
+ /// Batched per-head L2-norm over tokens.
+ /// is offset to the region base; is the per-token element
+ /// stride. grid = (numHeads, nTok). Bit-identical to nTok sequential
+ /// calls.
+ ///
+ public void GdnL2NormPerHeadBatched(Tensor data, long elementOffset, int numHeads, int headDim,
+ int rowStride, int nTok, float eps = 1e-6f)
+ {
+ EnsureImageKernels();
+ if (!_imageKernelsAvailable)
+ throw new NotSupportedException("NVRTC kernels are not available.");
+
+ nint dP = GetDevPtr(data) + (nint)(elementOffset * sizeof(float));
+ int pHD = headDim, pNH = numHeads, pRS = rowStride, pN = nTok;
+ float pE = eps;
+ nint* args = stackalloc nint[6]
+ {
+ (nint)(&dP), (nint)(&pHD), (nint)(&pNH), (nint)(&pE), (nint)(&pRS), (nint)(&pN)
+ };
+ int r = NvrtcInterop.LaunchKernel(_gdnL2NormPerHeadBatchedKernel, (uint)numHeads, (uint)nTok, 1, 256, 1, 1, 0, _stream, args, null);
+ if (r != 0) throw new InvalidOperationException($"cuLaunchKernel(gdn_l2_norm_per_head_batched) failed: {r}");
+ }
+
+ ///
+ /// Batched GQA-broadcast tile over tokens.
+ /// is offset to the region base; /
+ /// are per-token strides. Bit-identical to nTok sequential calls.
+ ///
+ public void GdnTileHeadsBatched(Tensor src, long srcOffset, Tensor dst, long dstOffset,
+ int srcHeads, int repeat, int headDim,
+ int srcStride, int dstStride, int nTok)
+ {
+ EnsureImageKernels();
+ if (!_imageKernelsAvailable)
+ throw new NotSupportedException("NVRTC kernels are not available.");
+
+ nint sP = GetDevPtr(src) + (nint)(srcOffset * sizeof(float));
+ nint dP = GetDevPtr(dst) + (nint)(dstOffset * sizeof(float));
+ int pSH = srcHeads, pR = repeat, pHD = headDim, pSS = srcStride, pDS = dstStride, pN = nTok;
+ nint* args = stackalloc nint[8]
+ {
+ (nint)(&sP), (nint)(&dP), (nint)(&pSH), (nint)(&pR), (nint)(&pHD),
+ (nint)(&pSS), (nint)(&pDS), (nint)(&pN)
+ };
+ int total = srcHeads * repeat * headDim;
+ uint grid = (uint)((total + 255) / 256);
+ int r = NvrtcInterop.LaunchKernel(_gdnTileHeadsBatchedKernel, grid, (uint)nTok, 1, 256, 1, 1, 0, _stream, args, null);
+ if (r != 0) throw new InvalidOperationException($"cuLaunchKernel(gdn_tile_heads_batched) failed: {r}");
+ }
+
+ ///
+ /// Fused sequential GDN recurrence scan over tokens: one launch
+ /// (one block per v-head) loops the positions internally, carrying the per-head state in
+ /// place. Bit-identical to nTok sequential calls — the
+ /// fused form of the per-token decode, NOT the parallel chunked-scan. Per-head input
+ /// strides let q/k come from the tiled [nTok × valueDim] buffers, v from the
+ /// silu'd conv output ( into a [nTok × convChannels]
+ /// buffer), z from a [nTok × valueDim] gate, alpha/beta from [nTok × numVHeads].
+ ///
+ public void GdnRecurrenceScan(
+ Tensor state, Tensor qAll, Tensor kAll, Tensor vAll,
+ Tensor alphaAll, Tensor betaAll, Tensor ssmA, Tensor dtBias,
+ Tensor normWeight, Tensor zAll, Tensor outputAll,
+ int numVHeads, int headDim, float normEps,
+ int qStride, int kStride, int vStride, int vHeadOff, int zStride, int oStride, int nTok)
+ {
+ EnsureImageKernels();
+ if (!_imageKernelsAvailable)
+ throw new NotSupportedException("NVRTC kernels are not available.");
+
+ nint sP = GetDevPtr(state);
+ nint qP = GetDevPtr(qAll), kP = GetDevPtr(kAll), vP = GetDevPtr(vAll);
+ nint aP = GetDevPtr(alphaAll), bP = GetDevPtr(betaAll);
+ nint aaP = GetDevPtr(ssmA), dbP = GetDevPtr(dtBias), nwP = GetDevPtr(normWeight);
+ nint zP = GetDevPtr(zAll), oP = GetDevPtr(outputAll);
+ int pHV = numVHeads, pD = headDim;
+ float pE = normEps;
+ int pQS = qStride, pKS = kStride, pVS = vStride, pVO = vHeadOff, pZS = zStride, pOS = oStride, pN = nTok;
+ nint* args = stackalloc nint[21]
+ {
+ (nint)(&sP), (nint)(&qP), (nint)(&kP), (nint)(&vP),
+ (nint)(&aP), (nint)(&bP), (nint)(&aaP), (nint)(&dbP),
+ (nint)(&nwP), (nint)(&zP), (nint)(&oP),
+ (nint)(&pHV), (nint)(&pD), (nint)(&pE),
+ (nint)(&pQS), (nint)(&pKS), (nint)(&pVS), (nint)(&pVO), (nint)(&pZS), (nint)(&pOS), (nint)(&pN)
+ };
+ uint sharedBytes = (uint)(8 * headDim * sizeof(float));
+ int r = NvrtcInterop.LaunchKernel(_gdnRecurrenceScanKernel,
+ (uint)numVHeads, 1, 1, (uint)headDim, 1, 1, sharedBytes, _stream, args, null);
+ if (r != 0) throw new InvalidOperationException($"cuLaunchKernel(gdn_recurrence_scan) failed: {r}");
+ }
+
public void FullSeqAttention(Tensor output, Tensor q, Tensor k, Tensor v,
int nTok, int nHeads, int headDim, float scale) =>
throw new NotSupportedException("CudaBackend.FullSeqAttention is not implemented (LLM path uses single-token Attention).");
@@ -3063,6 +3319,10 @@ private void ForceEagerJit()
_tqRotateQueryKernel, _tqKvAppendKernel, _tqAttentionKernel,
_siluInplaceKernel, _gdnConv1dDecodeKernel, _gdnL2NormPerHeadKernel,
_gdnTileHeadsKernel, _gdnRecurrenceDecodeKernel,
+ _gdnConv1dDecodeBatchedKernel, _gdnConv1dStateUpdateBatchedKernel,
+ _gdnL2NormPerHeadBatchedKernel, _gdnTileHeadsBatchedKernel, _gdnRecurrenceScanKernel,
+ _kvAppendBatchedKernel, _kvAppendBatchedBf16Kernel,
+ _fullSeqAttentionKernel, _fullSeqAttentionBf16Kernel,
];
foreach (nint k in kernels)
{
@@ -3144,6 +3404,17 @@ private void LoadKernelFunctions()
_gdnL2NormPerHeadKernel = GetKernelFunc("llm_gdn_l2_norm_per_head");
_gdnTileHeadsKernel = GetKernelFunc("llm_gdn_tile_heads");
_gdnRecurrenceDecodeKernel = GetKernelFunc("llm_gdn_recurrence_decode");
+
+ // qwen35moe GDN batched trunk + batched-query SDPA kernels (issue #114-B).
+ _gdnConv1dDecodeBatchedKernel = GetKernelFunc("llm_gdn_conv1d_decode_batched");
+ _gdnConv1dStateUpdateBatchedKernel = GetKernelFunc("llm_gdn_conv1d_state_update_batched");
+ _gdnL2NormPerHeadBatchedKernel = GetKernelFunc("llm_gdn_l2_norm_per_head_batched");
+ _gdnTileHeadsBatchedKernel = GetKernelFunc("llm_gdn_tile_heads_batched");
+ _gdnRecurrenceScanKernel = GetKernelFunc("llm_gdn_recurrence_scan");
+ _kvAppendBatchedKernel = GetKernelFunc("llm_kv_append_batched");
+ _kvAppendBatchedBf16Kernel = GetKernelFunc("llm_kv_append_batched_bf16");
+ _fullSeqAttentionKernel = GetKernelFunc("llm_full_seq_attention");
+ _fullSeqAttentionBf16Kernel = GetKernelFunc("llm_full_seq_attention_bf16");
}
///
diff --git a/src/SharpInference.Cuda/CudaTextKernels.cs b/src/SharpInference.Cuda/CudaTextKernels.cs
index 4548f01..58bc245 100644
--- a/src/SharpInference.Cuda/CudaTextKernels.cs
+++ b/src/SharpInference.Cuda/CudaTextKernels.cs
@@ -2958,5 +2958,430 @@ __device__ __forceinline__ int sharpi_int8_at(const unsigned int* __restrict__ b
output[hd_off + j] = o_normed * silu;
}
+
+// ════════════════════════════════════════════════════════════════════════════
+// Issue #114-B: batched GDN trunk kernels — collapse the per-position decode
+// launches into one launch each over all N prompt tokens. Every kernel runs the
+// IDENTICAL per-row / per-position arithmetic (and reduction order) as its
+// single-token counterpart above, so the batched trunk is BIT-IDENTICAL to the
+// per-token TrunkLayerSequential path. The recurrence/conv state evolves exactly
+// as the sequential loop; only the host launch overhead is removed.
+// ════════════════════════════════════════════════════════════════════════════
+
+// ── GDN: depthwise causal conv1d over a chunk (read-only state) ─────────────
+// Bit-identical to N sequential `llm_gdn_conv1d_decode` calls. Each (channel,
+// token) thread computes output[i,c] reading the chunk inputs + the carried
+// pre-chunk state (oldest-first [(K-1), channels]). Does NOT mutate state — the
+// state advance is a separate launch (`llm_gdn_conv1d_state_update_batched`) so
+// concurrent token blocks all read the same old state. Sum order matches the
+// single-token kernel exactly: current tap (weight[K-1]) first, then taps
+// k=0..K-2 oldest→newest.
+extern ""C"" __global__ void llm_gdn_conv1d_decode_batched(
+ const float* __restrict__ x, // [n_tok, channels]
+ const float* __restrict__ state, // [(K-1), channels] oldest-first, pre-chunk
+ const float* __restrict__ weight, // [K, channels]
+ float* __restrict__ output, // [n_tok, channels]
+ int channels, int kernel_size, int n_tok)
+{
+ int c = (int)(blockIdx.x * blockDim.x + threadIdx.x);
+ int i = (int)blockIdx.y;
+ if (c >= channels || i >= n_tok) return;
+
+ int retained = kernel_size - 1;
+ float x_c = x[(long)i * channels + c];
+ float sum = weight[(long)retained * channels + c] * x_c;
+ #pragma unroll 4
+ for (int k = 0; k < retained; k++) {
+ int p = i - retained + k; // chunk-relative position of tap k
+ float val = (p >= 0)
+ ? x[(long)p * channels + c]
+ : state[(long)(p + retained) * channels + c];
+ sum += weight[(long)k * channels + c] * val;
+ }
+ output[(long)i * channels + c] = sum;
+}
+
+// ── GDN: advance conv1d state past a chunk ──────────────────────────────────
+// After the sequential loop processes all n_tok tokens, the retained state holds
+// the last (K-1) inputs oldest-first. Reproduces that exactly: new_state[r] is the
+// chunk input at position (n_tok-(K-1)+r), or the carried old state when that index
+// is still before the chunk (n_tok < K-1). All sources are read into registers
+// before any write to tolerate the in-place aliasing of the small-N case.
+extern ""C"" __global__ void llm_gdn_conv1d_state_update_batched(
+ const float* __restrict__ x, // [n_tok, channels]
+ float* __restrict__ state, // [(K-1), channels] in/out
+ int channels, int kernel_size, int n_tok)
+{
+ int c = (int)(blockIdx.x * blockDim.x + threadIdx.x);
+ if (c >= channels) return;
+
+ int retained = kernel_size - 1;
+ float tmp[4]; // K-1 <= 4 for our models
+ #pragma unroll 4
+ for (int r = 0; r < retained; r++) {
+ int p = n_tok - retained + r;
+ tmp[r] = (p >= 0)
+ ? x[(long)p * channels + c]
+ : state[(long)(p + retained) * channels + c];
+ }
+ #pragma unroll 4
+ for (int r = 0; r < retained; r++)
+ state[(long)r * channels + c] = tmp[r];
+}
+
+// ── GDN: L2-norm per head, batched over n_tok rows ──────────────────────────
+// Bit-identical to n_tok sequential `llm_gdn_l2_norm_per_head` calls. grid =
+// (num_heads, n_tok); `data` is the region base (already offset to the Q or K
+// region by the host), `row_stride` is the per-token element stride (= conv
+// channels). One block per (head, token); same 256-thread tree reduction.
+extern ""C"" __global__ void llm_gdn_l2_norm_per_head_batched(
+ float* __restrict__ data,
+ int head_dim, int num_heads, float eps, int row_stride, int n_tok)
+{
+ __shared__ float sdata[256];
+ unsigned int tid = threadIdx.x;
+ unsigned int head = blockIdx.x;
+ int i = (int)blockIdx.y;
+ if ((int)head >= num_heads || i >= n_tok) return;
+
+ float* d = data + (long)i * row_stride + (long)head * head_dim;
+
+ float sum = 0.f;
+ for (int e = (int)tid; e < head_dim; e += 256) {
+ float v = d[e];
+ sum += v * v;
+ }
+ sdata[tid] = sum;
+ __syncthreads();
+ for (unsigned int s = 128; s > 0; s >>= 1) {
+ if (tid < s) sdata[tid] += sdata[tid + s];
+ __syncthreads();
+ }
+ float norm = sqrtf(sdata[0]);
+ float divisor = norm > eps ? norm : eps;
+ float inv_div = 1.0f / divisor;
+ for (int e = (int)tid; e < head_dim; e += 256)
+ d[e] *= inv_div;
+}
+
+// ── GDN: tile heads (GQA broadcast), batched over n_tok rows ────────────────
+// Bit-identical to n_tok sequential `llm_gdn_tile_heads` calls. `src` is the
+// region base (host-offset to Q or K region), `src_stride` the per-token source
+// stride (= conv channels), `dst_stride` the per-token dst stride (= value_dim).
+extern ""C"" __global__ void llm_gdn_tile_heads_batched(
+ const float* __restrict__ src,
+ float* __restrict__ dst,
+ int src_heads, int repeat, int head_dim,
+ int src_stride, int dst_stride, int n_tok)
+{
+ int dst_total = src_heads * repeat * head_dim;
+ int idx = (int)(blockIdx.x * blockDim.x + threadIdx.x);
+ int i = (int)blockIdx.y;
+ if (idx >= dst_total || i >= n_tok) return;
+ int j = idx % head_dim;
+ int h_dst = idx / head_dim;
+ int h_src = h_dst % src_heads;
+ dst[(long)i * dst_stride + idx] = src[(long)i * src_stride + (long)h_src * head_dim + j];
+}
+
+// ── GDN: fused sequential recurrence scan over a chunk ───────────────────────
+// One block per v-head, blockDim = head_dim. Loops the n_tok positions INTERNALLY,
+// running the exact passes of `llm_gdn_recurrence_decode` at each step and carrying
+// the per-head state matrix in global memory between steps. This is the bit-identical
+// fused form of N sequential `llm_gdn_recurrence_decode` launches — NOT the parallel
+// chunked-scan (which reorders the FP reductions). Each thread owns output column j
+// of its head's state, so the only cross-thread sharing is via shared memory with
+// the same __syncthreads barriers as the single-token kernel; the trailing barrier
+// makes the position boundary clean before the next step reloads shared inputs.
+//
+// Per-head input strides let q/k come from the tiled [n_tok, value_dim] buffers,
+// v straight from the silu'd conv output (v_head_off into the [n_tok, conv_ch]
+// buffer), z from the [n_tok, value_dim] gate, alpha/beta from [n_tok, num_v_heads].
+extern ""C"" __global__ void llm_gdn_recurrence_scan(
+ float* __restrict__ state, // [hv, d, d]
+ const float* __restrict__ q_all, // [n_tok, q_stride]; head h at h*d
+ const float* __restrict__ k_all, // [n_tok, k_stride]
+ const float* __restrict__ v_all, // [n_tok, v_stride]; head h at v_head_off + h*d
+ const float* __restrict__ alpha_all, // [n_tok, hv]
+ const float* __restrict__ beta_all, // [n_tok, hv]
+ const float* __restrict__ ssm_a, // [hv]
+ const float* __restrict__ dt_bias, // [hv]
+ const float* __restrict__ norm_weight,// [d]
+ const float* __restrict__ z_all, // [n_tok, z_stride]; head h at h*d
+ float* __restrict__ output_all, // [n_tok, o_stride]; head h at h*d
+ int hv, int d, float norm_eps,
+ int q_stride, int k_stride, int v_stride, int v_head_off,
+ int z_stride, int o_stride, int n_tok)
+{
+ int h = (int)blockIdx.x;
+ int j = (int)threadIdx.x;
+ if (h >= hv || j >= d) return;
+
+ extern __shared__ float smem[];
+ float* sK = smem;
+ float* sQ = sK + d;
+ float* sV = sQ + d;
+ float* sZ = sV + d;
+ float* sNormW = sZ + d;
+ float* sP = sNormW + d;
+ float* sD = sP + d;
+ float* sRed = sD + d;
+
+ sNormW[j] = norm_weight[j]; // layer-constant; each thread reads own j
+ long state_base = (long)h * (long)d * (long)d;
+
+ for (int i = 0; i < n_tok; i++) {
+ long qoff = (long)i * q_stride + (long)h * d;
+ long koff = (long)i * k_stride + (long)h * d;
+ long voff = (long)i * v_stride + v_head_off + (long)h * d;
+ long zoff = (long)i * z_stride + (long)h * d;
+ sK[j] = k_all[koff + j];
+ sQ[j] = q_all[qoff + j];
+ sV[j] = v_all[voff + j];
+ sZ[j] = z_all[zoff + j];
+ __syncthreads();
+
+ float alpha_x = alpha_all[(long)i * hv + h] + dt_bias[h];
+ float dt = alpha_x >= 20.0f ? alpha_x : __logf(1.0f + __expf(alpha_x));
+ float decay = __expf(dt * ssm_a[h]);
+ float b_sc = 1.0f / (1.0f + __expf(-beta_all[(long)i * hv + h]));
+
+ // Pass A: decay S, accumulate p[j] = Σ_i k[i]·S[i,j].
+ float p_local = 0.f;
+ for (int ii = 0; ii < d; ii++) {
+ long off = state_base + (long)ii * d + j;
+ float sij = state[off] * decay;
+ state[off] = sij;
+ p_local += sK[ii] * sij;
+ }
+ sP[j] = p_local;
+ __syncthreads();
+
+ float d_j = b_sc * (sV[j] - sP[j]);
+ sD[j] = d_j;
+ __syncthreads();
+
+ // Pass B: rank-1 update S[i,j] += k[i]·d[j], fused with readout o[j].
+ float o_local = 0.f;
+ for (int ii = 0; ii < d; ii++) {
+ long off = state_base + (long)ii * d + j;
+ float sij = state[off] + sK[ii] * d_j;
+ state[off] = sij;
+ o_local += sQ[ii] * sij;
+ }
+ o_local *= rsqrtf((float)d);
+
+ sRed[j] = o_local * o_local;
+ __syncthreads();
+ for (int s = d / 2; s > 0; s >>= 1) {
+ if (j < s) sRed[j] += sRed[j + s];
+ __syncthreads();
+ }
+ float scale = rsqrtf(sRed[0] / (float)d + norm_eps);
+ float o_normed = o_local * scale * sNormW[j];
+
+ float zv = sZ[j];
+ float silu = zv / (1.0f + __expf(-zv));
+ output_all[(long)i * o_stride + (long)h * d + j] = o_normed * silu;
+ __syncthreads(); // position boundary: next step reloads shared
+ }
+}
+
+// ════════════════════════════════════════════════════════════════════════════
+// Issue #114-B: batched-query SDPA for prompt prefill.
+// ════════════════════════════════════════════════════════════════════════════
+
+// ── KV cache append, batched over n_tok tokens (fp32 store) ─────────────────
+// Bit-identical to n_tok sequential `llm_kv_append` calls at positions
+// start_pos+i. grid = ((kv_dim+255)/256, n_tok).
+extern ""C"" __global__ void llm_kv_append_batched(
+ const float* __restrict__ k_all, const float* __restrict__ v_all,
+ float* __restrict__ k_cache, float* __restrict__ v_cache,
+ int kv_dim, int start_pos, int max_seq_len, int n_tok)
+{
+ int e = (int)(blockIdx.x * blockDim.x + threadIdx.x);
+ int i = (int)blockIdx.y;
+ if (e >= kv_dim || i >= n_tok) return;
+ long off = (long)(start_pos + i) * (long)kv_dim + (long)e;
+ k_cache[off] = k_all[(long)i * kv_dim + e];
+ v_cache[off] = v_all[(long)i * kv_dim + e];
+}
+
+// bf16-store variant (default KV dtype). Matches `llm_kv_append_bf16`.
+extern ""C"" __global__ void llm_kv_append_batched_bf16(
+ const float* __restrict__ k_all, const float* __restrict__ v_all,
+ unsigned short* __restrict__ k_cache, unsigned short* __restrict__ v_cache,
+ int kv_dim, int start_pos, int max_seq_len, int n_tok)
+{
+ int e = (int)(blockIdx.x * blockDim.x + threadIdx.x);
+ int i = (int)blockIdx.y;
+ if (e >= kv_dim || i >= n_tok) return;
+ long off = (long)(start_pos + i) * (long)kv_dim + (long)e;
+ k_cache[off] = (unsigned short)sharpi_fp32_to_bf16(k_all[(long)i * kv_dim + e]);
+ v_cache[off] = (unsigned short)sharpi_fp32_to_bf16(v_all[(long)i * kv_dim + e]);
+}
+
+// ── Full-sequence (batched-query) scaled dot-product attention ──────────────
+// Implements CudaBackend.FullSeqAttention for prompt prefill. grid = (num_heads,
+// n_tok): block (h, i) computes the attention output for query i (absolute
+// position start_pos+i) over the causal prefix [0, start_pos+i+1). Bit-identical
+// to n_tok sequential `llm_attention` calls (use_shared path) — same per-position
+// dot, same 256-thread tree softmax, same V-weighted sum. The host only dispatches
+// this when start_pos+n_tok ≤ MAX_STORED_SCORES, so every block stays on the
+// shared-scores fast path (no global scratch, no cross-block aliasing).
+extern ""C"" __global__ void llm_full_seq_attention(
+ const float* __restrict__ q_all, // [n_tok, num_heads*head_dim]
+ const float* __restrict__ k_cache,
+ const float* __restrict__ v_cache,
+ float* __restrict__ out_all, // [n_tok, num_heads*head_dim]
+ int num_heads, int num_kv_heads, int head_dim,
+ int start_pos, int max_seq_len, int n_tok)
+{
+ const int MAX_STORED_SCORES = 4096;
+ __shared__ float shared_scores[MAX_STORED_SCORES];
+ __shared__ float sdata[256];
+
+ unsigned int tid = threadIdx.x;
+ unsigned int h = blockIdx.x;
+ int i = (int)blockIdx.y;
+ if ((int)h >= num_heads || i >= n_tok) return;
+
+ int seq_len = start_pos + i + 1;
+ int kv_head = (int)h / (num_heads / num_kv_heads);
+ int kv_dim = num_kv_heads * head_dim;
+ float scale = rsqrtf((float)head_dim);
+ int q_dim = num_heads * head_dim;
+ const float* q = q_all + (long)i * q_dim;
+ float* out = out_all + (long)i * q_dim;
+ long q_off = (long)h * (long)head_dim;
+ long out_off = q_off;
+
+ for (int t = (int)tid; t < seq_len; t += 256) {
+ float dot = 0.f;
+ long k_off = (long)t * (long)kv_dim + (long)kv_head * (long)head_dim;
+ for (int dd = 0; dd < head_dim; dd++)
+ dot += q[q_off + dd] * k_cache[k_off + dd];
+ shared_scores[t] = dot * scale;
+ }
+ __syncthreads();
+
+ float local_max = sharpi_neg_inf();
+ for (int t = (int)tid; t < seq_len; t += 256)
+ local_max = fmaxf(local_max, shared_scores[t]);
+ sdata[tid] = local_max;
+ __syncthreads();
+ for (unsigned int s = 128; s > 0; s >>= 1) {
+ if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid + s]);
+ __syncthreads();
+ }
+ float max_val = sdata[0];
+ __syncthreads();
+
+ float local_sum = 0.f;
+ for (int t = (int)tid; t < seq_len; t += 256) {
+ float ev = __expf(shared_scores[t] - max_val);
+ shared_scores[t] = ev;
+ local_sum += ev;
+ }
+ sdata[tid] = local_sum;
+ __syncthreads();
+ for (unsigned int s = 128; s > 0; s >>= 1) {
+ if (tid < s) sdata[tid] += sdata[tid + s];
+ __syncthreads();
+ }
+ float inv_sum = 1.0f / sdata[0];
+ __syncthreads();
+
+ for (int t = (int)tid; t < seq_len; t += 256)
+ shared_scores[t] *= inv_sum;
+ __syncthreads();
+
+ for (int dd = (int)tid; dd < head_dim; dd += 256) {
+ float acc = 0.f;
+ for (int t = 0; t < seq_len; t++) {
+ long v_off = (long)t * (long)kv_dim + (long)kv_head * (long)head_dim;
+ acc += shared_scores[t] * v_cache[v_off + dd];
+ }
+ out[out_off + dd] = acc;
+ }
+}
+
+// bf16-read variant (default KV dtype). Matches `llm_attention_bf16`'s use_shared path.
+extern ""C"" __global__ void llm_full_seq_attention_bf16(
+ const float* __restrict__ q_all,
+ const unsigned short* __restrict__ k_cache,
+ const unsigned short* __restrict__ v_cache,
+ float* __restrict__ out_all,
+ int num_heads, int num_kv_heads, int head_dim,
+ int start_pos, int max_seq_len, int n_tok)
+{
+ const int MAX_STORED_SCORES = 4096;
+ __shared__ float shared_scores[MAX_STORED_SCORES];
+ __shared__ float sdata[256];
+
+ unsigned int tid = threadIdx.x;
+ unsigned int h = blockIdx.x;
+ int i = (int)blockIdx.y;
+ if ((int)h >= num_heads || i >= n_tok) return;
+
+ int seq_len = start_pos + i + 1;
+ int kv_head = (int)h / (num_heads / num_kv_heads);
+ int kv_dim = num_kv_heads * head_dim;
+ float scale = rsqrtf((float)head_dim);
+ int q_dim = num_heads * head_dim;
+ const float* q = q_all + (long)i * q_dim;
+ float* out = out_all + (long)i * q_dim;
+ long q_off = (long)h * (long)head_dim;
+ long out_off = q_off;
+
+ for (int t = (int)tid; t < seq_len; t += 256) {
+ float dot = 0.f;
+ long k_off = (long)t * (long)kv_dim + (long)kv_head * (long)head_dim;
+ for (int dd = 0; dd < head_dim; dd++)
+ dot += q[q_off + dd] * sharpi_bf16_to_fp32((unsigned int)k_cache[k_off + dd]);
+ shared_scores[t] = dot * scale;
+ }
+ __syncthreads();
+
+ float local_max = sharpi_neg_inf();
+ for (int t = (int)tid; t < seq_len; t += 256)
+ local_max = fmaxf(local_max, shared_scores[t]);
+ sdata[tid] = local_max;
+ __syncthreads();
+ for (unsigned int s = 128; s > 0; s >>= 1) {
+ if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid + s]);
+ __syncthreads();
+ }
+ float max_val = sdata[0];
+ __syncthreads();
+
+ float local_sum = 0.f;
+ for (int t = (int)tid; t < seq_len; t += 256) {
+ float ev = __expf(shared_scores[t] - max_val);
+ shared_scores[t] = ev;
+ local_sum += ev;
+ }
+ sdata[tid] = local_sum;
+ __syncthreads();
+ for (unsigned int s = 128; s > 0; s >>= 1) {
+ if (tid < s) sdata[tid] += sdata[tid + s];
+ __syncthreads();
+ }
+ float inv_sum = 1.0f / sdata[0];
+ __syncthreads();
+
+ for (int t = (int)tid; t < seq_len; t += 256)
+ shared_scores[t] *= inv_sum;
+ __syncthreads();
+
+ for (int dd = (int)tid; dd < head_dim; dd += 256) {
+ float acc = 0.f;
+ for (int t = 0; t < seq_len; t++) {
+ long v_off = (long)t * (long)kv_dim + (long)kv_head * (long)head_dim;
+ acc += shared_scores[t] * sharpi_bf16_to_fp32((unsigned int)v_cache[v_off + dd]);
+ }
+ out[out_off + dd] = acc;
+ }
+}
";
}
diff --git a/src/SharpInference.Engine/CudaHybridGdnForwardPass.cs b/src/SharpInference.Engine/CudaHybridGdnForwardPass.cs
index e9cc8b6..2962482 100644
--- a/src/SharpInference.Engine/CudaHybridGdnForwardPass.cs
+++ b/src/SharpInference.Engine/CudaHybridGdnForwardPass.cs
@@ -516,6 +516,29 @@ private void ThrowIfFaulted()
private Tensor? _gpuBtAttnOut; // [N × qDim] attn output (pre-O)
private Tensor? _gpuBtSGate; // [N × expertDim] shared-expert gate
private Tensor? _gpuBtSUp; // [N × expertDim] shared-expert up
+ // Issue #114-B fused-GDN-scan scratch (only allocated when BatchedGdnScanEnabled).
+ private Tensor? _gpuBtQkvConv; // [N × convChannels] post-conv1d + SiLU
+ private Tensor? _gpuBtQHead; // [N × valueDim] tiled GDN query heads
+ private Tensor? _gpuBtKHead; // [N × valueDim] tiled GDN key heads
+
+ // Issue #114-B: fuse the per-position GDN recurrence launches into one
+ // sequential-scan kernel + batched conv1d/L2norm/tile/silu over N. Default on;
+ // SHARPI_BATCHED_GDN_SCAN=0 falls back to the per-position View loop inside
+ // GdnBlockBatched (the pre-#114-B path). A/B-toggleable like SHARPI_BATCHED_TRUNK.
+ internal static bool BatchedGdnScanEnabled =
+ Environment.GetEnvironmentVariable("SHARPI_BATCHED_GDN_SCAN") != "0";
+
+ // Issue #114-B: batch the per-position KV-append + SDPA into one launch each
+ // (CudaBackend.AttentionBatched). Only used when the chunk stays on the
+ // shared-scores path (startPos+N ≤ 4096) and SnapKV Q-capture is inactive;
+ // otherwise AttnBlockBatched keeps the per-position loop. Default on;
+ // SHARPI_BATCHED_ATTN=0 forces the per-position path.
+ internal static bool BatchedAttnEnabled =
+ Environment.GetEnvironmentVariable("SHARPI_BATCHED_ATTN") != "0";
+
+ // One-shot guard so the >4096 batched-SDPA perf fallback is observable (it is
+ // bit-identical, just slower per-position) without spamming once per layer/chunk.
+ private bool _loggedBatchedAttnFallback;
public int VocabSize => _hp.VocabSize;
public int MaxSeqLen => _maxSeqLen;
@@ -1417,6 +1440,12 @@ private void EnsureBatchedTrunkScratch(int N)
_gpuBtAttnOut = A((long)N * qDim);
_gpuBtSGate = A((long)N * _expertDim);
_gpuBtSUp = A((long)N * _expertDim);
+ if (BatchedGdnScanEnabled)
+ {
+ _gpuBtQkvConv = A((long)N * _gdnConvChannels);
+ _gpuBtQHead = A((long)N * _gdnValueDim);
+ _gpuBtKHead = A((long)N * _gdnValueDim);
+ }
_btCap = N;
}
@@ -1427,6 +1456,7 @@ private void FreeBatchedTrunkScratch()
F(ref _gpuBtQkv); F(ref _gpuBtZ); F(ref _gpuBtAlpha); F(ref _gpuBtBeta); F(ref _gpuBtGdnOut);
F(ref _gpuBtQGate); F(ref _gpuBtQ); F(ref _gpuBtGate); F(ref _gpuBtK); F(ref _gpuBtV); F(ref _gpuBtAttnOut);
F(ref _gpuBtSGate); F(ref _gpuBtSUp);
+ F(ref _gpuBtQkvConv); F(ref _gpuBtQHead); F(ref _gpuBtKHead);
_btCap = 0;
}
@@ -1481,8 +1511,12 @@ private void TrunkLayerSequential(int layer, int N, int startPos, bool isAttn,
/// plus the attn-norm / post-attn-norm RmsNorms, per-head Q/K RMSNorm, RoPE,
/// Q‖gate split and GLU gate run as single batched launches over all N tokens.
/// The conv1d / delta-net recurrence and KV-append / scaled-dot-product attention
- /// stay per-position (they are positional), reading per-token slices via
- /// . Output is bit-identical to
+ /// were per-position over N (per-token slices via )
+ /// in #111; issue #114-B now batches them too by default — a fused sequential-scan
+ /// for the GDN recurrence and a batched-query SDPA — falling back to the per-position
+ /// View loops under SHARPI_BATCHED_GDN_SCAN=0 / SHARPI_BATCHED_ATTN=0
+ /// (and, for attention, when the chunk exceeds the shared-scores window or SnapKV is
+ /// active). Output is bit-identical to
/// : every batched kernel runs the same per-row /
/// per-element computation as its single-token counterpart, only collapsing the
/// N per-token launches into one each — which is what removes the host launch
@@ -1527,7 +1561,9 @@ private void TrunkLayerBatched(int layer, int N, int startPos, bool isAttn,
_gpu.Synchronize(); // drain all queued D2H: resid/norm/shared now host-valid
}
- /// Batched GDN block: projections over N tokens, per-position recurrence.
+ /// Batched GDN block: projections over N tokens; fused sequential-scan
+ /// recurrence + batched conv1d/L2norm/tile by default (issue #114-B), or the
+ /// per-position View loop under SHARPI_BATCHED_GDN_SCAN=0.
private void GdnBlockBatched(int layer, int N, Tensor norm, Tensor blockOut)
{
int convCh = _gdnConvChannels, valDim = _gdnValueDim, nVH = _gdnNumVHeads;
@@ -1544,6 +1580,42 @@ private void GdnBlockBatched(int layer, int N, Tensor norm, Tensor blockOut)
var scanState = _gpuGdnScanState[layer]!;
var convState = _gpuGdnConvState[layer]!;
+ // Issue #114-B: fuse the per-position conv1d + delta-net recurrence into
+ // one batched launch per stage + a single sequential-scan kernel. Output is
+ // bit-identical to the per-position View loop below (same per-position math,
+ // same reduction order; only the host launch overhead is removed).
+ if (BatchedGdnScanEnabled)
+ {
+ var qkvConvAll = _gpuBtQkvConv!;
+ var qHeadAll = _gpuBtQHead!;
+ var kHeadAll = _gpuBtKHead!;
+
+ // conv1d over all tokens (read-only state), then advance the state.
+ _gpu.GdnConv1dDecodeBatched(qkvAll, convState, _gpuSsmConv1d[layer], qkvConvAll,
+ convCh, _gdnConvKernel, N);
+ _gpu.GdnConv1dStateUpdateBatched(qkvAll, convState, convCh, _gdnConvKernel, N);
+ // SiLU over the whole [N × convCh] (matches the per-token full-convCh SiLU).
+ _gpu.SiLUInPlace(qkvConvAll);
+ // L2-norm the Q (offset 0) and K (offset kDim) regions, per head, per token.
+ _gpu.GdnL2NormPerHeadBatched(qkvConvAll, 0, _gdnNumKHeads, hd, convCh, N, eps: 1e-6f);
+ _gpu.GdnL2NormPerHeadBatched(qkvConvAll, kDim, _gdnNumKHeads, hd, convCh, N, eps: 1e-6f);
+ // Tile Q and K heads (GQA broadcast) into the [N × valueDim] head buffers.
+ _gpu.GdnTileHeadsBatched(qkvConvAll, 0, qHeadAll, 0, _gdnNumKHeads, _gdnKvRepeat, hd, convCh, valDim, N);
+ _gpu.GdnTileHeadsBatched(qkvConvAll, kDim, kHeadAll, 0, _gdnNumKHeads, _gdnKvRepeat, hd, convCh, valDim, N);
+ // Fused sequential scan: v read straight from the silu'd conv output's V
+ // region (vHeadOff = 2*kDim, stride convCh); q/k from the tiled head buffers.
+ _gpu.GdnRecurrenceScan(
+ scanState, qHeadAll, kHeadAll, qkvConvAll,
+ alphaAll, betaAll, _gpuSsmA[layer], _gpuSsmDtBias[layer], _gpuSsmNormW[layer],
+ zAll, gdnOutAll,
+ nVH, hd, normEps: 1e-6f,
+ qStride: valDim, kStride: valDim, vStride: convCh, vHeadOff: 2 * kDim,
+ zStride: valDim, oStride: valDim, nTok: N);
+
+ GpuMatMulBatched(blockOut, _gpuWSsmOut[layer], gdnOutAll, N);
+ return;
+ }
+
// Per-token conv1d + delta-net recurrence (positional → sequential). The
// conv/L2/tile scratch (_gpuGdnQkvConv / _gpuGdnQHead / _gpuGdnKHead /
// _gpuGdnVHead) is reused per token; only the batched-buffer inputs/output
@@ -1581,8 +1653,10 @@ private void GdnBlockBatched(int layer, int N, Tensor norm, Tensor blockOut)
GpuMatMulBatched(blockOut, _gpuWSsmOut[layer], gdnOutAll, N);
}
- /// Batched attention block: projections + Q/K norm + RoPE over N tokens,
- /// per-position KV-append + SDPA.
+ /// Batched attention block: projections + Q/K norm + RoPE over N tokens;
+ /// batched KV-append + batched-query SDPA by default (issue #114-B) when the chunk
+ /// stays on the shared-scores window (startPos+N ≤ 4096) and SnapKV is inactive,
+ /// else (or under SHARPI_BATCHED_ATTN=0) the per-position KV-append + SDPA loop.
private void AttnBlockBatched(int layer, int N, int startPos, bool snapKvActive, int wStart,
Tensor norm, Tensor blockOut)
{
@@ -1601,6 +1675,43 @@ private void AttnBlockBatched(int layer, int N, int startPos, bool snapKvActive,
_gpu.RoPEPartialBatched(qAll, startPos, _headDim, _ropeDim, _hp.RopeTheta, _numHeads, N, neox: true);
_gpu.RoPEPartialBatched(kAll, startPos, _headDim, _ropeDim, _hp.RopeTheta, _numKvHeads, N, neox: true);
+ // Issue #114-B: batch the KV-append + SDPA into one launch each when the
+ // chunk stays on the shared-scores fast path (startPos+N ≤ 4096) and SnapKV
+ // Q-capture is inactive. Bit-identical to the per-position loop below (the
+ // batched attention kernel clones llm_attention per (head, query) block).
+ // SnapKV active or long context → fall through to the per-position loop,
+ // which captures Q and supports the >4096 global-scratch attention path.
+ if (BatchedAttnEnabled && !snapKvActive && startPos + N <= 4096)
+ {
+ if (_kvDType == DType.BFloat16)
+ {
+ _gpu.KvAppendBatchedBf16(kAll, vAll, _gpuKCache[layer]!, _gpuVCache[layer]!, kvDim, startPos, _maxSeqLen, N);
+ _gpu.AttentionBatchedBf16(qAll, _gpuKCache[layer]!, _gpuVCache[layer]!, attnOutAll,
+ _numHeads, _numKvHeads, _headDim, startPos, _maxSeqLen, N);
+ }
+ else
+ {
+ _gpu.KvAppendBatched(kAll, vAll, _gpuKCache[layer]!, _gpuVCache[layer]!, kvDim, startPos, _maxSeqLen, N);
+ _gpu.AttentionBatched(qAll, _gpuKCache[layer]!, _gpuVCache[layer]!, attnOutAll,
+ _numHeads, _numKvHeads, _headDim, startPos, _maxSeqLen, N);
+ }
+ _gpu.SigmoidMulInPlace(attnOutAll, gateAll);
+ GpuMatMulBatched(blockOut, _gpuWO[layer], attnOutAll, N);
+ return;
+ }
+
+ // Falling back to the per-position attention loop. Bit-identical, but slower —
+ // surface the >4096 perf cliff once (the batched SDPA only covers the
+ // shared-scores window; SnapKV-active prefill also takes this path by design).
+ if (BatchedAttnEnabled && !snapKvActive && startPos + N > 4096 && !_loggedBatchedAttnFallback)
+ {
+ _loggedBatchedAttnFallback = true;
+ Console.Error.WriteLine(
+ $"[CudaHybridGdnForwardPass] batched-query SDPA disabled for this prefill: " +
+ $"startPos+N={startPos + N} exceeds the {4096}-position shared-scores window; " +
+ "using the per-position attention path (bit-identical, slower). #114-B follow-up: wave-based >4096 SDPA.");
+ }
+
// Per-position KV-append + scaled-dot-product attention (positional →
// sequential), plus SnapKV Q-capture for the trailing window.
for (int i = 0; i < N; i++)
diff --git a/tests/SharpInference.Tests.ForwardPass/CudaGdnBatchedTrunkTests.cs b/tests/SharpInference.Tests.ForwardPass/CudaGdnBatchedTrunkTests.cs
new file mode 100644
index 0000000..a4e4a93
--- /dev/null
+++ b/tests/SharpInference.Tests.ForwardPass/CudaGdnBatchedTrunkTests.cs
@@ -0,0 +1,393 @@
+using SharpInference.Core;
+using SharpInference.Cuda;
+
+namespace SharpInference.Tests.ForwardPass;
+
+///
+/// Issue #114-B per-kernel bit-exactness tests for the batched GDN trunk +
+/// batched-query SDPA kernels. Each batched kernel must produce results
+/// bit-identical to the N sequential single-token kernels it replaces —
+/// not just within tolerance. A divergence means a reduction was reordered (the
+/// failure mode the GDN/MTP byte-parity oracles trip on — see the K/V MatVecDual
+/// note). Mirrors ; silently skips without CUDA.
+///
+public sealed unsafe class CudaGdnBatchedTrunkTests
+{
+ private static CudaBackend? TryCreate()
+ {
+ if (!CudaBackend.IsAvailable()) return null;
+ try { return CudaBackend.Create(); }
+ catch { return null; }
+ }
+
+ private static float[] Rand(int n, Random rng)
+ {
+ var a = new float[n];
+ for (int i = 0; i < n; i++) a[i] = (float)(rng.NextDouble() * 2 - 1);
+ return a;
+ }
+
+ private static int Bits(float f) => BitConverter.SingleToInt32Bits(f);
+
+ private static void AssertBitId(string label, float[] batched, float[] reference)
+ {
+ Assert.Equal(reference.Length, batched.Length);
+ for (int i = 0; i < reference.Length; i++)
+ if (Bits(batched[i]) != Bits(reference[i]))
+ Assert.Fail($"{label}: index {i} batched={batched[i]} (0x{Bits(batched[i]):X8}) " +
+ $"!= sequential={reference[i]} (0x{Bits(reference[i]):X8}).");
+ }
+
+ // ── GDN conv1d (decode + state advance) ────────────────────────────────
+ [Fact]
+ public void GdnConv1dDecodeBatched_BitwiseMatchesSequential()
+ {
+ using var gpu = TryCreate();
+ if (gpu is null) return;
+
+ const int channels = 8192, kernelSize = 4, retained = kernelSize - 1;
+ // N=1 hits the maximal state-aliasing edge the state-update kernel guards against
+ // (all retained taps come from carried state); N
+ /// Issue #114-B: the batched GDN trunk's fused sequential-scan recurrence
+ /// (BatchedGdnScanEnabled) and batched-query SDPA
+ /// (BatchedAttnEnabled), both default-on, must be bit-identical to the
+ /// per-position View-loop fallback (the path taken when those flags are off — the
+ /// pre-#114-B reference). This is the only oracle that exercises the host glue in
+ /// GdnBlockBatched / AttnBlockBatched — the conv→silu→L2norm→tile→scan
+ /// strides/offsets and the KV-append/SDPA wiring — at the model level; the per-kernel
+ /// tests () prove each kernel in isolation but
+ /// not the argument plumbing that strings them together. Both arms run under batched
+ /// prefill + batched trunk (#110/#111); only the #114-B sub-paths differ.
+ ///
+ [Fact]
+ public void BatchedGdnScanAndAttn_BitwiseMatchesPerPosition_Carnice()
+ {
+ using var gpu = TryCreate();
+ if (gpu is null) return;
+ var path = FindCarnicePath();
+ if (path is null) return;
+
+ var prevCpuMoe = Environment.GetEnvironmentVariable("SHARPI_CPU_MOE");
+ Environment.SetEnvironmentVariable("SHARPI_CPU_MOE", "1");
+ bool prevBatchedPrefill = CudaHybridGdnForwardPass.BatchedPrefillEnabled;
+ bool prevBatchedTrunk = CudaHybridGdnForwardPass.BatchedTrunkEnabled;
+ bool prevScan = CudaHybridGdnForwardPass.BatchedGdnScanEnabled;
+ bool prevAttn = CudaHybridGdnForwardPass.BatchedAttnEnabled;
+ try
+ {
+ using var model = GgufModel.Open(path);
+ var hp = ModelHyperparams.FromGgufMetadata(model.Metadata, model);
+ if (!hp.IsMoE) return;
+ var tokenizer = GgufTokenizer.FromGgufModel(model);
+ var placement = new LayerPlacement(
+ GpuLayers: hp.NumLayers, CpuLayers: 0, GpuWeightBytes: 0, GpuKvBytes: 0,
+ RecommendedCtxSize: Math.Min(hp.ContextLength, 4096));
+ var tokens = tokenizer.Encode(
+ "The quick brown fox jumps over the lazy dog. " +
+ "Pack my box with five dozen liquor jugs. " +
+ "How razorback-jumping frogs can level six piqued gymnasts!");
+ Assert.True(tokens.Count >= 8);
+
+ // Hold the #110/#111 batched paths on; only toggle the #114-B sub-paths.
+ CudaHybridGdnForwardPass.BatchedPrefillEnabled = true;
+ CudaHybridGdnForwardPass.BatchedTrunkEnabled = true;
+
+ float[] RunWith(bool fused)
+ {
+ CudaHybridGdnForwardPass.BatchedGdnScanEnabled = fused;
+ CudaHybridGdnForwardPass.BatchedAttnEnabled = fused;
+ using var fwd = new CudaHybridGdnForwardPass(model, gpu, hp, placement);
+ return fwd.Prefill(tokens).ToArray();
+ }
+
+ float[] perPos = RunWith(false); // per-position recurrence + SDPA (View loops)
+ float[] fused = RunWith(true); // fused sequential-scan + batched-query SDPA
+
+ Assert.Equal(perPos.Length, fused.Length);
+ int firstDiff = -1;
+ for (int i = 0; i < perPos.Length; i++)
+ if (BitConverter.SingleToInt32Bits(perPos[i]) != BitConverter.SingleToInt32Bits(fused[i]))
+ { firstDiff = i; break; }
+ Assert.True(firstDiff < 0,
+ $"Fused GDN-scan + batched-query SDPA diverges from the per-position path at index {firstDiff}. " +
+ "GdnBlockBatched/AttnBlockBatched host wiring (strides/offsets) must be bit-identical to the View-loop fallback.");
+ Assert.Equal(Sampler.Greedy(perPos), Sampler.Greedy(fused));
+ }
+ finally
+ {
+ CudaHybridGdnForwardPass.BatchedPrefillEnabled = prevBatchedPrefill;
+ CudaHybridGdnForwardPass.BatchedTrunkEnabled = prevBatchedTrunk;
+ CudaHybridGdnForwardPass.BatchedGdnScanEnabled = prevScan;
+ CudaHybridGdnForwardPass.BatchedAttnEnabled = prevAttn;
+ Environment.SetEnvironmentVariable("SHARPI_CPU_MOE", prevCpuMoe);
+ }
+ }
+
///
/// Multi-chunk parity: prefill the prompt in two segments ([0,k) then
/// [k,N), the second with startPos=k) and assert the final-token