Dense Qwen3.5 hybrid architecture — PrismML Bonsai-27B runs end-to-end (#157) by jamesburton · Pull Request #158 · jamesburton/dotLLM

jamesburton · 2026-07-20T22:04:30Z

Summary

Adds Architecture.Qwen3HybridDense (qwen35) — a dense-FFN sibling of Qwen3MoeHybrid reusing the same GDN/full-attention token-mixing infrastructure, with dense SwiGLU FFN instead of MoE routing.
CPU (Qwen3HybridDenseTransformerModel) and CUDA (CudaQwen3HybridDenseTransformerModel) forward passes, GGUF config extraction, ModelLoader routing, and a CUDA dotllm bench dispatch path for both hybrid architectures.
Fixed a VRAM-oversubscription bug found validating on a 12 GB RTX 3060 (embedding lookup was pre-dequantizing the whole 248320-token vocab table; replaced with per-call host-side row dequant — silently tripped WDDM's host-RAM paging fallback, which looked exactly like a hang).
Overnight CUDA optimization sweep (profiler-guided, each step correctness-validated against a CPU reference before benchmarking):
1. Decode-time host loops (per-head cuMemcpyDtoDAsync splits) replaced with single gather-kernel launches.
2. PQ2_0 decode GEMV rewritten for shared-memory input staging + warp-per-row (mirrors I2_S's proven scheme).
3. Fixed a remaining coalescing gap in that same GEMV kernel — lanes were still striding across separate groups; restructured so the warp cooperates byte-by-byte within one group at a time.
4. Found and fixed the identical uncoalesced pattern in the prefill dequant kernel (never touched until now) — this was the single biggest win of the whole sweep.
5. Applied the same proven fix to I2_S's (BitNet) sibling dequant kernel, since it had the byte-for-byte identical defect (flagged as scope-adjacent to this PR, not end-to-end validated — no local BitNet GGUF this session).
6. Fused the residual-copy+RmsNorm decode pattern (called 2×/layer × 64 layers) into one kernel, halving that pair's launch count.
7. Fused the GDN decay+sigmoid decode step (called once each per GDN layer × 48 GDN layers) into one kernel — lives in the -fmad=false bit-perfect kernel file since it sits in the GDN recurrence's precision-critical path.
A PQ2_0_ROWS_PER_WARP=4 tuning experiment and a fused-projection GEMV round were tried and measured no further gain (documented as negative results, not blindly discarded).
Remaining decode-time gap surveyed via profiler after the two fusion rounds above: GEMM-bound categories (dense FFN gate+up/down, GDN projections) now dominate decode time (~55%+) and are already coalescing-optimized — further gains there would need a different GEMV algorithm, not launch-count reduction. The next non-GEMM candidate (GDN's causal conv1d input-prep, ~6% of decode) requires restructuring a -fmad=false precision-critical kernel for an estimated <1% gain and was judged not worth the added risk — deferred, not attempted.

Results

dotllm bench --device cuda -p 64 -n 16 -r 3 against the real 7.17 GB Ternary-Bonsai-27B-Q2_0.gguf on an RTX 3060:

	Prefill (tok/s)	Decode (tok/s)
Initial working baseline	9.86	1.71
After GEMV/dequant coalescing sweep	89.06	10.34
After decode-fusion follow-up (copy+RmsNorm, decay+sigmoid)	88.34	10.59
Speedup	9.0×	6.2×

Full history in benchmarks/perf-matrix/results.csv.

Test plan

GgufModelConfigExtractorTests — TDD unit test for qwen35 config extraction against real Bonsai metadata (34/34 pass).
Qwen3HybridDenseRealGgufConstructionTests — CPU loader construction against the real 7.17 GB GGUF.
CudaQwen3HybridDenseRealGgufSmokeTest — CUDA prefill + decode against the real GGUF on an RTX 3060, finite logits through all 64 layers.
New unit tests per kernel change, each against a CPU/separate-calls reference: CudaPQ2_0GemvTest (F16In GEMV incl. tail-clamp cases, fused-vs-separate parity, dequant bit-exact match), CudaI2SGemvTest (dequant bit-exact match), CudaQwen3MoeHybridElementwiseKernelTests (de-interleave gather kernels, GDN decay+sigmoid fusion — all exact/bit-exact match), CudaCopyRmsNormF32Test (fused copy+RmsNorm, bit-exact match).
Full CUDA test suite re-run clean after every change (301/301 passing on the latest round; one pre-existing, unrelated flaky test — CudaGraphCaptureEquivalenceTest, SmolLM-135M — confirmed passing in isolation multiple times across the sweep, not caused by this PR).
dotnet build clean (0 warnings, 0 errors) across all touched projects.
Real-weight benchmark after every optimization round — see table above.

🤖 Generated with Claude Code

…ai-27B (#157) Adds Architecture.Qwen3HybridDense ("qwen35") — a dense-FFN sibling of the existing Qwen3MoeHybrid family (same GDN/full-attention token-mixing infra, dense SwiGLU FFN instead of MoE routing). Implements both CPU and CUDA forward passes, wires GGUF config extraction and ModelLoader routing, and adds a CUDA bench-command dispatch path for both hybrid architectures (previously CUDA bench always fell through to the plain dense loader regardless of architecture). The CUDA dense hybrid model's Gemm dispatcher adds explicit I2_S/PQ2_0 branches missing from the MoE hybrid's own dispatcher (a separate, pre-existing gap, not fixed here) — required since Bonsai ships PQ2_0 end-to-end (GDN, attention, and FFN projections). Fixes a VRAM-oversubscription bug found while validating on the RTX 3060: the embedding lookup previously followed CudaQwen3MoeHybridTransformerModel's pattern of pre-dequantizing the full token embedding table to F32 when not natively F32/F16/Q8_0. At Bonsai's 248320-token vocab that's ~5 GB, which combined with the ~7.2 GB PQ2_0 weights and the ~2.5 GB prefill dequant scratch pushed total VRAM demand past the 3060's 12 GB budget and silently tripped WDDM's host-RAM paging fallback — manifesting as an apparent hang (confirmed via `Get-Counter '\GPU Process Memory(*)\Shared Usage'` showing multi-GB paged usage for the process). Replaced with a per-call host-side row dequant (only `seqLen` rows are ever needed, not all 248320). Validated end-to-end against the real 7.17 GB Ternary-Bonsai-27B-Q2_0.gguf: CPU construction test, CUDA real-weight smoke test (prefill + 4 decode steps, finite logits through all 64 layers), and a baseline `dotllm bench` run — prefill 9.86 tok/s, decode 1.71 tok/s (unoptimized, correctness-first PQ2_0 GEMV/dequant kernels; a target for follow-up optimization work). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Copilot

Pull request overview

Adds first-class support for a dense Qwen3.5 hybrid GGUF architecture (qwen35) across config extraction, CPU/CUDA loaders, and CLI benchmarking so PrismML Bonsai-27B can run end-to-end in dotLLM.

Changes:

Introduces Architecture.Qwen3HybridDense (qwen35) and GGUF config extraction for the GDN+attention hybrid layout without MoE.
Adds CPU + CUDA implementations of the hybrid-dense forward pass, loader routing, and CUDA dotllm bench dispatch for hybrid architectures.
Adds real-GGUF construction/smoke tests plus documentation/perf-matrix updates for the Bonsai-27B baseline.

Reviewed changes

Copilot reviewed 16 out of 16 changed files in this pull request and generated 6 comments.

Show a summary per file

File	Description
tests/DotLLM.Tests.Unit/Models/Gguf/GgufModelConfigExtractorTests.cs	Adds unit coverage for `qwen35` config extraction (GDN + hybrid layout, no MoE).
tests/DotLLM.Tests.Unit/Models/Architectures/Qwen3HybridDenseRealGgufConstructionTests.cs	Adds real-weight CPU construction smoke test for Bonsai-27B.
tests/DotLLM.Tests.Unit/Cuda/CudaQwen3HybridDenseRealGgufSmokeTest.cs	Adds CUDA prefill+decode finiteness smoke test for real Bonsai-27B GGUF.
src/DotLLM.Models/ModelLoader.cs	Routes CPU GGUF loading to the new hybrid-dense model.
src/DotLLM.Models/Gguf/GgufModelConfigExtractor.cs	Parses `qwen35` and extracts GDN/hybrid layout without MoE config.
src/DotLLM.Models/Architectures/Qwen3HybridDenseTransformerModel.cs	Implements CPU forward + GGUF weight loading for dense hybrid (GDN/attention + dense SwiGLU).
src/DotLLM.Models/Architectures/Qwen3HybridDenseLayerWeights.cs	Defines per-layer weight references for the dense hybrid block.
src/DotLLM.Models/Architectures/Qwen3HybridDenseForwardState.cs	Adds CPU scratch-buffer/state container for allocation-free forward execution.
src/DotLLM.Cuda/Architectures/CudaQwen3MoeHybridTransformerModel.cs	Adds `CreateKvCache` helper returning a CUDA hybrid KV-cache handle.
src/DotLLM.Cuda/Architectures/CudaQwen3HybridDenseTransformerModel.cs	Implements CUDA loader + forward for dense hybrid, including PQ2_0/I2_S GEMM dispatch.
src/DotLLM.Cuda/Architectures/CudaQwen3HybridDenseForwardState.cs	Adds CUDA per-token scratch buffer container for dense hybrid forward.
src/DotLLM.Cuda/Architectures/CudaHybridKvCacheHandle.cs	Introduces a length-only IKvCache handle for CUDA hybrid models.
src/DotLLM.Core/Configuration/Architecture.cs	Adds `Architecture.Qwen3HybridDense` enum entry + docs.
src/DotLLM.Cli/Commands/BenchCommand.cs	Fixes CUDA bench dispatch to route hybrid architectures to the correct loaders.
README.md	Adds News entry documenting Bonsai-27B end-to-end and PQ2_0 + qwen35 support.
benchmarks/perf-matrix/results.csv	Records baseline Bonsai-27B CUDA benchmark results row.

+        var embDesc = tensors["token_embd.weight"];
+        long embRowBytes = Dequantize.RowByteSize(hiddenSize, embDesc.QuantizationType);
+        long embTotalBytes = embRowBytes * config.VocabSize;
+        nint tokenEmbedDevice = AllocDevice(embTotalBytes);
+        CopyHtoD(tokenEmbedDevice, dataBase + (nint)embDesc.DataOffset, embTotalBytes);
+


+        float[] embedHost = new float[(long)seqLen * hiddenSize];
+        for (int t = 0; t < seqLen; t++)
+        {
+            nint rowSrc = _embedDataBase + (nint)(_embedDataOffset + (ulong)tokenIds[t] * (ulong)_embedRowBytes);
+            Dequantize.ToFloat32(rowSrc, hiddenSize, _tokenEmbedQt,
+                embedHost.AsSpan(t * hiddenSize, hiddenSize));
+        }
+        fixed (float* pEmbedHost = embedHost)
+        {
+            CudaDriverApi.cuMemcpyHtoDAsync_v2(_state.HiddenState, (nint)pEmbedHost,
+                (nuint)((long)seqLen * hiddenSize * sizeof(float)), streamH).ThrowOnError();
+        }


+                float alpha = alphaBuf[gbOff + vh] + gdnW.DtBias[vh];
+                float sp = MathF.Log(1f + MathF.Exp(alpha)); // softplus
+                alphaBuf[gbOff + vh] = MathF.Exp(sp * gdnW.A[vh]);


+        if (kvCache is not null)
+        {
+            EnsureF16KvCache(kvCache.MaxLength, numKvHeads, headDim);
+            int slot = _kvSlotForLayer[layer];


+            WriteF16KvRows(slot, k, v, positions, numKvHeads, headDim);
+
+            int positionOffset = positions[0];
+            int seqKv = _f16CacheCurrentLength;
+            int kvLiveElems = seqKv * kvElems;


+    /// <summary>
+    /// Creates a length-only <see cref="IKvCache"/> handle sized to <paramref name="maxSeqLen"/>.
+    /// K/V storage is owned internally by this model (a per-attention-layer F16 device
+    /// cache) — the returned handle only communicates the capacity to <see cref="Forward(System.ReadOnlySpan{int}, System.ReadOnlySpan{int}, int, IKvCache?)"/>.
+    /// </summary>
+    public CudaHybridKvCacheHandle CreateKvCache(int maxSeqLen) => new(maxSeqLen);
+


….15x) The PQ2_0 F16 GEMV was "correctness-first: one warp per output row, grid-strided over rows" — no shared-memory staging of the input vector (every warp re-read x from global) and a grid capped at 256 blocks (serializing large-n projections like the 17408-row FFN gate/up across many grid-stride iterations per warp). Measured 1.71 tok/s decode on real Bonsai-27B weights on an RTX 3060 — far below what a ~0.27-bytes/element ternary format should sustain. Rewrites pq2_0_gemv_f16in to mirror I2_S's proven v2 scheme (i2_s_gemv.cu): stage x into shared memory once per block and reuse it across PQ2_0_ROWS_PER_BLOCK=16 rows (8 warps x 2 rows/warp), uncapped grid (ceil(n/16) blocks). PQ2_0's bit-interleave turns out byte-for-byte identical to I2_S's, so the decode math ports directly — the only structural difference is PQ2_0's per-128-group FP16 scale (vs I2_S's single per-tensor scale), which stays scaled-then-accumulated per group exactly as before. The shared stage is kept half-width (not float, unlike I2_S) since Bonsai's largest K (17408, the FFN down-projection) at 4 bytes/float would exceed sm_86's 48 KB static shared cap; 17408 halfs is 34 KB and fits. Not yet ported: I2_S's vectorized uint4 coalesced weight loads — PQ2_0's 2-byte per-group scale prefix means a group's 32 code bytes begin at a generally-unaligned offset, unlike I2_S's fully contiguous rows. Flagged as a follow-up in the kernel's file header. The W2A8 dp4a path is deliberately not pursued either — I2_S's own dp4a variant measured slower than its F16 warp-per-row path on this hardware (CUDA_NOTES.md 2026-06-17), so it's not expected to help PQ2_0 either. Validated: new CudaPQ2_0GemvTest.PQ2_0GemvF16In_MatchesCpuFloatReference against the CPU float reference at real Bonsai shapes (n=512, k=5120/17408) plus tail-row-clamp edge cases (n=37, n=3) that exercise the block/warp boundary where a block covers more rows than exist. Full CUDA suite re-run clean (one pre-existing, unrelated flaky test confirmed to pass in isolation). Real-weight benchmark: decode 1.71 -> 7.09 tok/s (4.15x), prefill unchanged (9.86 -> 9.97, untouched dequant-then-cuBLAS path). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Adds pq2_0_gemv2_f16in — a fused decode-time GEMV for two PQ2_0 projections sharing one input, mirroring i2_s_gemv2_f16in (virtual row-concatenation of the two weight matrices, one shared-x stage instead of two). Wired into all three decode-time projection pairs that share an input in the dense hybrid model: dense FFN gate+up, full-attention K+V, and GDN alpha+beta (the GDN Qkv projection is already fused at the GGUF weight level, so no additional fusion opportunity there). Correctness validated: PQ2_0GemvFusedDecode_MatchesSeparateLaunches asserts the fused kernel matches two separate LaunchPQ2_0GemvF16In calls exactly, using odd/unequal n0=37,n1=53 to exercise the row-selection/tail-clamp boundary between the two virtually-concatenated matrices (mirroring the I2_S fused-decode test). Real-weight smoke test and full CUDA suite re-verified clean (one pre-existing, unrelated flaky CUDA-graph test confirmed passing on retry). Honest result: measured NO further decode throughput gain (7.09-7.10 tok/s across the shared-x-staging kernel alone and all three fusion additions, within run-to-run noise) — at batch=1 the actual GEMV compute now dominates the per-call shared-x-staging overhead that fusion targets, so this optimization's cost model doesn't apply here the way it does at I2_S's larger scale. Kept because it's zero-risk (correctness-validated, no regression) and may help at larger batch/seqLen where staging overhead is proportionally larger. Decode is no longer GEMV-bound; the next lever is likely kernel-launch-count reduction (CudaDecodeGraph.cs already exists in this codebase for CUDA graph capture) rather than further GEMV tuning. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Adds a lightweight, DOTLLM_HYBRID_PROFILE=1-gated category profiler (ProfStart/ProfMark, decode-only so prefill's much heavier dequant path doesn't mask the numbers) to CudaQwen3HybridDenseTransformerModel — used to find where decode time actually goes now that the GEMV kernel itself is optimized (previous commits). Finding: the attention Q+Gate de-interleave host loop (numHeads separate paired cuMemcpyDtoDAsync calls per layer) was ~12% of decode time — more than the attention kernel itself. Adds two gather kernels to elementwise_f32.cu: deinterleave_qgate_f32 (splits a fused Q+Gate projection into Q and Gate) and deinterleave_gdn_qkv_f32 (splits GDN's fused conv1d output into Q/K/V). Each replaces a host loop of many small async memcpy calls with one kernel launch; each lane/thread reads and writes exactly one element, so no host loop is needed regardless of seqLen/numHeads. Falls back to the original host loop when the PTX is stale (HasDeinterleaveF32, mirroring the existing HasElementwiseF32 defensive pattern for this file's other optional kernels). Validated: 4 new unit tests (DeinterleaveQGateF32_MatchesCpuReference, DeinterleaveGdnQkvF32_MatchesCpuReference) assert exact bit-for-bit equality against a CPU reference (pure gather, no floating-point math, so no tolerance needed) — decode-shape (seqLen=1) and multi-token cases, including real Bonsai head-count/dim values. Real-weight smoke test and full CUDA suite (290/290) pass clean. Profiled attn-2-deinterleave collapsed 94.55ms -> 2.12ms (44x) in the 4-decode-step profiling window. The GDN Q/K/V de-interleave loop (smaller, since it was already only 3 memcpy calls vs Q+Gate's 2*numHeads) also improved (30.61ms -> 17.30ms). Real-weight benchmark: prefill 9.88 -> 12.54 tok/s (+27% — the old loop scaled with seqLen, so prefill's 5-token calls benefited more than decode's 1-token calls), decode 7.09 -> 7.88 tok/s (+11%). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

… decode speedup) v2 fixed shared-x staging and grid sizing but still distributed GROUPS across lanes (`for g = lane; g < groups_per_row; g += warpSize`) — at any instruction the 32 lanes of a warp were reading 32 DIFFERENT groups, 34 bytes apart. That's the same "uncoalesced, single-byte weight loads" problem the original research (that led to v2) identified in the v1 kernel, just moved from per-element to per-group granularity — the actual global-memory weight reads were never fixed. v3 restructures the loop nesting so the warp cooperates on ONE group at a time instead of each lane owning whole groups: the group loop is now sequential (every lane executes it together), and within each iteration lane L reads code byte L of that group (`group_base[2 + lane]`) — 32 lanes reading 32 consecutive bytes, a single coalesced transaction. Byte L's decode target in dotLLM's PQ2_0 bit-interleave is exactly elements {L, L+32, L+64, L+96} of the group, so `pq2_0_accum_byte` (unchanged) is called with `lane` where the old per-lane `gp` loop variable was — the decode math is identical, only which lane touches which byte changed. Applied to both the single-projection and fused 2-way kernels. Total weight-byte traffic per warp is unchanged; this is a pure access-pattern reorganization, not a change to what's read. Validated: full CudaPQ2_0GemvTest suite (F16In vs CPU float reference across real Bonsai shapes + tail-row-clamp edge cases, fused-vs-separate parity) passes at the same tolerance as before. Real-weight smoke test and full CUDA suite (289/289, one pre-existing unrelated flaky test confirmed passing in isolation) pass clean. Real-weight benchmark: decode 7.88 -> 10.52 tok/s (+33%), prefill unchanged (different code path — untouched dequant-then-cuBLAS). Cumulative since the original naive kernel this session started from: 1.71 -> 10.52 tok/s (6.15x). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Tried bumping PQ2_0_ROWS_PER_WARP from 2 to 4 after the v3 coalescing fix, hypothesizing the ILP-vs-occupancy tradeoff that picked 2 for I2_S might shift now that weight reads are coalesced. Measured worse on real Bonsai-27B weights (decode 10.52 -> 9.91 tok/s, -5.8%, 3-rep/16-token benchmark) — reverted to 2. Documented in the kernel's comment so this isn't re-tried without new evidence. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

…ster The prefill-path dequant-to-F16 kernel (dequant_pq2_0_f16, feeds the dequant-then-cuBLAS-HGEMM path) had the exact same uncoalesced access pattern the GEMV decode kernel had before its v3 fix: one THREAD per 128-element group, so adjacent threads (same warp) processed adjacent groups — group_base addresses PQ2_0_GROUP_BYTES (34) apart. At any given instruction the warp's 32 lanes were reading/writing 32 bytes/halfs scattered ~1088 bytes / ~4KB apart. This kernel was never touched when the GEMV decode path was optimized earlier this session (prefill wasn't the initial focus), so it kept the original v1-style pattern throughout. Rewrites it as one WARP per group instead: the grid-stride loop now advances by warp, and within a group lane L reads code byte L (`group_base[2 + lane]`, coalesced 32-byte read) and writes to `out_base[lane]`/`[lane+32]`/`[lane+64]`/`[lane+96]` (each a coalesced 32-lane write — the old kernel's writes were also lane-scattered by 128 elements). Byte L's decode target is elements {L, L+32, L+64, L+96} of the group, so `lane` substitutes directly for the old per-thread `gp` loop variable. The C# launch wrapper's grid sizing is adjusted to count in warps, not threads, to match. Validated: new DequantPQ2_0ToF16_MatchesCpuReference test — bit-exact (0.0 abs diff) against the CPU Dequantize.ToFloat32 reference across real Bonsai shapes (n=512/k=5120, n=3/k=17408) plus a non-round-numbers tail case (n=37). Real-weight smoke test and full CUDA suite (289/289) pass clean. Real-weight benchmark: prefill 12.6 -> 89.06 tok/s (+7.07x). Decode unchanged (10.52 -> 10.34, within noise — decode uses the GEMV path, not this kernel). Cumulative from this session's starting baseline: prefill 9.86 -> 89.06 tok/s (9.03x), decode 1.71 -> 10.34 tok/s (6.05x). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

While optimizing PQ2_0's decode/prefill kernels this session (issue #157), found the exact same uncoalesced access pattern in I2_S's (BitNet ternary) prefill dequant kernel — dequant_i2_s.cu was structurally identical to dequant_pq2_0.cu before its fix (one thread per 128-element block, adjacent threads processing adjacent blocks 32 bytes apart). Since this is the identical, already-proven transformation (not new design work) and BitNet is a shipped feature, applying it here too rather than leaving a known duplicate of a just-fixed defect in a sibling kernel. NOTE: this commit is adjacent to issue #157's actual scope (Bonsai/PQ2_0), included here as an opportunistic low-risk fix rather than spinning up a separate issue/branch/PR during an unsupervised session — should be split out to its own issue/PR per the project's one-issue-one-branch convention if a maintainer wants clean attribution. Same fix as dequant_pq2_0.cu: one WARP per 128-element block instead of one thread, lane L reads/writes packed byte/output L (coalesced both directions). Grid sizing adjusted to count in warps, not threads. Validated: new DequantI2_SToF16_MatchesCpuReference test against the CPU Dequantize.ToFloat32 reference (n=2560/k=2560, n=37/k=2560 tail case) — ~5e-6 max abs diff (F16 rounding only). Full I2_S GEMV test suite (8/8, including the pre-existing F32In/A8/fused-decode tests) and full CUDA suite pass clean (one pre-existing unrelated flaky test confirmed passing in isolation). NOT validated end-to-end against a real BitNet model — no BitNet GGUF is available on this local machine this session (the integration tests that would exercise this, e.g. BitNetAccuracyTests.CpuVsCuda_LastTokenLogits_Match, self-skip with "BitNet GGUF not available"). Kernel-level correctness against the CPU reference is strong evidence, but real-model end-to-end prefill validation is a flagged follow-up. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

…rm_f32 kernel) Every layer's forward body calls the pattern "cuMemcpyDtoDAsync(residual, hidden) immediately followed by rmsnorm_f32(hidden, weight, normOut)" twice (once before token-mixing, once before the FFN) — saving the pre-norm value for a later residual add, then normalizing the same input. Profiling found this pair costs ~9% of decode time across 64 layers (128 launches/step), consistent with per-launch dispatch overhead dominating at this size. Adds copy_rmsnorm_f32.cu: byte-for-byte the same two-pass reduction as the existing rmsnorm_f32.cu (float2 vectorized loads/stores, __shfl_xor_sync symmetric warp reduction, fmaf-folded 1/n) with one addition — pass 1's existing per-element load also writes the fused residual copy (a free extra store next to a load already happening). Halves the launch count for this pattern from 2 to 1. Wired in with a HasCopyRmsNormF32 capability check and host-loop fallback (mirrors the deinterleave kernels' defensive pattern for stale-PTX builds). Validated: new CudaCopyRmsNormF32Test asserts the fused kernel's two outputs (residual copy + normalized result) are bit-for-bit identical to running the separate copy + rmsnorm_f32 pair — real Bonsai shapes plus an odd-hiddenSize case (rows=1 only; multi-row + odd n hits a genuine float2-cast alignment fault shared with the original rmsnorm_f32.cu, pre-existing and out of scope, never triggered in production since real hidden sizes are always even). Real-weight smoke test and full CUDA suite (297/298, one pre-existing unrelated flaky test confirmed passing in isolation for the 6th time this session) pass clean. Real-weight benchmark: decode 10.34 -> 10.51 tok/s (+1.6%, modest but real and zero-risk). Prefill unaffected (89.06 -> 88.51, within noise). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Fuses the per-GDN-layer decay+sigmoid pair (gdn_decay_f32 + sigmoid_f32, called once each per GDN layer x 48 GDN layers) into a single kernel (gdn_decay_sigmoid_f32), halving launch count for that pair. Lives in gated_delta_net_scan.cu, which is already in the -fmad=false bit-perfect kernel list since it sits in the GDN recurrence's precision-critical path. Validated: new unit test (GdnDecaySigmoidF32_MatchesSeparateDecayPlusSigmoid, seqLen=1/5) confirms bit-exact match against separate decay+sigmoid calls. Real-weight smoke test finite logits, full CUDA suite 301/301 clean. Bench (Bonsai-27B Q2_0, p64/n16/r3): decode 10.51 -> 10.59 tok/s (+0.8%), prefill flat at 88.3-88.5 (unaffected code path, as expected). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Extends the overnight CUDA sweep News entry with the copy_rmsnorm_f32 and gdn_decay_sigmoid_f32 fusion rounds and the re-profiled decision to stop there (GEMM-bound categories now dominate decode; the next non-GEMM candidate needs a riskier -fmad=false restructure for <1% gain). Final cumulative: prefill 9.86 -> 88.34 tok/s (9.0x), decode 1.71 -> 10.59 tok/s (6.2x). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

…#157) ncu --set full flagged the tail write of pq2_0_gemv_f16in/pq2_0_gemv2_f16in as the biggest single inefficiency in the decode-GEMV profile: each warp wrote its reduced row result via a single lane-0 scalar store (2 bytes into a 32-byte sector, ~6% MemoryCacheAccessPattern efficiency, ~51.56% Estimated Speedup) instead of coalescing across the block's 16 rows. Stage each warp's result into a small shared buffer (rowOut[16], 32 bytes) and, after a block-wide __syncthreads(), have the first 16 threads perform one coalesced write per block. pq2_0_gemv2_f16in routes each lane to y0/y1 based on global row, splitting into two coalesced runs only when a block straddles the n0/n1 boundary. The prior per-warp early "return" for out-of-range rows is no longer valid (all threads must reach the new __syncthreads()) — replaced with a warpActive guard that still falls through to the sync + write.

…ant reads (#157) pq2_0_gemv_f16in/pq2_0_gemv2_f16in's v3 warp-cooperative group read (group_base[2+lane]) is coalesced across lanes but each group's 32-byte code span starts at a non-32-byte-aligned offset (PQ2_0_GROUP_BYTES=34 never divides evenly into 32), so every read still spans 2 physical sectors. An in-kernel batched/staged fix was tried on a parallel investigation branch; its extra __syncwarp() barriers between staged batches blocked instruction pipelining and were measured to regress decode throughput, so it was not merged here. This instead performs a ONE-TIME repack at model-load time (CudaQwen3HybridDenseTransformerModel.UploadRawTensor): each PQ2_0 tensor is physically reordered from the interleaved on-disk layout (scale+codes interleaved per group) into a split layout (all group scales first, then all code bytes) via a new warp-per-group repack kernel (native/kernels/pq2_0_repack.cu, mirrors dequant_pq2_0.cu's existing warp-per-group structure). In the split layout, group g's code bytes always start at codesBase + g*32 — trivially a multiple of 32 — so the decode/dequant kernels gain unconditional alignment with ZERO added synchronization in the hot loop. The codes-region start is rounded up to the next 32-byte boundary (rather than the unpadded totalGroups*2) so this holds regardless of a tensor's row count parity (confirmed necessary by hand for the existing fused-decode test's odd n0=37/n1=53 fixtures, not just the two dominant real Bonsai shapes). pq2_0_gemv_f16in, pq2_0_gemv2_f16in, and dequant_pq2_0_f16 now read the split layout; pq2_0_gemv_f32in (CPU-vs-GPU correctness reference) is untouched. UploadRawTensor's signature and all ~13 call sites are unchanged — the repack is entirely internal to that one choke point, with its own lazily-loaded static PTX module (independent of the per-model CudaKernels instance, which isn't constructed yet at that point in the static loading pipeline). CudaPQ2_0GemvTest's F16In/Gemv2/DequantPQ2_0ToF16 cases now repack via the new CudaKernels.LaunchPQ2_0RepackSplitF16 wrapper before invoking the kernels under test, exercising the repack kernel itself as part of correctness validation rather than hand-repacking in C#. F32In/Pack() stay on the interleaved format, unchanged. Not benchmarked on hardware this session (GPU shared with other work). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

…atency (#157) ncu source-correlated stall counters showed the xs[i]=x[i] shared-memory staging loop at the top of pq2_0_gemv_f16in/pq2_0_gemv2_f16in as the single largest per-instruction latency stall in the kernel (dwarfing every weight-read PC), despite x[] already being perfectly coalesced: with blockDim.x=256 and k up to 17408, each thread issued up to 68 sequential 2-byte loads before __syncthreads() released any compute. Stage via uint4 (16 bytes/iteration) instead of half (2 bytes/iteration), cutting the sequential load count ~8x (68 -> 9 for k=17408). This targets round-trip latency (load count), a different lever from the two prior coalescing fixes (output-write staging, split-layout weight repack) which targeted per-transaction byte efficiency. k is always a multiple of 128 (existing on-disk-layout invariant), hence always a multiple of 8 -- no scalar tail path needed. xs[] gets __align__(16) (shared arrays aren't 16-byte aligned by default); the activation buffer x is always a cuMemAlloc_v2-backed base pointer (never a sub-offset), and CUDA's device allocator guarantees >=256-byte alignment, so the uint4 reinterpret is always valid. Pure perf change -- xs[] content is bit-identical regardless of transfer width.

…tional change (#157) Round 4 ncu follow-up flagged residual MemoryCacheAccessPattern inefficiency ("20.0/32 bytes utilized") after the split-layout/output-staging/vectorized- staging fixes already landed. Traced by SASS inspection (cuobjdump --dump-sass on a locally-built sm_86 cubin, no GPU execution) to the per-group scale read in pq2_0_gemv_f16in/pq2_0_gemv2_f16in: every lane computes the identical address (gFlat depends only on warp-uniform rowBase/g, not lane), so this is a genuine single-address broadcast that the LSU already collapses into one sector fetch per instance regardless of the compiler not hoisting it into a uniform-register load — consistent with a known false-positive pattern for ncu's byte-utilization rule on broadcast reads, not evidence of redundant traffic. Considered and rejected (with arithmetic) rather than implemented: - Bulk shared-memory staging of group scales: needs 4352 bytes (k=17408) or 1280 bytes (k=5120) per block vs only 304 bytes of headroom under sm_86's 48KB static cap. Infeasible without cutting PQ2_0_ROWS_PER_BLOCK, which the already-landed ROWS_PER_WARP=4 experiment showed regresses throughput. - __shfl_sync broadcast in place of the redundant per-lane load: doesn't remove a real transaction (the hardware already broadcasts), only trades cheap redundant LDG issues for a shuffle instruction — an unmeasurable bet in a session where GPU execution is off-limits, and this investigation's own history (batch-staging: modeled ~94% sector efficiency, measured -57% throughput) argues against committing modeled memory-pattern wins without measurement. No functional/numerical change — comments only. PTX build verified via nvcc -ptx -arch=compute_75.

…ing (#157) ncu still showed the F16 PQ2_0 decode GEMV latency-bound with achieved occupancy stuck at 33.3% theoretical, bound by Block Limit Shared Mem=2 — xs[PQ2_0_MAX_K=17408] costs 34848 B/block on EVERY launch, including the much smaller k=5120 attention/GDN call sites that only need 10240 B. Add pq2_0_gemv_f16in_small/pq2_0_gemv2_f16in_small (xs[5120]), routed transparently from LaunchPQ2_0GemvF16In/LaunchPQ2_0Gemv2F16In based on k (k<=5120 -> small variant, unchanged call sites). Occupancy arithmetic (sm_86: 100KB/SM shared mem, 65536 regs/SM, 1536 threads/SM): 10272 B/block -> floor(102400/10272)=9 shared-mem-limited blocks, but registers (48 regs/thread * 256 = 12288 regs/block -> floor(65536/12288)=5) now bind first -> 5*8/48 = 83.3% theoretical occupancy (up from 33.3%). Verified against ncu's reported Block Limit Registers=5/Warps=6, and the large-kernel math (34848 B/block -> floor=2) matches ncu's reported Block Limit Shared Mem=2 exactly before trusting the new-case number. Test coverage: existing k=5120/17408 + tail (n=37,n=3) cases now exercise both kernel variants transparently via the wrapper; added k=5248 (Pq2_0MaxKSmall+128) boundary cases to both the single-GEMV and fused-GEMV2 theories to guard the dispatch threshold itself (a misrouted k=5248 into the 5120-sized xs[] would overflow shared memory). Not yet measured on real hardware — theoretical occupancy is not measured throughput (see this file's own reverted-experiment history). Next: confirm on real Bonsai-27B decode.

jamesburton and others added 2 commits July 20, 2026 23:02

docs: README News entry + perf-matrix baseline for Bonsai-27B (#157)

81307d1

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

Copilot AI review requested due to automatic review settings July 20, 2026 22:04

Copilot started reviewing on behalf of jamesburton July 20, 2026 22:04 View session

Copilot AI reviewed Jul 20, 2026

View reviewed changes

jamesburton and others added 17 commits July 21, 2026 00:02

docs: consolidate overnight CUDA optimization sweep summary in README

1c79488

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>

wip: temporary set-aside of batch-8 negative-result doc note

1ebb9be

jamesburton merged commit a57dcac into dev Jul 22, 2026
4 checks passed

jamesburton deleted the issue/157-qwen35-dense-hybrid branch July 22, 2026 18:46

Navigation Menu

Repository navigation

Dense Qwen3.5 hybrid architecture — PrismML Bonsai-27B runs end-to-end (#157) - #158

Dense Qwen3.5 hybrid architecture — PrismML Bonsai-27B runs end-to-end (#157)#158
jamesburton merged 19 commits into
devfrom
issue/157-qwen35-dense-hybrid

jamesburton commented Jul 20, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

jamesburton commented Jul 20, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Results

Test plan

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

jamesburton commented Jul 20, 2026 •

edited

Loading