Chunking Qwen3.5's gated DeltaNet for 1.5x faster prefill on Apple Silicon

Qwen3.5 is a hybrid model. In the 0.8B variant, 6 of the 24 layers are ordinary attention and the other 18 are gated DeltaNet layers, a form of linear attention that keeps a fixed-size state matrix per head instead of a growing KV cache. In Alloy, our kernel compiler and inference stack for Apple Silicon, the model prefills a 4096-token prompt at 5988 tok/s and decodes at 293 tok/s on an M4 Max.

A delta-rule layer rewrites its state after every token, which makes the computation serial along the sequence: awkward for a GPU, and worse for a compiler that assumes the model is a pure dataflow graph. This post walks through how Alloy runs it: why the whole layer ships as one custom op, how decode runs the recurrence serially in a single kernel, how prefill breaks the seriality with a chunked two-stage kernel, and what the chunked kernel measures against the serial one.

Results

qwen3.5:0.8b, Q8_0, synthetic prompts on an M4 Max. The serial arm forces the handler's fallback kernel for prefill; everything else is identical:

At production depth, where both kernels do 4096 tokens of real work, the chunked kernel is worth 1.54x end to end. The 512-token row measures something different: the serial kernel cannot take part in Alloy's padded-plan machinery, so a short prompt through it pays the full padded scan, and the gap balloons to 7.8x. The mechanics behind both numbers are covered below. Decode does not move when the arm switches, since decode never runs the chunked path.¹ Inside the default prefill chunk the two DeltaNet kernels are the largest single item at 34.0% of the chunk's 671.5ms of GPU time, and at decode the recurrence costs 0.44ms of a 4.0ms step.²

What is the gated delta rule in Qwen3.5?

Each linear-attention head carries a state matrix S with one row per key dimension and one column per value dimension, 128 by 128 held in fp32.³ Per token, after a short causal convolution and an L2 normalization of Q and K, the layer applies a decay to the state, computes a correction from what the state currently predicts for this token's key, writes the correction back, and reads the output:

The state is written by every token and read by every later one, so token t cannot start until token t-1 has finished.

Why can't torch.compile run a recurrence like this?

The first reason is structural. Alloy executes the model through Hugging Face's transformers modeling code under torch.compile, and the transformers DeltaNet layer updates its convolution and recurrent state through in-place writes on cache tensors. Ahead-of-time autograd functionalizes the graph before any backend sees it. Probed on a single GDN forward, our backend receives 1,202 ops with no mutation among them and three outputs: the layer's actual output plus the two new state values.⁴ The write has become a Python epilogue that copies those outputs into the cache after each call. A compiled layer does advance its state this way, but only while Python runs around every call, and Alloy's fast path removes Python from the loop: after the first runs, a model executes as a recorded C++ dispatch plan. So Alloy ships the entire layer body, convolution through gated normalization, as one custom op whose state tensors are declared mutable. The write is a kernel inside the graph, and the plan records it like any other.

The second reason is cost. Traced op by op, the chunked delta-rule math turns into a long tail of small per-chunk, per-head matmuls, copies, and concatenations. No fusion pass reassembles work at that granularity. The op boundary has to enclose the whole recurrence, and what sits behind the boundary has to be written by hand.

The strongest stock alternative, PyTorch's own MPS backend, does not close the cost problem either. We measured it on the pure delta-rule functions from transformers at this model's dimensions, cache machinery stripped away, which sets the state question aside entirely and hands stock PyTorch its best case.⁵ Compiling the chunked prefill formulation does not get through codegen. Inductor fuses the chunk math into a single generated kernel with more inputs than Metal permits a kernel to bind:

