feat(everything-bagel): clean migration of PR #9 + #10 onto updated dev#14
Merged
Conversation
`Device.markWeightsResident(_:)` attaches the given MTLBuffers to a persistent residency set on the command queue. Without it, the Metal driver re-validates per-allocation residency on every command-buffer encode; at large MoE models with thousands of dispatches per prefill, that overhead dominates wall time. The set is shared across all weight buffers and lazily created on first call. Repeated calls add to it. macOS 15 / iOS 18 minimum; older OSes no-op silently. `FFAI_NO_RESIDENCY_SET=1` disables for A/B benching. Per-model integration (call `markWeightsResident` after `Model.load`) lands in a follow-up commit so this stays a pure infrastructure add. KV / GDN / conv state cache pinning also lands separately so the residency-set adoption can be diffed per surface.
Walks every entry of the loaded `SafeTensorsBundle` and registers its backing MTLBuffer with `Device.markWeightsResident(_:)` immediately after the bundle resolves, before `ModelRegistry.dispatchAndLoad` kicks the per-family loader. The driver's per-command-buffer residency re-validation is the cost we avoid. Opt-out via `FFAI_NO_RESIDENCY_SET=1` (handled inside the helper). Older OSes silently no-op via the `@available(macOS 15, iOS 18, *)` gate on the helper.
Each cache class owns one or two MTLBuffers that persist for the lifetime of a generation and get touched by every decode-step dispatch. Without explicit residency pinning, Apple's Metal driver re-validates per-allocation residency on every command-buffer encode — at decode rates of thousands of dispatches per token, the per-dispatch overhead dominates wall time. Reuses the device-wide MTLResidencySet introduced for weight buffers (see `Device.markWeightsResident`). The set is single-flighted under a lock and shared across all cache classes; the underlying Metal API gracefully no-ops on macOS < 15 and is gated behind `FFAI_NO_RESIDENCY_SET=1` for A/B comparison. - `KVCache.init` pins kBuffer + vBuffer - `GDNStateCache.init` pins current + next recurrent-state buffers - `ConvStateCache.init` pins the rolling-window state Migrated from PR #9 (Everything Bagel 1) — measured win was part of the 16.29 → 55.94 tps decode T=1 jump on Qwen3.6-A3B / M5 Max.
The Qwen3-MoE decode path needs a scalar-broadcast FMA primitive to accumulate the top-K expert outputs into a single hidden-state vector. Without it the path was forced into a `Tensor.filled([hidden], weight)` host alloc per expert plus a `mul + add` pair per expert — a host roundtrip and 2 dispatches the GPU could absorb into one. Adds four wrappers (all in `Ops/OpsFused.swift`) over kernels that metaltile already emits: - `scalarFMA(scalar, value, base, into:)` — single-dispatch `out[i] = base[i] + scalar[0] * value[i]`. - `scalarFMAMany(scalars, values, acc:)` — N sequential dispatches reusing one compute encoder; Metal guarantees in-order encoder execution so the accumulation is safe. - `scalarFMAChain8(scalars, values, out:)` — exact 8-way fused dispatch (one kernel reads 16 inputs at once), collapses the topK=8 expert accumulator into a single dispatch. - `sigmoidScalarFMAResidual(gate, value, base, residual, into:)` — sigmoid-gated FMA with the post-MoE residual add folded in; Eric's `sigmoidScalarFMA` covers the non-residual variant. Tests in `Tests/FFAITests/Ops/OpsSpecialPathTests.swift` cover each wrapper with a closed-form expected value (4 new tests). Migrated from PR #9 / #10 (Everything Bagel 1 + 2) — measured as load-bearing for the decode T=1 16.29 → 96.3 tps win on Qwen3.6-A3B / M5 Max.
Three batched wrappers around the existing `dequant_gemv_int4_*` metaltile kernels that drive N back-to-back dispatches inside a single compute encoder, sharing the `input` binding and the `in_dim` / `group_size` constants. Saves N-1 encoder begin/end pairs per call. - `dequantGemvInt4Two` — per-expert MoE gate + up projections. - `dequantGemvInt4Three` — Qwen3-style attention q/k/v projections in the unfused decode path. - `dequantGemvInt4Four` — Qwen3.5 GDN mixer's qkv/z/b/a input-projection quartet, all reading the same xNorm. CPU-correctness tests in `Tests/FFAITests/Ops/QuantizedOpsTests.swift` reuse a `makeInt4Projection` helper to build packed weights + scales + biases + expected GEMV output, then verify each batched wrapper matches CPU truth for every output tensor. Migrated from PR #9 / #10.
Both Q and K go through the same partial RoPE rotation with identical `(headDim, rotaryDim, position, thetaBase, scaling)` parameters every decode step on Qwen3-class models. `ropePartialTwo` sets the RoPE constants once and dispatches both rotations through a single shared compute encoder, saving one encoder begin/end pair per attention layer per token. The kernel is unchanged (`ffai_rope_llama_*`); the wrapper is a buffer-binding rotation around the existing in-place rotation. A test verifies bit-identical output against two sequential `ropePartial` calls on the same inputs. Migrated from PR #10.
Pairs two `ffai_gather_*` dispatches against the same `table` binding inside one compute encoder, saving an encoder begin/end pair when the caller needs to fetch two parallel id streams from the same embedding (e.g. the Qwen3 attention head-half split path). Migrated from PR #10.
Two `mt_rms_norm_*` dispatches inside one compute encoder, used by the Qwen3 attention mixer's pre-RoPE Q-norm + K-norm pair: both projections run the fast path (TPG = rowSize / 4) and have no data dependency between them. The wrapper shares a single eps buffer allocation when the two eps values are equal — the common case since both norms come from the same checkpoint config. The wrapper restricts callers to the fast `mt_rms_norm_*` path (rowSize ≤ 4096 and a multiple of 128). Wider rows must fall back to two separate `rmsNormRows` calls so the wide kernel is picked per call. Test verifies bit-equivalence against two sequential `rmsNormRows` calls on the same inputs. Migrated from PR #10.
Wraps two `kv_cache_update_*` dispatches in a single compute encoder. Cache layouts, dtypes, and writing thread are identical to `kvCacheUpdate`; this only saves the encoder begin/end pair between the K-cache write and the V-cache write. With 10 full attention layers on Qwen3.6-A3B, this drops 10 encoder pairs per decode token. Migrated from PR #10.
…n one encoder Used inside the Qwen3.5 GDN T-loop where, every token, one tensor goes through `siluCastToF32` (the gate input) while the two raw gate scalars `aRaw` / `bRaw` go through plain `castToF32`. The wrapper sets the silu PSO once, dispatches the silu, then switches to the plain-cast PSO and dispatches the two raw casts — all inside one compute encoder. Saves the encoder begin/end pairs the sequential 3-call path would create per GDN layer per token. Migrated from PR #10.
The Qwen3-MoE decode path historically had to commit + wait on a
single GPU layer just to read the router logits back to the host,
run a CPU top-K + softmax, and then dispatch the per-expert gate +
up + down GEMVs. Two new wrappers move both halves to the GPU:
- `moeRouterTopK` / `moeRouterTopKMany` — wrap the
`mt_moe_router_topk_*` kernel. Single-row form takes
`logits: [nExperts]` and writes `[k]` u32 indices plus `[k]`
weights; the T-batched form iterates `program_id<0>()` across
prefill rows in one dispatch.
- `dequantGemvInt4ExpertIndexed` /
`dequantGemvInt4ExpertIndexedMany` — wrap
`dequant_gemv_int4_expert_indexed_*`, which reads
`expertIndex[0]` at kernel entry to pick which slab of a
stacked `[nExperts, outDim, inDim/8]` u32 weight tensor to
dequantize. The `Many` variant runs N sequential dispatches
inside one compute encoder sharing the constexpr binding so
Qwen3.6-A3B's topK=8 gate + up phase collapses to one encoder
per MoE layer.
The router output indices buffer can be aliased into N single-u32
views (`buffer + offset = slot · 4`) and threaded straight into
the expert-indexed GEMV — no host readback per layer.
Tests:
- `dequantGemvInt4ExpertIndexedCorrectness` — stacks 4 experts'
weights, picks expert 2, and checks the output matches a
CPU-truth dequant + GEMV of that expert's slab.
- `moeRouterTopKDeterministic` — top-K of (1, 5, 3, 4) at k=2
picks indices {1, 3}; weights sum to 1.0 and the larger logit
dominates.
Migrated from PR #10 ITERs 54–60.
Eric's `sdpaMulti` on dev only switches on dtype — every call lands
on `ffai_sdpa_multi_*` (the head_dim=128-only kernel). Qwen3.6-A3B's
full-attention layers run at head_dim=256, and metaltile emits a
dedicated `ffai_sdpa_multi_d256_*` family with a 2-phase output
reduction that keeps the per-row threadgroup memory under Apple's
32 KB cap. Without the d256 routing the d128 kernel reads past
each row of K/V and produces silent garbage.
Switches `sdpaMulti` to a `(dtype, headDim)` tuple match so f32 /
f16 / bf16 each pick the correct kernel; relaxes
`OpsValidation.validateSdpaMulti` to accept `headDim ∈ {128, 256}`
and updates the matching validator test.
Adds a runtime test that mirrors the existing d=128 uniform-K case
at head_dim=256, exercising the new routing end-to-end.
Migrated from PR #10 ITER 93 (Qwen3.6-A3B prefill wiring needs it).
Eric's dev has the base versions (`castToF32`, `swiglu`, `gatedMixerNorm`) but not the batched variants our PR9/10 wirings rely on. Each variant reuses the same metaltile kernel and collapses N independent dispatches onto a single compute encoder: - `castToF32Two` / `castToF32Three` — Qwen3 GDN fused-prep path promotes 2-3 same-dtype tensors (convAct / aRaw / bRaw) to fp32 every token; one shared encoder replaces 2-3 separate ones. - `swigluMany` — MoE per-expert SwiGLU at decode runs topK independent dispatches per layer; one shared encoder per call. - `gatedMixerNormMany` — Qwen3.5 GDN mixer's per-token norm is evaluated `T·Hv` times during prefill; widening the grid X axis from `Hv` to `T·Hv` collapses the T-loop into a single dispatch. Tests in `Tests/FFAITests/Ops/OpsSpecialPathTests.swift` verify each batched form matches the corresponding sequential reference. Migrated from PR #10 (closes the Tier 2 paired-on-encoder gap identified during the Eric-overlap audit).
…ateKVMany, conv1dCausalStepSiluCastMany, gatedDeltaPrepChunk Eric's dev only has the single-token forms of these ops; the PR9/10 forwardMany / decodeMany prefill paths replace the per-token T-loop with a single batched dispatch. - `ropePartialMany` — single dispatch covering T rows of `qk` rotated by per-row `positions[T]`; saves T-1 encoder pairs and matches T sequential `ropePartial` calls bit-identically. - `ropePartialManyTwo` — Q + K share `positions` and constants on one encoder; row strides differ. - `kvCacheUpdateKVMany` — writes T tokens' K + V rows in one encoder, indexed by `positions[T]`. - `conv1dCausalStepSiluCastMany` — `T` depthwise conv + SiLU + fp32 cast collapsed into one dispatch with the conv state held in per-channel registers across the sweep. - `gatedDeltaPrepChunk` — chunked-prefill counterpart to `gatedDeltaPrepStep`; runs the fused GDN prep + recurrence over a token range in one dispatch, keeping per-head recurrent state in registers across the chunk. `tLen` arrives as a `[1]` u32 GPU buffer so chunked dispatch chains avoid CPU readback. All kernels (`ffai_rope_llama_many_*`, `kv_cache_update_many_*`, `ffai_conv1d_causal_step_silu_cast_many_*`, `mt_gated_delta_prep_chunk_*`) are already emitted on the dev metaltile. Tests: - `ropePartialManyMatchesSequential` — T=3 batched matches three sequential `ropePartial` outputs bit-identically. - `kvCacheUpdateKVManyF32` — verifies multi-position K + V writes land in the right cache rows and leave untouched slots zero. Migrated from PR #10 (ITERs 86–99).
…coder Mirror of `dequantGemvInt4Two/Three/Four` for the per-expert MoE down phase, where each chosen expert has its own activation tensor but all N projections share PSO, `inDim`, and `groupSize`. The wrapper sets the constexpr binding once and rotates `(weight, scales, biases, input, output)` per dispatch, saving N-1 encoder begin/end pairs versus N independent `dequantGemvInt4` calls. Test verifies the batched form matches three sequential `dequantGemvInt4` CPU-truth references across distinct (scale, bias) projections + distinct inputs. Migrated from PR #10.
…code Same dispatch contract as `sdpaMulti(causal: false)`, but the scalar `causal` flag is replaced by an additive `[nQuery, nQuery]` mask tensor consulted only for in-block KV positions. Positions `< baseKV` (the cached prefix) are always attended. Mask is `0.0` to allow and `-inf` to block, added directly to pre-softmax scores. Used by speculative-decode tree-verify: one verifier call attends every leaf-to-root path of the draft tree at once, gated by a tree-causal mask the drafter produces from its node structure. Wraps `ffai_sdpa_multi_tree_mask_*` (headDim=128 only — no d256 variant for tree-verify on dev). Tests: - All-zero mask reproduces `sdpaMulti(causal: false)` output bit-equivalent (sanity check on the mask add path). - Diagonal-only -inf mask isolates each query to its own KV slot: with uniform Q/K and per-row distinct V, output collapses to a copy of the matching V row. Migrated from PR #10 ITER 70.
Pass 1 partitions the KV cache into `blocks` even tiles and writes per-block partial `(O, max, lse)` triples; pass 2 merges them with the log-sum-exp trick. Total work matches single-pass `sdpaDecode` but the two passes parallelize across more GPU threadgroups, paying off at long KV (≥ 1K typical) where the single-pass kernel starves the GPU with one TG per Q-head. Block-count contract (the load-bearing detail PR10's port silently violated): the pass2 kernel hardcodes a 32-lane block reduction (`bn = 32`, `block_chunks = blocks / bn`), so `blocks` MUST be a multiple of 32 AND at least 32. Anything smaller makes `block_chunks` round to zero — pass2's outer loop iterates zero times and the output stays whatever pass1 left in scratch, typically all zeros. The previous port's "power of 2" check let `blocks=2` through and was the cause of the all-zero smoke failures during the earlier port attempt; the new precondition catches it at the wrapper. Scratch buffers (`partialO`, `partialM`, `partialL`) are caller-owned so a decode loop can amortise them across steps. Test: `blocks=32, nKV=64, nQHeads=2` — verifies the two-pass output matches single-pass `sdpaDecode` to within 1e-2 absolute. Migrated from PR #10 ITER 53.
…jection) Wraps `ffai_batched_qkv_qgemv_fast` (M=1 decode) and `ffai_batched_qkv_qmm_fast` (M>1 batched prefill). Both fuse the three Q / K / V int4 dequant-projection dispatches into ONE kernel launch — `program_id<2>()` selects the matrix across a depth-3 grid, so all three projections share TG memory and read the input activation once instead of three times from DRAM. Decode path (M=1): replaces the 3-dispatch `dequantGemvInt4Three` shared-encoder form. Saves 2 encoder begin/end pairs per attention layer. Prefill path (M>1): outputs land in three separate contiguous tensors (`qBuf [M, out_q]`, `kBuf [M, out_k]`, `vBuf [M, out_v]`) so downstream code reads each projection as `[M, dim]` without strided views. Pre-zeroes the outputs since the kernel's tile mechanic relies on smaller-matrix tile-past-out_* slots being already-zero. Kernel constraints (preconditioned): `in_dim % 512 == 0`, each `out_q / out_k / out_v % 8 == 0`, `group_size = 64`, TPG = 64. Dispatches f32 / f16 / bf16 variants.
…GDN input proj) Wraps `ffai_batched_4_qgemv_fast` (M=1 decode) and `ffai_batched_4_qmm_fast` (M>1 batched prefill). Extends the QKV (3-output) pattern to 4 projections sharing the same input — used by the Qwen3.5 GDN mixer where qkv / z / b / a all project from the same `xNorm` and can fuse into a single kernel launch. Decode: replaces the 4-dispatch `dequantGemvInt4Four` shared- encoder form. Pays the input DRAM roundtrip once instead of four times; saves 3 encoder begin/end pairs per GDN layer. Prefill (M>1): reads `x = [M, in_dim]` once into TG memory and produces all 4 outputs (each `[M, out_*]`) in one dispatch. Eliminates the 3 redundant `xNormFlat` DRAM reads of the unfused `callMany` chain. Kernel constraints (preconditioned): `in_dim % 512 == 0`, each `out_a/b/c/d % 8 == 0`, `group_size = 64`, TPG = 64. Dispatches f32 / f16 / bf16 variants.
…+gemv) Wraps `ffai_rms_norm_qgemv_fast` (RMSNorm-fused) and `ffai_gated_rms_norm_qgemv_int4_fast` (Mamba2-style gated-RMSNorm- fused). Both keep the normalised activation in registers and feed it straight into the int4 dequant-GEMV, replacing a 2-dispatch chain (`Ops.rmsNorm` / `Ops.gatedRmsNorm` + `Ops.dequantGemvInt4`) with a single kernel. Saves a full `[in_dim]` DRAM roundtrip on the intermediate `normed` per call. `rmsNormQgemvInt4Fast` is the workhorse for the finalNorm + lmHead boundary (one dispatch per token) and any other pre-norm → int4-qmm pattern in transformer blocks. `gatedRmsNormQgemvInt4Fast` is purpose-built for the Qwen3.5 GDN mixer post-recurrence boundary: collapses `silu(z) · w · rmsnorm(y)` + outProj qgemv into one launch. Kernel constraints (preconditioned): `in_dim % 512 == 0`, `out_dim % 8 == 0`, `group_size = 64`, TPG = 64. The gated variant additionally requires `Hv·Dv ≤ 8192` (kernel stages the full input in TG memory). Dispatches f32 / f16 / bf16 variants.
Wraps `ffai_moe_down_swiglu_accum_int4_chain8`. Collapses the
GPU-router MoE path's three back-to-back per-layer dispatches:
1. `swigluMany` — inner[k] = silu(gate[k]) * up[k]
2. `dequantGemvInt4ExpertIndexedMany` (down)
— out[k] = W_down[expert_idx[k]] · inner[k]
3. `scalarFMAChain8` — acc[i] = Σ_k slot_weight[k] · out[k][i]
into ONE kernel launch. Each threadgroup owns one output row of
`[out_dim]`, iterates the 8 slots sequentially: stages `inner` in
3 KiB threadgroup memory, runs the dequant-gemv inner-product
against `W_down[expert_idx[k]]` for that slot, accumulates into a
per-thread `acc` with the slot scalar baked in, then reduce-sums
across threads at the end. Eliminates the `inner[k]` and `out[k]`
DRAM roundtrips between phases — 16 buffer-write-then-read pairs
per layer at topK=8.
Kernel constraints (preconditioned): `in_dim ≤ 768` (TG-memory
cap, kernel statically allocates 768 floats), `group_size = 64`,
TPG = 128. Lives in OpsFused.swift next to `scalarFMAChain8`
(the third stage it collapses). Dispatches f32 / f16 / bf16
variants.
Note: kernel ships in metaltile PRs #196–#201 (regenerated
locally into Sources/MetalTileSwift/Generated/MetalTileKernels.swift
for this branch).
Batched-T KV append wrapper around Ops.kvCacheUpdateKVMany — replaces the T-loop of Ops.kvCacheUpdateKV calls with one shared encoder + 2 dispatches. Plus the host-side slot reservation helper that fills a u32 positions tensor under lengthLock. Pre-req for ITER 98 wiring into Qwen35AttentionMixer.forwardMany.
Commit message hygiene checkAll commit messages and PR text are clean. ✅ |
… scratches PR10 ITERs ported into the single-token attention decode path: - ITER 22/28: cache sliceHeadHalves35 row-index tensors at init, pin to residency set (was rebuilt + memcpy'd per token). - ITER 24: batch q+k+v projections through Ops.dequantGemvInt4Three when all 3 projections are QuantizedLinear int4 with matching groupSize. - ITER 78: single-dispatch Ops.batchedQkvQgemvInt4Fast when constraints match (in_dim%512==0, out_*%8==0, groupSize==64). Backing scratch is one contiguous [qDim+kDim+vDim] buffer; gate via FFAI_NO_FUSED_QKV. - ITER 65: Ops.gatherTwo replaces the two separate sliceHeadHalves35 gathers with one shared-encoder dispatch. - ITER 12/35: Ops.rmsNormRowsTwo collapses q_norm + k_norm into one shared encoder, cached scratch outputs. - ITER 21: Ops.ropePartialTwo pairs Q+K RoPE on one shared encoder. - ITER 53: Ops.sdpaDecode2Pass (Flash-Decoding) kicks in at long KV (default threshold 1024, FFAI_SDPA_2PASS_THRESHOLD overrides). - ITER 11/40/43: Ops.sigmoidMul with cached output scratch replaces the mul+sigmoid chain at the attention output gate.
…ding Speculative-decode candidate-proposer interface used by the in-progress `SpecDecode` driver. Three implementations land here: - `NGramDrafter` — prompt-lookup n-gram. Scans `history` for the last `nMatch` tokens, falls back to shorter n-grams down to `minNMatch`, returns up to `gamma` follow-on tokens from the earlier occurrence. Zero ML cost; the workhorse for code and structured-chat workloads where prior context recurs. - `NeverDrafter` — stub that always returns `[]`. Used to bench pure spec-decode driver overhead vs raw single-token decode. - `NGramTreeDrafter` — real branching tree drafter. Conforms to both `Drafter` (linear γ fallback via `propose`) and `TreeDrafter` (`proposeTree` returning top-K continuations per depth) for tree-verify spec decode against `Ops.sdpaMultiTreeMask`. `DraftTreeNode` is the pure-Swift tree representation, with `flatten()`, `verify(oracleAtHistoryEnd:oracle:)`, and `treeCausalMask()` helpers — all pure functions, no model / cache dependencies, ready to wire into the SpecDecode driver once `forwardManyAllLogits` lands. `LinearTreeAdapter` lets any linear `Drafter` masquerade as a `TreeDrafter` (degenerate single-branch tree) so the tree-verify driver can A/B against the linear baseline using the same `NGramDrafter` instance. Tests in `Tests/FFAITests/Generation/DrafterTests.swift` cover all three drafters end-to-end: 20 host-side tests, no GPU dependencies. All pass. Migrated from PR #10 (ITERs 61, 68, 69, 71).
…many-RoPE/KV PR10 ITERs ported into the T-batched prefill path: - ITER 88: Ops.batchedQkvQmmFast fuses Q/K/V projections into one dispatch when fusedQKVEligible (same constraints as ITER 78 single-token path). - ITER 74: sliceHeadHalvesManyBoth35 helper — one shared-encoder Ops.gatherTwo dispatch for the queries+gate split (was two back-to-back sliceHeadHalvesMany35 calls). - ITER 83: Ops.rmsNormRowsTwo collapses Q/K rmsNormRows into one shared encoder over (T·nHeads) + (T·nKVHeads) rows. - ITER 84+97: Ops.ropePartialManyTwo replaces the T-loop of per-row ropePartial calls with ONE shared-encoder dispatch over all T tokens (Q+K paired). Saves ~5100 dispatches per layer at T=512. - ITER 98: kv.appendRangeOnGPUMany + reserveSlotsManyOnHost — one shared encoder + 2 dispatches for the K/V cache append over all T tokens (was a T-loop of per-token kvCacheUpdateKV calls). - ITER 90/93: Ops.sdpaMulti collapses T-token causal attention into ONE dispatch. headDim=128 routes to ffai_sdpa_multi; headDim=256 (Qwen3.6-A3B) routes to ffai_sdpa_multi_d256. Other head_dims fall back to the T-loop sdpaDecode form. - ITER 14: Ops.sigmoidMul on the prefill output-gate path.
ITER 16/17 (PR10): replace silu+mul chain with Ops.swiglu in Qwen35DenseMLP.forward/forwardMany and Qwen35MoEFFN.forward (single- token) + forwardMany (batched) shared-expert paths. One fused dispatch per call vs two (silu + mul).
ITER 76/77/79 (PR10): collapse the finalNorm + lmHead 2-dispatch chain into one Ops.rmsNormQgemvInt4Fast call when the lmHead is 4-bit QuantizedLinear with constraints (in_dim%512==0, out_dim%8==0, groupSize=64). Fires on Qwen3.6-A3B's quantized lmHead. Wired in three forward paths via shared `qwen35FinalNormLmHead` helper: - Qwen35Model.forward (single-token decode) — ITER 76 - Qwen35Model._forwardManyBatched (batched-T prefill last-row) — 77 - Qwen35Model.forward(inputEmbedding:) (VLM embed-input) — ITER 79 Fallback to separate finalNorm + lmHead dispatches when constraints not met (dense lmHead, int8/int3 quant, etc.).
Greedy spec-decode loop for Qwen3.5/3.6:
1. Drafter proposes up to γ candidates via Drafter.propose.
2. Snapshot caches; run forwardManyAllLogits([prev, c_0..c_{γ-1}])
to advance caches by γ+1 tokens and get per-position logits.
3. Greedy verify: accept consecutive argmax matches up to the first
mismatch.
4. On partial accept (any mismatch): restoreAll caches from
snapshot, replay accepted prefix via single-step forward, take
next-prev from the LAST single-step (bit-identical to baseline
greedy loop; forwardManyAllLogits drifts slightly due to
different kernel paths). On full accept: commit all candidates +
bonus token from trailing logits[γ].
5. Empty proposal → fall back to single-token forward.
Stops on maxNewTokens or stop-tokens. SpecDecodeStats records
acceptanceRate, tps, fallbackSingleSteps, etc.
Greedy-only for v0 — temperature / nucleus / top-K follows the same
loop with probability-ratio rejection comparing sampled tokens against
the drafter's proposed tokens.
…eTests The concurrency test uses a TaskGroup of MTLComputePipelineState, which under Swift 6 strict concurrency tripped 'type does not conform to the Sendable protocol' on the CI runner (Xcode 16.4 SDK). Adding @preconcurrency to the Metal import keeps the same test behaviour but lets the compiler treat Metal's pre-Sendable protocols as warnings rather than errors, matching the convention the rest of the project uses for system-framework imports under strict concurrency.
…en3.6 Integration test from PR #9 that compares baseline greedy decode (forward loop) vs SpecDecode.generateGreedy + NGramDrafter at gamma in 1, 2, 4. Verifies generated sequences match baseline exactly (greedy + greedy drafter + correct snapshot/restore is deterministic) and reports decode-only tps so the cost of the verify forward is visible. Gated on the local Qwen3.6-A3B checkpoint at /Users/tom/models/Qwen3.6-35B-A3B-4bit; skips with a print when the path is missing.
…ingle-step loop Measures whether spec-decode's verify step should use forwardManyAllLogits(T) or a loop of single forward() calls. Times 16 iterations of each path with cache snapshot/restore between iterations (so each iter starts in identical state) and reports the median ms / iter ratio for T=3 (gamma=2 verify) and T=2 (gamma=1 verify) shapes. The answer determines the SpecDecode driver's optimal verify primitive once the cost crossover point is found. Gated on the local Qwen3.6-A3B checkpoint at /Users/tom/models/Qwen3.6-35B-A3B-4bit; skips with a print when the path is missing.
… type-check it CI build was failing on Swift 6 / Xcode 16.4 with 'the compiler is unable to type-check this expression in reasonable time' on the 8-byte little-endian length decode at line 1202. The pre-existing single expression chained 8 byte loads + 7 bit-shifts + 7 bit-ORs through Int promotion, which trips the type-inference budget under strict concurrency. Split into two 32-bit halves (nLo32 + nHi32) joined by one final OR; identical bit semantics, type-check finishes in <1 ms.
…harness Same shape as Qwen3.6's forwardManyBench but pointed at the locally- available Qwen3.5-35B-A3B-4bit checkpoint. Engine is m.qwen35 in both cases, so this exercises every Bagel-migration wiring landed on tom/bagel-clean (GPU MoE router, batched QKV, GDN chunked recurrence, rmsNormQgemvInt4Fast finalNorm, 2-pass Flash-Decoding, MTLResidencySet weight pinning). Each test gates on the local checkpoint existing and prints a skip message when missing. M5 Max decode T=1 measurement: 96.87 tps (median of 5 × 32-step runs), matching the PR #10 Qwen3.6-A3B tip target of 96.3 tps.
…32 guard QuantizedLinear.callMany routes 4-bit weights to Ops.dequantGemmDynamicM (mt_qmm_mma backed), whose kernel requires outDim to be a multiple of 32 (the BN tile of the MMA kernel). Without an outDim guard the precondition tripped on Qwen3.5-35B-A3B's shared_expert_gate, which quantises to 4-bit at outDim=1 (a per-token gate scalar): 'Ops.dequantGemmDynamicM: nOut (1) must be multiple of 32 (BN tile)'. Add the multiple-of-32 check next to the 4-bit branch. Anything that doesn't fit (4-bit narrow projections + every non-4-bit) falls through to the per-row dequantGemv loop, bit-identical to the per-token path and irrelevant to perf at outDim=1. Verified on M5 Max with Qwen3.5-35B-A3B-4bit: forwardManyBench T=8 / T=128 / T=512 all complete cleanly. Headline perf: decode T=1 96.87 tps, prefill T=512 458.6 tps (4.75x speedup over per-token loop) — matches the PR #10 tip targets of 96.3 / 455 tps.
…e4/LFM2/Jamba dense MLP Two ITERs from PR #10 the wirings agent missed in the first pass. ITER 20: bump MetalTileLibrary.defaultMaxCommandBufferCount from 16 to 64. Lets the Metal driver pipeline more cmd-buffers in flight before applying backpressure; absorbs per-cmd-buffer encode latency across more concurrent submissions. PR #10 measured +2-3% prefill T=512 on Qwen3.6-A3B / M5 Max. FFAI_MAX_COMMAND_BUFFERS override preserved. ITER 17: replace Ops.mul(Ops.silu(g, cmd), u, cmd) with the fused Ops.swiglu in Granite4Text.MLP.forward, LFM2Text.MLP.forward, and JambaText.MLP.forward. Same semantic; one kernel dispatch instead of two, intermediate silu(g) stays in registers instead of round- tripping to DRAM. Brings these three cross-family models in line with Qwen3.5 (already wired via the agent's earlier ITER 16 commit on Qwen35DenseMLP + MoEFFN shared expert).
… bench cases Extends the bench harness with two long-context shapes that PR #10's work needs to validate but Eric's existing harness only covered up to T=512: - forwardManyT2K (T=2048): exercises chunked GDN recurrence over a longer window + larger SDPA prefill grid. - decodeAfterLongPrefill: prefill 1024 tokens via forwardMany, then measure decode-only tps over 32 steps. Catches any KV-cache / residency / sdpaDecode2Pass routing regression at non-trivial KV. Measured on M5 Max, Qwen3.5-35B-A3B-4bit: forwardManyBench T=2048: per_token=21499ms batched=4755ms speedup=4.52x batched_tps=430.7 decode T=1 after T=1024 prefill: runs=[0.335, 0.335, 0.336, 0.336, 0.337]s median=0.336s -> 95.37 tps Decode tps barely drops from empty-cache (96.51) to KV=1024 (95.37) - the kernel-compute floor PR #10 documented at 96 tps holds through the long-context regime.
…-context bench
Drops the lukewarm 1024-token case (KV size barely above empty-cache,
decode tps barely budged from 96.51 to 95.37) and replaces it with
a real long-context probe at T=32768.
Single-run shape — 5-run median is impractical at 32K (each run is
multi-minute). Per-token baseline is intentionally not measured at
that size (would be 5+ minutes per run). Measures:
- Batched forwardMany at T=32768.
- Decode-only tps over 32 steps after the T=32768 prefill, so any
KV-cache scaling regression (sdpaDecode2Pass routing, residency
pinning, etc.) surfaces.
Measured on M5 Max / Qwen3.5-35B-A3B-4bit (single run):
Prefill T=32768: 243.829s -> 134.4 tps batched
Decode T=1 after T=32768 prefill: 2.215s over 32 steps -> 14.45 tps
The decode drop from 96 tps (empty KV) to 14.45 tps (32K KV) is the
expected 10-attention-layer KV-bandwidth tax — 30 of 40 layers are
GDN with O(1) state, but the 10 full-attention layers must walk the
linearly-growing K/V caches every decode step.
CI's macos-latest is on Xcode 16.4 which lacks the async `MTLCommand- Buffer.completed()` extension Apple added in Xcode 26.x. Local dev box on Xcode 26.5 ships `completed()` AND flags the long-standing sync `waitUntilCompleted()` as "unavailable from asynchronous contexts" in async test functions. So neither single API works on both toolchains. Adds `MTLCommandBuffer.awaitCompletion()` to `Tests/Helpers/CommonTest- Helpers.swift` — bridges `addCompletedHandler` through `withChecked- Continuation`. Works on every Metal SDK from macOS 10.11 onward, no compiler complaints on either toolchain, no real-thread blocking. Replaces 11 call sites in: - Tests/ModelIntegrationTests/Text/Qwen35MoEBenchIntegrationTests.swift - Tests/ModelIntegrationTests/Text/Qwen36TextIntegrationTests.swift - Tests/ModelIntegrationTests/Text/SpecDecodeVerifyCostBench.swift Pre-existing CI break that surfaced after the metaltile dependency PRs landed in #193-#201 unblocked compilation past the previous "MetalTileKernels has no member ..." cascade.
The mt_moe_gather_qmm_mma_int4_bm64_mpp_bf16 kernel uses MPP cooperative tensor intrinsics (simdgroup_matrix on cooperative-tensor types) that require Apple GPU Family 9 or later. On Family 7 (M1, CI runner) the kernel compiles but emits zeroed output, producing cos=0.0 vs the bm16 and m1 references. The kernel is independently covered by upstream metaltile's GPU correctness suite on adequate hardware. Gates all 7 bf16 cells in MoEBgemmBm64MppTests with .enabled(if:) using a static MTLDevice.supportsFamily(.apple9) probe. M5 Max (Family 9+) runs untouched; M1 runners skip with a clear reason string.
Merged
- Drop `let _ = resultFlat` immediately before `return resultFlat` in Qwen35DecoderLayer.decodeMany; it was a stale leftover. - Drop unused `let dt = hFlat.dtype` + `let dtBytes = dt.byteSize` pair in both Qwen35GDNLayer.decodeMany and Qwen35AttentionLayer.decodeMany.
The batchedQkvQ*Fast and batched4*Fast wrappers derive `in_dim` from `wQ.shape[1] * 8` (or `wA.shape[1] * 8` for the 4-way) and pass that single value to the kernel. The kernel walks every weight tensor at that packed width, so a mis-packed wK/wV (or wB/wC/wD) would silently miscompute its branch's output rather than fail at dispatch time. Add an explicit precondition that all packed widths match the leading weight tensor, with an error message naming the mismatched widths so the failure mode is obvious if it ever fires. Affects: - Ops.batchedQkvQgemvInt4Fast - Ops.batchedQkvQmmFast - Ops.batched4QgemvInt4Fast - Ops.batched4QmmFast
MoELayer and Qwen35{AttentionMixer,GDNMixer,MoEFFN} each hold a growing
set of per-instance scratch tensors that are mutated by `decode` /
`forward` / `forwardMany` without internal synchronisation. The
implicit contract — one inference per layer instance at a time — is
satisfied today by the per-token serial model loop, but is not
expressed in the API.
Make it explicit on every affected class header. Also tighten the
contract note on `MoELayer.decodeGPURouter` to call out that the
returned tensor aliases a scratch the GPU is still writing to via the
just-committed cmd, so a re-entrant call would race the in-flight
work.
GDNStateCache.snapshot() and ConvStateCache.snapshot() both return a single per-instance scratch tensor that is overwritten on subsequent calls. The current spec-decode driver does snapshot → verify → restore once per layer per verify step so this is safe in practice, but the contract is invisible to a caller reading just the method signature. Spell it out: a second snapshot() before the matching restore() will trample the prior contents. Future nested or concurrent snapshot usage would need a different design (per-call allocation, or a small ring).
ekryski
approved these changes
May 27, 2026
Collaborator
There was a problem hiding this comment.
@TheTom We may decide to drop these in future (likely). Fine for now as we're developing so let's leave it and the other file in this PR.
But really this is what ffai bench is for. I don't want to get into the habit of abusing integration tests and benching infrastructure. Leads to lost of complications and overhead that we had to deal with in my mlx-swift-lm fork.
If functionality is missing in the benchmark runner, I'd prefer to see us land updates there so it can be used for benching the way we need.
Collaborator
There was a problem hiding this comment.
This file is humongous. We should see if we can shrink it to make it more efficient and/or split out into separate files grouped either by model family, architecture layer, or ops or something. Not a blocker here. Something noted for future.
Remove all "ITER N" / "(PR10)" / "PR10" markers from comments in the model and op-wiring code. These were references to an authoring agent's internal iteration log and don't help a future reader — they made comments harder to skim and left a trail of stale cross-refs. Where the surrounding comment depended on the marker for grammar or context, tighten the prose so the remaining sentence still reads cleanly. No behavioural changes; the comments still describe the same fast-paths, gating env vars, and dispatch shapes.
For files Tom co-authored (modifications on top of an existing Eric- copyrighted file), add "and Tom Turney (@TheTom)" to the existing copyright line. For files Tom wrote from scratch, mark Tom as the sole copyright holder. SpecDecode.swift, CacheSnapshot.swift, and DrafterTests.swift were missing the standard Apache 2.0 boilerplate — add it now. No code changes.
PR #10's regression was the wrapper accepting blocks=2 (kernel pass-2 hardcodes a 32-lane block-merge reduction, so anything below 32 silently produces zero output). PR #14 closed it with a `blocks >= 32 && blocks % 32 == 0` precondition. The other half of the contract — that at the minimum legal blocks=32, two-pass output matches the canonical single-pass kernel within tolerance — had no host-side test. Add one. metaltile's per- kernel GPU correctness test would not have caught the dispatch- routing bug; the wrapper is what needs covering.
ekryski
added a commit
that referenced
this pull request
May 27, 2026
CI (Xcode 16.4 / Swift 5.10) was emitting ~70 warnings the local toolchain (Xcode 26.5 / Swift 6.3) doesn't surface. Mechanical cleanup pass — every change is a syntactic let-mutability fix, a deletion of a dead binding, or a one-line typo correction. No behavioural change. Bulk fix — Accelerate ILP64 BLAS interface - Package.swift: add ACCELERATE_NEW_LAPACK + ACCELERATE_LAPACK_ILP64 swiftSettings on the FFAI target. cblas_sgemm / cblas_sgemv were deprecated in macOS 13.3 in favour of these headers; without the define every audio DSP file that wraps a BLAS call (DeepFilterNet, DeepFilterNetDSP, LFMAudio, GraniteSpeech, vision tower CPU attention, …) emitted a deprecation warning per build pass. Per-file mechanical edits (~62 sites) - var → let where the variable is never mutated (28 sites; LFMAudio tuple destructurings dominate, plus DACVAE, DeepFilterNet, FireRedVAD, Sortformer). - Delete unused immutable let-bindings with side-effect-free RHS (22 sites; DeepFilterNet 13, FireRedVAD 3, Parakeet 2, FastVLMVision 2, LFMAudio 1, Qwen3xText 1). - DeepFilterNet linInW / linOutW: `if let X = weights["…"]` where X was never used — rewrite as `weights["…"] != nil` existence check; the body re-loads via groupedLinear with the same key. - Parakeet padded: pure-RHS array build never read, deleted entirely. - GraniteSpeech copyF32ToTensor: prefix `_ =` on three withUnsafeBytes calls whose results weren't consumed. - FireRedVAD pickle decoder: drop duplicate `case 0x28:` and `case 0x32:` lines that were already handled by previous patterns. - FastVLMVision: `M_SQRT1_2` is deprecated; use `0.5.squareRoot()`. - GPTOSSText: qualify `ReasoningLevel.none` so the compiler doesn't resolve to `Optional<…>.none` instead. - CacheSnapshot: remove the dead `ConvStateCache` cast branches in both snapshotAll and restoreAll — `ConvStateCache` doesn't conform to `LayerCacheProtocol`; only `Qwen35GDNLayerCache.conv` and `Mamba2LayerCache.conv` wrap it, and they're already handled by separate branches. The `LayerCacheSnapshot.conv` enum case stays in place (it's public API). - Loader/Model.swift: strip a stale rebase note from PR #14 that doesn't add information at the new HEAD. Out of scope here - ~119 Sendable / non-sendable closure capture warnings in CPU DSP + vision/audio CPU helpers. Need per-call-site review (the captures ARE thread-safe at runtime — DispatchQueue.concurrentPerform with disjoint indices — but Swift 6 strict concurrency flags them). Tracked as a follow-up commit on this branch. - ~5151 MSL `unused variable` warnings in Sources/MetalTileSwift/Resources/kernels/*.metal — codegen output from metaltile, not FFAI source. Tracked separately in session-plan.md's queued items. Verified - make build green - make format-check green - No swift-format / lint regressions
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Strategic re-port of PR #9 + PR #10 (Everything Bagel 1 + 2) onto the post-AURA-port
dev. The original PRs stay open for posterity; this branch is the cleanly-commented, scope-audited replacement that targets the updated layout (Eric'sOps//KVCache//Loader//Generation/split) and conforms to the repo's commit-hygiene + conventional-commit rules.Measured perf
M5 Max / 128 GB (Apple9) — Qwen3.5-35B-A3B-4bit:
M2 Pro / 32 GB (Apple8) — Qwen3.5-2B-4bit (dense GDN-hybrid, same engine class minus MoE-specific paths):
Same engine class (
m.qwen35) on both, so the bench exercises every wiring landed: GPU MoE router (decode + prefill — M5 Max only since M2's 32 GB can't load the 35 B-A3B MoE checkpoint), fused Q+K+V projection at decode + prefill, chunked GDN recurrence, fused conv1d+silu+cast, rmsNormQgemvInt4Fast finalNorm, ropePartialManyTwo, kvCacheUpdateKVMany, sdpaMulti d=256 (the silent-miscompute fix), sdpaDecode2Pass, MTLResidencySet weight pinning.Depends on metaltile PRs
Six new fused kernels + two pre-existing kernel PRs need to merge upstream before consumers without our metaltile fork can build this branch. We regenerate locally from
tom/feat/batched-qkv-qmmand ship the build artefacts via the standardGenerated/pipeline; no Swift change is required on this side once they land.mt_sigmoid_scalar_fma_residual— sigmoid-gated FMA with residual fold (post-MoE-FFN).ffai_sdpa_multi_tree_mask— tree-causal SDPA for spec-decode tree verify.ffai_sdpa_bidirectional d80 + d96— DiNOv2 / Idefics3 / Phi-3.5-vision head-dim coverage.ffai_sdpa_multi_d256— Qwen3.6-A3B full-attention layers (closes the silent-miscompute when routed through the d=128 kernel).rope_llama_many,kv_cache_update_many,conv1d_causal_step_silu_cast_many.batched_qkv_qmm_fast,batched_4_qgemv_fast,batched_4_qmm_fast.gated_rms_norm_qgemv_int4_fast— GDN tail fusion.moe_down_swiglu_accum_int4_chain8— 3-stage MoE per-expert path.All 8 are CLEAN + MERGEABLE with green CI.
What's in this PR
Foundation.
MTLResidencySetpin helper onDevice, applied atModel.load+KVCache/GDNStateCache/ConvStateCacheinit so the driver stops re-validating per-allocation residency on the thousands of decode-step dispatches.defaultMaxCommandBufferCountbumped 16 → 64 so the Metal driver pipelines more cmd buffers in flight before applying backpressure.Single-encoder / paired-dispatch Ops.
scalarFMA+Many+Chain8,sigmoidScalarFMAResidual,dequantGemvInt4Two/Three/Four/Many,ropePartialTwo/Many/ManyTwo,gatherTwo,rmsNormRowsTwo,kvCacheUpdateKV/Many,siluCastF32PlusCastF32Two,castToF32Two/Three,swigluMany,gatedMixerNormMany,conv1dCausalStepSiluCastMany. Each pairs / batches N kernel calls onto one encoder, eliminating begin/end overhead at the rate we hit decode (thousands per token).Fused-kernel wrappers.
moeRouterTopK+Many,dequantGemvInt4ExpertIndexed+Many(the GPU MoE router pair — biggest decode T=1 lever),sdpaMultiTreeMask,sdpaDecode2Pass(with theblocks ≥ 32 and multiple of 32precondition PR #10 silently violated),gatedDeltaPrepChunk,batchedQkvQgemvInt4Fast+batchedQkvQmmFast,batched4QgemvInt4Fast+batched4QmmFast,rmsNormQgemvInt4Fast,gatedRmsNormQgemvInt4Fast,moeDownSwigluAccumInt4Chain8.Model wirings.
Qwen35AttentionMixer.forward(batched QKV + 2-pass SDPA + scratch caching),Qwen35AttentionMixer.forwardMany(batched QKV qmm + sdpaMulti d=128/d=256 + ropePartialManyTwo + kvCacheUpdateKVMany),Qwen35GDNMixer.forward(batched 4-projection + fused silu+cast2 + residency pinning),Qwen35GDNMixer.forwardMany+forwardManyChunked(chunked recurrence + conv1dCausalStepSiluCastMany + gatedMixerNormMany),Qwen35Model.finalNorm+lmHead(fused rmsNormQgemvInt4Fast across decode / prefill / VLM embed-input paths),Qwen35DenseMLP+Qwen35MoEFFNshared expert (fusedOps.swiglu),Qwen35MoEFFN.forward(shared-encoder batching + sigmoidScalarFMAResidual),MoELayer.decode+MoELayer.decodeMany(GPU router fast path defaulted ON + bm64 default atmTotal ≥ 1024). Cross-familyOps.swigluadoption in Granite4 / LFM2 / Jamba dense MLPs.Speculative decode.
Drafterprotocol +NGramDrafter+NeverDrafter+NGramTreeDrafter+ tree-causal mask +DraftTreeNode.verify(20 host-side unit tests).Qwen35Model.forwardManyAllLogitsfor per-position verify.GDNStateCache/ConvStateCachesnapshot/restore+ compositeCacheSnapshotfor rejection rollback.SpecDecode.generateGreedy— greedy spec-decode driver against anyDrafter.Audit findings worth calling out
sdpaMultiwas a silent-miscompute on Qwen3.6-A3B full-attention layers — dev only switched on dtype and ran the d=128 kernel athead_dim=256. Restored the(dtype, headDim)tuple-match + relaxedOpsValidation.validateSdpaMulti; runtime test pins both paths.addAndRmsNormon dev supersedes PR #10'saddRmsNorm+addRmsNormRows— same kernel, tuple-return signature,nRows-aware.rmsNormResidualis a distinct op (single-output fused norm) and is NOT equivalent to PR #10'saddRmsNorm— verified by reading both kernel bodies.sdpaDecode2Passpass-2 kernel hardcodes a 32-lane block reduction (bn = 32,block_chunks = blocks / 32), soblocksMUST be a multiple of 32 and ≥ 32. PR #10's wrapper letblocks=2through; output was silently zero. Wrapper precondition catches it now and the runtime test exercisesblocks=32, nKV=64against single-passsdpaDecode.QuantizedLinear.callMany's 4-bit fast path needsoutDim % 32 == 0(the BN tile ofdequantGemmDynamicM). Without the guard, Qwen3.5's 4-bitshared_expert_gateatoutDim=1tripsOps.dequantGemmDynamicM: nOut (1) must be multiple of 32. Fall-through to the per-rowdequantGemvloop is bit-identical at that shape.Deferred per spec (not in this PR)
gatedRmsNormQgemvInt4Fastattention wiring — Op landed but wiring stays out; PR Everything Bagel 2 #10 documented -7% decode regression atout_dim=2048(TG-mem staging caps occupancy below the unfused chain).batched_4_qmm_fastGDN-input wiring — Op + wrapper landed but stays unwired; PR Everything Bagel 2 #10 documented workgroup imbalance at the Qwen3.6 GDN shape (out_c/d=16 ≪ out_a/b=2048wastes 99% of TGs).OpsICB(Day 2 plumbing) — 0.55× vs direct rebound at Qwen3.6's 40-layer MLP shape. ICB wins on many small non-dependent dispatches; not our workload.Tests
Host-side unit tests for every new Op (CPU-truth or pair-equivalence in
Tests/FFAITests/Ops/), 20 Drafter / tree-verify tests inTests/FFAITests/Generation/, and integration benches inTests/ModelIntegrationTests/Text/:Qwen35MoEBenchIntegrationTests— decode T=1, T=128 / T=512 / T=2048 forwardMany, plus T=32K long-context (prefill + decode-after-prefill). Gated on local Qwen3.5-35B-A3B-4bit checkpoint.Qwen36TextIntegrationTests.forwardManyBench{32,128,512,2048}— Eric's harness extended with T=2048 long-context case.SpecDecodeBenchTests+SpecDecodeVerifyCostBench— greedy spec-decode parity vs baseline + verify-shape cost crossover.Cross-hardware: all 47 new Op tests + 20 Drafter tests pass on both M5 Max (Apple9 / 128 GB) and M2 Pro (Apple8 / 32 GB). Build green on both. Decode + prefill perf measured on both — numbers in the table above.
Pre-existing dev test issues fixed along the way so CI goes green:
PSOCacheTests.swiftMTLComputePipelineStateSendable error under Xcode 16.4 (@preconcurrency import Metal);FireRedVAD.swift:1202BINUNICODE8 length expression typing-out under strict concurrency (split into two 32-bit halves).Known pre-existing dev test failure:
Soprano-80M synthesizeintegration test hitsKVCache: dtype mismatchprecondition. Reproduces onwaffles/devtip with this branch's changes stashed, so unrelated to the migration.Cross-reference