Chapter 18 — The SwiGLU feed-forward block

status: polished · path: Muse Glimmer, pinned Muser tree

Prerequisites: Ch 2 (SIMD groups, simd_sum, threadgroup memory), Ch 5 and Ch 6 (blocks, scales, the Q4_K super-block), Ch 12 (the dual-eps tail that produces this chapter’s input), Ch 13 (the matvec family and the pinned ggml kernels), Ch 17 (attention just closed into the residual stream). The FFN and its gated activation are taught from zero here.

18.1 What it computes

Attention has just folded its result back into the residual stream and handed the layer a vector. What does the layer do with it next? Half the layer’s work is still ahead, and it is the half that owns most of the weights on disk — so before any kernel, the question to settle is what shape that computation has and what it does to the vector it is given.

Each Muse Glimmer layer does two things to the residual stream: attention (Ch 13–Ch 17) mixes information across tokens; the feed-forward network (FFN) transforms the vector within a token — one position at a time, no cross-token mixing at all. Concretely it is two matvecs and a nonlinearity between them:

x : [6656]                        (hidden_dim — the ffn_norm'ed stream, Ch 12)
      │  gate projection   W_gate : [6656 → 19968]   (Q4_K)
g : [19968]                       (intermediate_dim)
      │  SiLU·Hadamard with the up branch
h : [19968]                       ("ffn_mid")
      │  down projection  W_down : [19968 → 6656]   (Ch 19)
out : [6656]

The vector is blown up 6,656 → 19,968 — a 3× wider “thinking space” (the checkpoint declares muse-glimmer.feed_forward_length, config.rs:180) — and squeezed back down in Ch 19. The widening is where most of the model lives: the two projections this chapter covers are 74.76 MB each of Q4_K per layer (§18.7), against 48.4 MB for the entire attention block (Ch 17 Figure 17.2).

The combination rule between the two branches is SwiGLU — Figure 18.1 shows its wiring:

h = silu(W_gate · x)  ⊙  (W_up · x)

Read the symbols: W_gate · x and W_up · x are two independent [6656 → 19968] matvecs over the same input x; silu(·) is the activation of §18.4 applied only to the gate branch; ⊙ is the element-wise (Hadamard) product, (a ⊙ b)[j] = a[j] × b[j], no mixing across the 19,968 coordinates.

18.2 Why it exists — the “thinking” half of the layer

Attention is the block’s mechanism for gathering context: it decides which earlier tokens this token looks at. The FFN is the mechanism for reasoning over what was gathered: a learned, position-independent transformation applied to the vector attention assembled. Every token at every layer walks the same FFN weights; what differs is the vector it walks in with. That is the “thinking” framing — and the size framing is starker: at ~30 B parameters, the FFN pair-plus-down is roughly 224–259 MB of the ~273–307 MB per-layer weight read (§18.7), so when Ch 1 said decode is ~99 % reading weights, this is where most of the weights are.

Why gated? A plain FFN (out = down(act(up(x))), the original transformer’s shape [arxiv:1706.03762]) applies its activation unconditionally per element — for feature j, “should this feature fire?” and “what does it carry?” are the same number. A gated FFN splits them into two learned projections: the gate branch answers should it fire, the up branch answers what does it carry. The silu turns the gate into a smooth on/off ramp. This structural claim is standard SwiGLU motivation [arxiv:2002.05202]; the quality delta for Muse Glimmer specifically is [unverified] here — this book does not retrain the model to A/B its own architecture. The cost side, however, is exact arithmetic: gating means two widening matrices instead of one, ~50 % more FFN parameters than the ungated shape — and it is why W_gate and W_up both stream past every token.

flowchart LR
    x["x [6656]<br/>ffn_norm'ed stream"]
    wg["W_gate<br/>[6656→19968] Q4_K"]
    wu["W_up<br/>[6656→19968] Q4_K"]
    g["g [19968]"]
    u["u [19968]"]
    sg["silu(g)"]
    had["⊙ Hadamard"]
    out["ffn_mid [19968]"]
    x --> wg --> g --> sg --> had
    x --> wu --> u --> had
    had --> out

Figure 18.1: The SwiGLU dataflow. The fusion this chapter covers collapses the two projections, the silu, and the Hadamard into one kernel; the down projection that consumes ffn_mid is Ch 19.

18.3 The matrix operation, by hand

Trusting a fused kernel gets easier once you have done its arithmetic by hand. So: what exactly is computed for a single output coordinate, and where is the redundancy that a fusion could take away?

For output element j (0 ≤ j < 19,968):

g_j = Σ_{i=0..6655} W_gate[j, i] · x[i]      one dot product over hidden_dim
u_j = Σ_{i=0..6655} W_up  [j, i] · x[i]      another, independent, same x
h_j = silu(g_j) · u_j                        the gated combination

Two dot products over the same x, then a pointwise combine. A toy worked example with hidden = 2, one output element, invented numbers:

x = [1.0, −2.0]
W_gate row = [0.5, 0.25]  →  g = 0.5·1 + 0.25·(−2) = 0.0
W_up   row = [2.0, −1.0]  →  u = 2.0·1 + (−1.0)·(−2) = 4.0
silu(0.0) = 0.0 · σ(0.0) = 0.0 · 0.5 = 0.0
h = 0.0 · 4.0 = 0.0          ← gate fully closed: nothing flows

Had the gate row been [1.0, 0.0] instead, g = 1.0, silu(1.0) ≈ 0.731, and h ≈ 0.731 · 4.0 ≈ 2.92 — the same up-value flows, scaled by how open the gate is. That decoupling is the whole idea.

The observation that motivates a fused kernel: both dot products read the same x[k]. In an unfused pair of matvecs, x is fetched twice and both 19,968-wide intermediates (g and u) make a round trip through memory. One kernel that loads x[k] once and updates two accumulators in lockstep eliminates both — Figure 18.2 shows the loop shape.

18.4 SiLU — the sigmoid linear unit

Two branches go in and one comes out, and the function on the gate branch is the only nonlinearity in the whole block. It decides how much of the up branch survives. So it is worth a section of its own: what is it, and how does it behave at the edges where it will actually be asked to decide?

The activation on the gate branch is SiLU (a.k.a. Swish, [arxiv:1710.05941]):

silu(x) = x · σ(x)  =  x / (1 + e^(−x))        (σ from Ch 17 §17.1)

Behavior at the extremes — like ReLU at the ends, unlike it in the middle:

Table 18.1: SiLU vs ReLU at five points (computed by hand from the formula). Large positive x passes through nearly unchanged; large negative x is suppressed — but small negative x lets the gate subtract a little, not merely go silent.

Two properties matter downstream. First, SiLU is smooth and its gradient is nonzero everywhere (ReLU has a dead zone for x < 0) — whether that is the reason gated FFNs ship SiLU is [unverified] here; we inherit the checkpoint’s choice. Second, it is cheap but not free: one exp, one add, one divide per element. At 19,968 elements per layer it runs 19,968 times — fully parallel, no reduction, and in every kernel this chapter quotes it is folded into the final write.

The CPU oracle states the whole combination in two loops (crates/muser-engine/src/reference.rs:511-513):

#![allow(unused)]
fn main() {
for (a, b) in ffn_a.iter_mut().zip(ffn_b.iter()) {
    *a = silu_fast(*a) * *b;
}
}

The helper it calls is silu_fast(x) = x / (1.0 + (−x).exp()) — this section’s formula written as a one-liner, inherited from Ferrite along with the rest of the shader lineage. We kept the provenance: crates/muser-engine/src/quant/helpers.rs:70-74, tracked through the extraction manifest [docs/extraction-manifest.md].

18.5 The Metal kernel — ffn_q4k_gate_up_silu_4r2s

There are two gate+up routes in the tree, and which one runs is the chapter’s real story (§18.6). The fused kernel first — it is the cleanest expression of the SwiGLU fusion, a wholesale port of Ferrite’s accepted 897a6256b kernel, and both the call site and the wrapper say so in their own comments (decode.rs:5823-5825, ffn.rs:7-8). Below: the signature, the x-load, and the lockstep MAC, verbatim; the Q4_K decode helper it calls is summarized after.

// crates/muser-engine/src/shaders/ferrite/ffn_fused_tail.metal:485
// ── ffn_q4k_gate_up_silu_4r2s ─────────────────────────────────────────────
//
// 4-row-per-TG fused gate+up variant using the V4 x-load pattern.
// 64 threads (2 SIMDs × 32), 4 output rows per TG (2 per SIMD).
// Each thread loads x into local registers yl[16]/yh[16] via stride-4
// block sub-groups (same as matvec_q4k_f32_v4), reusing x for both
// gate and up weight rows.
//
// Activation selected at PSO build time via FC_FFN_ACTIVATION.
// dispatch_thread_groups( (ceil(i_dim/4), 1, 1), (64, 1, 1) )
//
kernel void ffn_q4k_gate_up_silu_4r2s(
    device const uchar* W_gate [[ buffer(0) ]],
    device const uchar* W_up   [[ buffer(1) ]],
    device const float* x      [[ buffer(2) ]],
    device       float* out    [[ buffer(3) ]],
    constant     uint&  rows   [[ buffer(4) ]],
    constant     uint&  cols   [[ buffer(5) ]],
    uint tgid [[ threadgroup_position_in_grid ]],
    uint lid  [[ thread_index_in_simdgroup ]],
    uint sgid [[ simdgroup_index_in_threadgroup ]])
{
    const uint n_blocks    = cols / 256u;
    const uint block_bytes = 144u;
    const uint row_bytes   = n_blocks * block_bytes;

    // 2 rows per SIMD, 2 SIMDs = 4 rows per TG
    const uint base_row = tgid * 4u + sgid * 2u;
    if (base_row >= rows) return;

    // V4 thread partitioning: 4 sub-groups of 8 threads for block stride
    const uint ix = lid / 8u;   // 0..3 — block stride index
    const uint it = lid % 8u;   // 0..7 — position within block
    const uint iq = it / 4u;    // 0 or 1 — half-block selector
    const uint ir = it % 4u;    // 0..3 — quarter within half

    // x-vector pointer: each thread reads 8 positions per block (stride-4 blocks)
    device const float* xp = x + ix * 256u + 64u * iq + 8u * ir;

    float yl[16], yh[16];
    float gate_sumf[2] = {0.f, 0.f};
    float up_sumf[2]   = {0.f, 0.f};

    for (uint ib = ix; ib < n_blocks; ib += 4u) {
        // Load x slice into registers (8 elements × 4 positions)
        float4 sumy = {0.f, 0.f, 0.f, 0.f};
        for (uint i = 0u; i < 8u; i++) {
            yl[i]     = xp[i];       sumy[0] += yl[i];
            yl[i + 8] = xp[i + 32];  sumy[1] += yl[i + 8];
            yh[i]     = xp[i + 128]; sumy[2] += yh[i];
            yh[i + 8] = xp[i + 160]; sumy[3] += yh[i + 8];
        }

        // Gate weight: 2 rows starting at base_row + sgid*2
        device const uchar* gate_blk = W_gate + (ulong)base_row * (ulong)row_bytes
                                       + (ulong)ib * block_bytes;
        q4k_v4_dual_row_mac(gate_blk, row_bytes, yl, yh, sumy, iq, ir, gate_sumf);

        // Up weight: same 2 rows
        device const uchar* up_blk = W_up + (ulong)base_row * (ulong)row_bytes
                                     + (ulong)ib * block_bytes;
        q4k_v4_dual_row_mac(up_blk, row_bytes, yl, yh, sumy, iq, ir, up_sumf);

        xp += 4u * 256u;  // advance by 4 blocks (stride)
    }

    // Reduction across all 32 threads in each SIMD
    const float gr0 = simd_sum(gate_sumf[0]);
    const float gr1 = simd_sum(gate_sumf[1]);
    const float ur0 = simd_sum(up_sumf[0]);
    const float ur1 = simd_sum(up_sumf[1]);

    if (lid == 0u) {
                                    out[base_row]      = apply_activation(gr0, ur0);
        if (base_row + 1u < rows)   out[base_row + 1u] = apply_activation(gr1, ur1);
    }
}

Line by line:

• Bindings (496–505). Two weight buffers (W_gate slot 0, W_up slot 1), the input x (slot 2), the output out (slot 3), and rows/cols as inline constants (19,968 / 6,656 on this model).
• Row ownership (511–513). base_row = tgid·4 + sgid·2: one threadgroup of 64 threads = 2 SIMD groups, each SIMD group owns two whole output rows. Four rows per threadgroup; the grid of §18.6 covers 19,968 rows.
• The V4 lane decomposition (515–522). The 32 lanes split into 4 sub-groups of 8 (ix), each striding a different Q4_K super-block (ib += 4); within a sub-block, iq/ir select which of the eight 32-element runs this lane dots. This is the same x-in-registers pattern as the QKV matvec family of Ch 13 — the header says so (“same as matvec_q4k_f32_v4”).
• The register x-cache (524–536). Each thread pulls its 32 x-slice elements (yl[16], yh[16]) and their quarter-sums sumy into registers once per block iteration. The four sumy partial sums exist for the Q4_K min-term subtraction (w = d·sc·nib − dmin·m, Ch 6).
• The lockstep MAC (538–548). This is the fusion. The same register-held yl/yh/sumy feed two calls to q4k_v4_dual_row_mac — once against the gate rows, once against the up rows at the same base_row:

                       ┌─► gate_sumf[r] += dequant(W_gate[base_row+r, k]) · x[k]
  x[k] (registers)  ───┤
                       └─► up_sumf[r]   += dequant(W_up  [base_row+r, k]) · x[k]

Figure 18.2: The lockstep MAC. Every x element is loaded from device memory once (into yl/yh) and consumed by both the gate and the up accumulator for two rows each — one load, four dot-product contributions.

q4k_v4_dual_row_mac itself (shaders/ferrite/_q4k_helpers.metal:34-88) decodes two consecutive Q4_K rows’ block: it unpacks the 6-bit scale strip with the 0x3F3F/0x0F0F/0xC0C0 masks, accumulates nibble·x products into four float4 lanes per row, and folds them into the d·(Σ…) − dmin·(Σ…) super-block epilogue — the deferred-scaling schedule of Ch 13, here applied to two rows and two matrices at once.

• Reduction and activation (551–560). simd_sum collapses each accumulator across the 32 lanes; lane 0 of each SIMD group writes the two finished rows via apply_activation(g, u). There is no cross-SIMD combine and no threadgroup_barrier — whole-row ownership per SIMD group means each group’s simd_sum is the final answer for its rows. The ownership implies synchronization freedom, the same argument as Ch 17 §17.7.

apply_activation is compiled, not branched. The helper is selected at PSO build time through a Metal function constant:

// crates/muser-engine/src/shaders/ferrite/ffn_fused.metal:15
// 0 = SiLU (default), 1 = GELU
constant uint FC_FFN_ACTIVATION [[function_constant(32)]];
constant bool HAS_FC_FFN_ACTIVATION = is_function_constant_defined(FC_FFN_ACTIVATION);

// Activation helper — compiled away at pipeline creation time (zero runtime cost)
inline float apply_activation(float g, float up) {
    if (HAS_FC_FFN_ACTIVATION && FC_FFN_ACTIVATION == 1u) {
        // … (GELU branch elided) …
    } else {
        // SiLU (default): g * sigmoid(g) * up
        return (g / (1.0f + exp(-g))) * up;
    }
}

Build the PSO without slot 32 and the dead-code eliminator removes the GELU path entirely; the shipped Muse kernel contains only the SiLU line — silu(g) · u, the chapter’s formula, as the kernel’s last instruction (ffn_fused.metal:14-30, the Ch 4 function-constant mechanism).

18.6 The Rust dispatch — and the opt-in flag story

Two questions decide what this section owes you. How does the host launch that kernel? And — since the tree carries both routes — which one actually runs when the engine serves a token? The first has a short answer. The second is this chapter’s war story.

The wrapper first:

#![allow(unused)]
fn main() {
// crates/muser-engine/src/metal/encode/ffn.rs:7
/// Ferrite 897a6256b Q4_K SiLU gate+up route: four rows per threadgroup,
/// two SIMD groups, with the input vector shared by both projections.
pub fn encode_ffn_q4k_gate_up_silu_4r2s(
    &self,
    encoder: &ComputeCommandEncoderRef,
    gate_weights: GpuByteView<'_>,
    up_weights: GpuByteView<'_>,
    input: &GpuBuffer,
    output: &GpuBuffer,
    intermediate_dim: usize,
    hidden_dim: usize,
) {
    let row_bytes = hidden_dim / 256 * 144;
    debug_assert_eq!(gate_weights.len(), intermediate_dim * row_bytes);
    // … (up/input/output length asserts elided) …
    self.bind(encoder, "ffn_q4k_gate_up_silu_4r2s");
    encoder.set_buffer(0, Some(gate_weights.metal()), gate_weights.offset() as u64);
    encoder.set_buffer(1, Some(up_weights.metal()), up_weights.offset() as u64);
    encoder.set_buffer(2, Some(input.metal()), 0);
    encoder.set_buffer(3, Some(output.metal()), 0);
    set_value(encoder, 4, &(intermediate_dim as u32));
    set_value(encoder, 5, &(hidden_dim as u32));
    encoder.dispatch_thread_groups(
        MTLSize::new(intermediate_dim.div_ceil(4) as u64, 1, 1),
        MTLSize::new(64, 1, 1),
    );
}
}

In prose: grid (19,968 ÷ 4, 1, 1) = (4,992, 1, 1) threadgroups of (64, 1, 1) threads — 4,992 × 4 = 19,968 rows covered exactly. The two weight views are bound at their offsets into the mmap’d GGUF (Ch 3); rows/cols ride as inline constants.

Now the fork. We had a kernel that was better on paper in every dimension we knew how to count — fewer dispatches, fewer bytes moved, the same algebra — and it arrived as a port of something Ferrite had already accepted and run. The obvious move was to make it the FFN path and move on. We expected to do exactly that. Instead: the fused kernel is opt-in, and the default is the unfused control. The gate at the call site shows both branches side by side, which is why it is worth reading whole:

#![allow(unused)]
fn main() {
// crates/muser-engine/src/decode.rs:5819
if self.ferrite_ffn_gate_up
    && layer.ffn_gate.layout.dtype == GgmlType::Q4_K
    && layer.ffn_up.layout.dtype == GgmlType::Q4_K
{
    // Port Ferrite 897a6256b wholesale: the four-row/two-SIMD-
    // group kernel reads the normalized input once for both Q4_K
    // projections and writes the final SiLU(gate) * up row.
    dispatch(command, |encoder| {
        self.kernels.encode_ffn_q4k_gate_up_silu_4r2s(
            encoder,
            layer.ffn_gate.view(&self.mapped_weights),
            layer.ffn_up.view(&self.mapped_weights),
            &self.activations.post_norm,
            &self.activations.ffn_gate,
            cfg.intermediate_dim,
            cfg.hidden_dim,
        );
    });
} else {
    // Exact upstream-matvec control and non-Q4_K fallback.
    dispatch(command, |encoder| {
        self.encode_projection(
            encoder,
            &layer.ffn_gate,
            &self.activations.post_norm,
            &self.activations.ffn_gate,
            1,
        );
        self.encode_projection(
            encoder,
            &layer.ffn_up,
            &self.activations.post_norm,
            &self.activations.ffn_up,
            1,
        );
    });
    dispatch(command, |encoder| {
        self.kernels.encode_silu_mul(
            encoder,
            &self.activations.ffn_gate,
            &self.activations.ffn_up,
        );
    });
}
}

ferrite_ffn_gate_up is set from the environment — MUSER_FERRITE_FFN_GATE_UP (decode.rs:1334) — so by default the engine takes the else branch: two pinned ggml matvecs (W_gate, W_up, the exact kernel_mul_mv_q4_K_f32 of Ch 13) plus one pointwise activation kernel. That third kernel is muser_silu_mul_inplace, which does gate[i] = silu(gate[i]) · up[i] — the same formula as the fused tail, materialized:

// crates/muser-engine/src/shaders/muse_reference.metal:4
kernel void muser_silu_mul_inplace(
    device float *gate [[buffer(0)]],
    device const float *up [[buffer(1)]],
    constant uint &count [[buffer(2)]],
    uint index [[thread_position_in_grid]]) {
    if (index < count) {
        float value = gate[index];
        gate[index] = (value / (1.0f + exp(-value))) * up[index];
    }
}

dispatched one-thread-per-element over 19,968 (ffn.rs:38-60).

So why is the better kernel the one that does not run? Because when the fused tail was put on the serving route, the numbers coming out of the model stopped matching the numbers the pinned comparator produced. Not wrong — different, and different by more than the public tolerance allows. That result is written down where it belongs, in the routing comment that governs serving decode:

#![allow(unused)]
fn main() {
// crates/muser-engine/src/decode.rs:2085
// The legacy one-token graph uses Ferrite fused residual/norm and
// gate-up kernels whose rounding diverges from the source-pinned
// llama Metal graph enough to breach public logprob tolerance.
// The one-row batch graph dispatches the exact pinned kernels and
// has the same KV transition, so it is the serving correctness
// path until each fused kernel independently passes full-logit
// parity.
}

That is the contract discipline in one comment. The fused kernel’s reduction order — two interleaved accumulators folded per super-block — is mathematically SwiGLU, but it is not bit- the same as llama.cpp’s graph of independent mul_mv nodes plus a pointwise silu-mul, and Muser’s public commitment is logprob parity with the pinned comparator.

Say it the other way round, because this is the sentence the rest of the book keeps coming back to. Floating-point addition is not associative, so changing when you add changes what you get; a contract written against a specific graph is therefore a contract against a specific order of additions, not against the algebra those additions approximate. A kernel that computes the same function by a different schedule is, by that measure, a different kernel — and the measure is the one the public promise is written in.

The lesson we kept is the uncomfortable one: fusion is never free, even when it is free, because what it spends is exactness, and here exactness is the product. The fused kernel therefore lives on as a qualified-off fast path, runnable under the flag for the teacher-forced/diagnostic route (encode_token) and never on the serving default. The dtype guards in the if are a smaller version of the same care — any non-Q4_K FFN tensor falls back automatically, and on the release artifact ffn_gate/ffn_up are Q4_K. We kept that evidence too: ffn_gate/up 6656->19968 q4k, crates/muser-bench/src/m16.rs:171-175.

18.7 The access pattern — the largest weight read in the layer

Where does the time go in this block? Not into the silu, and not into the Hadamard: both are pointwise, and both vanish into the write. It goes into dragging two very large matrices past the arithmetic units, once for every token the model emits. It is worth knowing exactly how large, and worth knowing what the fusion actually buys against that background — because the answer decides whether the flag we just described gave up something expensive or something cheap.

All arithmetic derived from the verified shapes, shown step by step:

  W_gate: 19,968 rows × (6,656 / 256 = 26 blocks) × 144 B = 19,968 × 3,744
        = 74,760,192 B ≈ 74.76 MB        (Q4_K: 0.5625 B/element)
  W_up  : same shape, same dtype          = 74,760,192 B ≈ 74.76 MB
                                          ────────────────────────
  gate + up pair per layer                ≈ 149.5 MB   ← read once per token
  (for scale: the whole attention block   ≈  48.4 MB;  o_proj alone 15.34 MB,
   Ch 17 Figure 17.2; the down projection is Ch 19: 74.76 MB Q4_K /
   109.03 MB Q6_K)

Activation traffic, fused vs unfused (per layer):

  FUSED (one kernel):
    read  x [6656] f32                    26,624 B  (registers thereafter)
    read  W_gate + W_up                 149,520,384 B
    write ffn_mid [19968] f32             79,872 B

  UNFUSED (two matvecs + silu_mul — the default):
    read  x twice                          53,248 B
    read  W_gate + W_up                  149,520,384 B
    write g [19968], u [19968]            159,744 B   ← intermediates born
    read  g + u                           159,744 B   ← …and read back
    write ffn_mid                          79,872 B
                                          ──────────
    extra activation traffic vs fused:    ≈ 345,856 B ≈ 338 KiB/layer

Hold the two ledgers side by side and the same lesson falls out that the ancestor’s FFN chapter [ferrite-book Ch 17] drew, re-derived here for this geometry. Reading x once is mostly a cache effect — x is 26 KiB and would have been sitting in cache for the second matvec anyway. The hard saving is the one that never shows up as a load at all: in the unfused route the two 19,968-wide intermediates are born, written, and read back; in the fused route they simply never exist.

Then scale it up before deciding how much to care. 338 KiB/layer × 52 layers ≈ 18.0 MB/token of avoided activation traffic — about 12 % of one layer’s FFN weight read. Real, then, but second-order; and on the serving route it is deliberately not taken (§18.6). What is neither second-order nor optional is the weight stream underneath it: 149.5 MB per layer, 7.78 GB across 52 layers, irreducible at Q4_K bitrate no matter which route dispatches it. That is the number Ch 1’s bandwidth argument leans on.

18.8 Tradeoffs

Fused 4r2s vs the unfused control — bytes versus bits. The fusion saves ~338 KiB/layer of activation round-trips and one dispatch: three closures become one. The control gives up both of those and buys back llama.cpp’s exact per-node arithmetic, and with it the public logprob contract. Muser ships the control as the default and gates the fusion behind MUSER_FERRITE_FFN_GATE_UP, and the reason lives in the source, where the fused kernels’ rounding divergence “breache[s] public logprob tolerance” against the source-pinned llama Metal graph (decode.rs:5819-5836 and :1334 for the flag, decode.rs:2085-2091 for the reason).

That verdict was not invented here. The dispatch-gap investigation had already made the discipline explicit for a whole family of these fusions: they were “removed, not hidden behind a tolerance” [docs/decode-dispatch-gap-20260815.md, Rejected hybrid postmortem]. A flag is the gentlest form of that same judgement — the code survives, the default does not move. What we still cannot tell you is what the flag would be worth end to end: no retained A/B quotes a tok/s delta for this specific flag [unverified]. The burden the fused route has to clear is full-logit parity first, per the comment, and at the pin that gate is not recorded as passed.

Two accumulators in lockstep vs two kernels. The 4r2s design doubles down on the V4 pattern of Ch 13: the register x-cache is shared across two matrices and two rows simultaneously — one x load feeds gate_sumf[0..1] and up_sumf[0..1]. The cost is register pressure (yl[16], yh[16], sumy, four running accumulators) and a kernel that only exists for Q4_K (the q4k_v4_… helper is Q4_K-specific; Q5_K/Q6_K tensors would need their own decoders — hence the dtype guard at decode.rs:5820-5821). The payoff is the traffic ledger of §18.7.

The design did not arrive in that shape, either. An older 4-SIMD-group variant — ffn_q4k_gate_up_silu_4sg, one row per threadgroup, 128 threads — sits in the same shader file at ffn_fused_tail.metal:295, with a threadgroup-x-cache sibling at :392. Both are the same idea at a coarser grain of x-reuse, and both were kept rather than deleted; the 4r2s port at :496 is the one Muser wired.

SiLU vs GELU vs ReLU. The choice is the checkpoint’s, not the engine’s; the function-constant mechanism (§18.5) exists so one .metal source could serve either at zero runtime cost. Muse Glimmer is SiLU — every route (fused apply_activation, control muser_silu_mul_inplace, CPU oracle silu_fast) implements x·σ(x).

The normed-quant tail variants — present but unwired. The shader library carries a family that goes one step further than 4r2s: fuse the norm into the gate+up read, so that the FFN would consume the raw residual and normalize in-kernel. The kernel is ffn_q4k_gate_up_silu_normed, at shaders/ferrite/ffn_fused_normed_quant.metal:296, with Q5_K siblings at :1 and :190. We went looking for whoever calls it, and found nobody: at the pinned commit no Rust wrapper binds any kernel from that file (verified: no reference in crates/muser-engine/src). It is not on a live path — retained as Ferrite-lineage research material.

The reason that belongs in a tradeoffs section rather than a footnote is that it names the ceiling of the whole fusion strategy. Its fate is the story of §18.6 taken one step further: the more arithmetic you fold across a norm boundary, the harder bit-exactness becomes (Ch 19 §19.9 makes that tradeoff precise).

18.9 Where the gap lives

The gate+up stage is not the gap — but its fusion is one of the casualties of the exactness contract. In the one-token closure accounting, the FFN gate-up and swiglu closures are “common math” — identical counts in the production and legacy graphs (the 406 = 406 row of [docs/decode-dispatch-gap-20260815.md]); the +196-closure gap lives in norm boundaries, SWA staging, KV publication, and one copy, not here. The FFN’s connection to that story is the reverse direction: the fused kernel this chapter teaches is part of the legacy route’s fusion set, and the serving graph pays extra closures (two matvecs + silu_mul instead of one) precisely to keep llama.cpp’s node-for-node arithmetic (decode.rs:2085-2091). When Ch 35 and Ch 40 audit what was measured and rejected, this is a standing example: a structurally sound bandwidth win, held out of serving by a logprob tolerance, exactly as the campaign’s fail-closed culture requires.

The FFN is half closed: ffn_mid [19968] holds the gated, activated intermediate. One projection remains — the squeeze back to 6,656, the residual add, and the fused tail that hands the next layer its normed input. That tail is also where this book’s central tradeoff — dispatch count versus the logprob contract — gets priced to the last ULP. It is Ch 19.

References

• crates/muser-engine/src/shaders/ferrite/ffn_fused_tail.metal:485-561 — ffn_q4k_gate_up_silu_4r2s, the fused Q4_K SwiGLU kernel (primary source; the _4sg ancestor variant at :295, _4sg_tgcache at :392).
• crates/muser-engine/src/shaders/ferrite/ffn_fused.metal:14-30 — FC_FFN_ACTIVATION function constant + apply_activation (PSO-build-time SiLU/GELU selection).
• crates/muser-engine/src/shaders/ferrite/_q4k_helpers.metal:34-88 — q4k_v4_dual_row_mac, the dual-row Q4_K MAC both projections share.
• crates/muser-engine/src/metal/encode/ffn.rs:7-36 — encode_ffn_q4k_gate_up_silu_4r2s (grid (i_dim/4, 1, 1) × 64 threads); :38-60 encode_silu_mul (control-route activation).
• crates/muser-engine/src/shaders/muse_reference.metal:4-13 — muser_silu_mul_inplace, the pointwise control kernel.
• crates/muser-engine/src/decode.rs:5819-5862 — the flag/dtype gate, the fused dispatch, and the unfused control; :1334 the MUSER_FERRITE_FFN_GATE_UP env read; :2085-2091 the serving-exactness comment; :5165-5185 the packed decode group’s unfused FFN.
• crates/muser-engine/src/reference.rs:493-516 — the CPU oracle’s gate/up/silu·mul order; quant/helpers.rs:70-74 silu_fast.
• crates/muser-engine/src/config.rs:180 — intermediate_dim from muse-glimmer.feed_forward_length.
• crates/muser-bench/src/m16.rs:171-175 — ffn_gate/up 6656->19968 q4k (release-artifact dtype evidence).
• crates/muser-engine/src/shaders/ferrite/ffn_fused_normed_quant.metal:1,190,296 — the normed-input fused FFN family; no Rust binder at the pin (§18.8).
• [docs/decode-dispatch-gap-20260815.md] — closure accounting (common math) and the rejected-hybrid postmortem’s remove-don’t-tolerate discipline.
• [docs/extraction-manifest.md] — silu_fast and the shader lineage from Ferrite 83cfd55…/a85048a9….
• Ch 13 — the V4 lane decomposition and the pinned ggml matvec the control route dispatches.
• Ch 6 — the Q4_K super-block and dequant formula behind the MAC helper.
• Ch 1 — the per-token weight-read arithmetic this block dominates.
• [ferrite-book Ch 17] — the ancestor’s SwiGLU chapter; the fused-vs-control honesty pattern and the intermediate-buffer byte ledger ported here.
• [arxiv:1706.03762] — Vaswani et al., Attention Is All You Need (the ungated ReLU FFN).
• [arxiv:1710.05941] — Ramachandran et al., Searching for Activation Functions (Swish/SiLU).
• [arxiv:2002.05202] — Shazeer, GLU Variants Improve Transformer (SwiGLU).