InductorError: SyntaxError: failed to compile
    kernel void generated_kernel(
        device float* out_ptr30,
        constant float* in_ptr0,
        constant float* in_ptr1,
        ...
program_source:897:17: error: number of constant buffers exceeds maximum supported (31)

Eager MPS, the fallback, runs the chunk math at 19.7ms per layer against the Alloy kernel's 12.7ms, with every integration cost excluded: served for real, the eager path would break the compiled graph around all 18 layers and give up the padded-plan machinery. Decode does compile, at 0.12ms per layer for a recurrence Alloy runs in 0.024ms; across 18 layers that difference alone would grow the 4.0ms decode step by over 40%.

How does decode run a sequential update on a GPU?

At one token per step there is nothing to parallelize across time, so the decode kernel accepts the serial order and parallelizes over the state instead. One thread owns one column of one head's state, a 128-float slice that stays in registers for the whole call. Per token, the thread applies the decay, accumulates its share of the three dot products the update needs in a single pass over the slice, and applies the rank-1 correction. Between tokens of the same call the state never touches device memory.

Run this way, the whole recurrence is 0.44ms of a 4.0ms decode step. The quantized weight projections around it cost several times more, so the recurrence is not the decode bottleneck. What the kernel protects is the state itself: it stays resident and it stays fp32, which keeps a long generation's accumulated rank-1 updates from drifting.

How do you parallelize a recurrence for prefill?

Prefill hands the layer thousands of tokens at once, and pushing them through a serial update leaves most of the GPU idle. The standard escape is the chunked formulation of linear attention, which Alloy maps onto Metal in two kernels.

Stage one cuts the sequence into chunks of 8 tokens and resolves each chunk's internal causality in closed form, independently and in parallel across every chunk and head. The closed form exists because of a structural fact: within a chunk, token-to-token interactions form a strictly lower-triangular matrix $A$, and a strictly triangular matrix is nilpotent, $A^C=0$ for chunk size $C$. The inverse that unrolls the in-chunk recurrence is therefore a finite series,

which repeated squaring evaluates in ${\log }_{2}C$ rounds of small matrix products on the GPU's simdgroup matrix units. From $T$, stage one precomputes a set of per-chunk summaries.

Stage two is the part that remains serial, but its recurrence steps once per chunk, not once per token. With stage one's summaries in hand, chunk $c$ reduces to

where $U$, $W$, $M$, $Q^g$, $\hat{K}$, and the decay $\gamma$ all come from stage one. The first line is the token-level correction $V-S^{\top }K$ lifted to a whole chunk, and the shape of the whole system is the same delta rule with the serial chain $C$ times shorter and every link a dense matrix product instead of a vector rewrite. The scan walks the chunks in a runtime loop, holds its state tile in the 32KB of threadgroup memory Metal provides, and blocks the value dimension so the tile fits.

Stage two also carries the two decisions that are about Metal rather than math. Its chunk loop is not unrolled, because an unrolled loop emits shader code proportional to the sequence length and Metal's compile time grows superlinearly with shader size; at production chunk sizes an unrolled scan stops compiling in reasonable time. And both kernels put the sequence on the launch grid's first axis, which lets Alloy's padded-prompt machinery shrink a dispatch to the tokens a request actually filled. The serial kernel cannot be shrunk this way. Its grid indexes the state, not the sequence, so a short prompt inside a padded plan pays the full padded length. The 7.8x at 512 tokens prices exactly that: the serial kernel's incompatibility with the padded plan, not the chunk math. The 1.54x at 4096 tokens, where padding is a rounding error, is the algorithmic gap between the two formulations.

Why keep the serial kernel at all?

Decode runs it every step: at sequence length 1 there is no chunk to build, and the serial kernel is the decode path by construction. Ordinary prompts never fall back to it, because prefill buckets are chunk-aligned by construction and the chunked kernels mask past the real prompt length.

Speculative decoding is the deliberate exception. Draft verification runs the serial kernel even at chunk-friendly widths, because the chunked kernel's triangular inverse reorders the f32 arithmetic, and the reordering shifts logits enough to flip an occasional near-tie token. Verification has to reproduce exactly what decode would have computed, so it uses the same serial numerics decode commits.

Do linear-attention models need hand-written kernels?

On this evidence, yes, at both ends of the sequence-length range. torch.compile alone leaves the state write hanging on a Python wrapper and the chunk math either unbuildable or slow on Metal, and neither has a pass-level fix; the recurrence has to live behind one op with kernels underneath. The prefill kernel returns 1.54x at production depth; short prompts gain 7.8x in the shipped system, most of it because the fallback cannot cooperate with padded plans. The serial kernel stays for decode and for verification, which need its exact numerics.

Both kernels ship in Alloy: alloy serve -m qwen3.5:0.8b runs the hybrid with them in place, and alloy profile qwen3.5:0.8b prints the per-kernel breakdown these numbers came from.

The MoE half of this model family has its own kernel story: qwen3.6:35b picks 8 of 256 experts per token, and grouping tokens by expert on a GPU without atomic operations turns out to be a counting-sort problem. We'll write about this next.

Footnotes

1. Decode across the arms and depths: 301.3 vs 307.0 tok/s at depth 512, 293.3 vs 285.9 at depth 4096 (chunked vs serial arm, tg128, reps 5).
2. Consistency check on the profile: the chunk's kernel-time sum is 671.5ms, and the bench wall clock for the same 4096 tokens is 684ms (4096 over 5988.4 tok/s), so the per-kernel table accounts for nearly all of the wall and the kernels barely overlap.
3. The GGUF weight layout has a trap here: llama.cpp stores the value heads grouped so that value head h pairs with key head h modulo the key-head count, not the repeat-interleaved pairing the transformers code uses. Pair them the transformers way on these weights and the model emits fluent nonsense.
4. Probe backend on a standalone Qwen3_5GatedDeltaNet at the 0.8b dimensions, cache primed eagerly, one compiled forward: the functionalized graph handed to the backend contains 1,202 ops, zero in-place ops, and three outputs (layer output, conv state, recurrent state). The same harness shows the epilogue working: eager and compiled runs advance the cache identically, including the recorded-plan call.
5. torch_chunk_gated_delta_rule and torch_recurrent_gated_delta_rule from transformers, at B=1, 16 heads, 128-dim keys and values, fp32, the functions' default chunk size of 64, torch 2.11 on the same M4 Max. Ten timed iterations after warmup, MPS synchronized around the loop. The setup is deliberately generous to stock PyTorch: the isolated function skips the cache handling entirely, receives exactly chunk-aligned input, and pays none of the graph breaks or padded-plan integration a served model would.