perf(seedless): concurrent MTLComputeCommandEncoder with an explicit per-layer dependency table — blocked on #143

[penta2himajin]

Status

Blocked on #143. Do not start until P3 reports a bubble share above the decision threshold. This issue is the expensive option in the dispatch-count workstream and the only one that can silently break losslessness, so it is gated on measurement by design.

The observation

Every compute encoder in swift/Sources/QwispCore/ is created with makeComputeCommandEncoder() — the Metal default MTLDispatchType.serial. dispatchType and memoryBarrier do not appear anywhere in the module. So all ~654 dispatches/token (#143) are totally ordered, including pairs that are structurally independent.

Independent work that exists in the current op sequence and is currently serialized:

• MoE block: the routed-expert chain (gqmm4_swiglu_rows → gather d), the shared-expert chain (gqmm4_swiglu_rows → shared down qmm), and the shared-gate qmm8 are three chains that only converge at shared_gate_combine_rows_fold.
• GDN: z is produced by the in-proj demux and consumed only at gdn_norm_gate, so it is live across the whole conv/prep/recurrence chain; the recurrence itself (gated_delta_step, 32×Dv×Hv threads) is low-occupancy and has nothing to overlap with under the serial encoder.
• Attn: the v-cache write is independent of q/k prep; only sdpa needs all three.

At M=1 every one of these kernels is a matvec that underutilizes the GPU on its own — which is exactly the regime where overlapping independent kernels, rather than making any single kernel faster, is the available lever.

The change

Switch the decode/verify forward encoders to .concurrent and insert memoryBarrier(scope:) only at true producer→consumer edges.

This is encode-order refactoring, not kernel surgery — no kernel source changes, no arithmetic changes. Prohibition 3 is satisfied in letter. It is nonetheless the highest-risk item in this workstream, for the reason below.

The risk, stated plainly

A missing barrier does not fail loudly. It produces a race whose outcome depends on scheduling, so it can pass RAWTESTS, pass a 512/512 lossless run, and then corrupt one token in production a week later. The usual "tests are green" evidence is not sufficient evidence here.

Consequence for how this must be built:

1. Write the dependency table first, as a reviewable artifact — every buffer in GdnScratch / AttnScratch / MoEScratch / hBuf / cache buffers, with its producer and consumer dispatches per layer type. The table is the deliverable that gets reviewed; the code is downstream of it. Ad-hoc "add barriers until it works" is the failure mode to avoid.
2. Scratch-buffer aliasing is the sharp edge. The scratch structs are shared across all 40 layers precisely because serial execution made that safe (GPUScratch comment: 全層で再利用＝serial 実行ゆえ安全). Under .concurrent, layer i+1's producer can legally start before layer i's consumer has read the same scratch slot. Either a barrier at every layer boundary (which caps the achievable win — worth estimating before implementing, since it may already fall below the perf: Step 0 — decode dispatch inventory + serial-vs-concurrent encoder bubble probe #143 threshold) or per-layer scratch (which costs memory, and on the streaming tier memory pressure is itself a known failure mode, cf. Requesting more testing with long context + swap thrashing regression #86 / PR fix(server): strict streaming right-sizes C to the machine's memory budget #70).
3. Start with the MoE block only, single layer, flag-gated. That is where the three independent chains are, it is the biggest share of the per-layer dispatch count, and it bounds the blast radius.

Gates

• G-A (locked, model-free): concurrent vs serial bit-identical across a synthetic multi-layer stack, M ∈ {1, 2, 8}, both layer types, repeated ≥100× per configuration — a race that shows up once in 50 runs must be caught by the gate, not by a user. Bumps RAWTESTS total = N.
• G-B (lossless, model): spec-vs-greedy 512/512 × 4 regimes, run more than once; flag-off byte-unchanged (QWISP_CONCURRENT_ENCODE=0 default, and it stays default-off until G-D and a full review round are both clean).
• G-C (wiring): reviewer walks the dependency table against the encode order line by line, independently of the author. This is the load-bearing gate for this issue — treat it as such.
• G-D (speed, driver-decided): tok/s and profLastGPUMs before/after on resident and streaming, plus the isolated MoE-block A/B from perf: Step 0 — decode dispatch inventory + serial-vs-concurrent encoder bubble probe #143 P3 as the sanity anchor. If the real-block win is materially below what P2's synthetic bound suggested, report the discrepancy rather than shipping on the synthetic number.

Honest priority note

Resident already decodes at 85–91 tok/s and the estimated roofline is ~250–300 tok/s (#143 U3, and that estimate is itself unconfirmed). If P4 shows the engine is nearer the roofline than the estimate suggests, or P3 shows the bubbles are small, this issue should be closed not-planned rather than attempted — the cheap wins in #144 and #145 do not depend on it, and a race introduced into the bit-proven decode core is a far worse outcome than leaving 10% on the table.

Refs

#143 (Step 0 — hard prerequisite) · #144, #145 (independent, cheaper dispatch-count reductions) · notes/01 §2-4 (the sync-floor framing, from the superseded MLX era) · SeedlessFusedForward.GPUScratch / MoEScratch / GdnScratch / AttnScratch (the aliasing surface) · #86, PR #70 (streaming-tier memory pressure, relevant to per-layer scratch)

[penta2himajin]

Closing not-planned — prior NO-GO, and it fails #143's own decision rule before it starts (Claude + Codex cross-check, 2026-07-25)

The issue is well-constructed — it gates itself on #143, names its own risk honestly, and includes a self-close condition. Applying that condition is what closes it.

1. Prior art: explicit NO-GO

notes/10-moe-combine-fold-spec.md §0 (2026-07-05 recon, marked 正典):

(a) gather_d+combine / (c) route qmm8+top8 / **§3 concurrent** = **NO-GO**(grid 8×/32× collapse or 帯域律速)

Concurrent scheduling was already rejected. notes/09:14 separately estimated the routed∥shared overlap at only 0.5–1 ms. Neither was cited here.

One fair objection first: the stated reason "帯域律速" is loose. The engine achieves 141.9–151.9 GB/s = 35.5–38.0% of the M1 Max's 400 GB/s (corrected numbers now in #143), so it is not bandwidth-saturated. But the conclusion survives on better grounds — notes/09 §1 decomposes the gap as weight floor ~3.6 ms · dispatch intervals ~1.5 ms · inefficient M=1 GEMV ~8 ms. The engine is occupancy-limited, not bandwidth-limited and not dispatch-limited. Overlapping kernels that are each individually occupancy-starved, and that all contend for the same memory system, does not convert that ~8 ms into throughput.

2. Confirmed observations (the issue is factually right about the mechanism)

• dispatchType and memoryBarrier appear nowhere in swift/Sources/QwispCore/; all 90 encoders are default makeComputeCommandEncoder() = .serial.
• Scratch is allocated once per forward and shared across all 40 layers (SeedlessFusedVerify.swift:3066, :3296). The aliasing hazard the issue flags is real.

3. But the win is capped below the issue's own threshold

Because scratch is shared and layer i's residual writes hBuf which layer i+1 immediately reads for its input norm (SeedlessFusedVerify.swift:3358 → :3403), a barrier is mandatory at every layer boundary. Under that constraint a generous dispatch-DAG upper bound — assuming equal-duration kernels, unlimited resources, and zero bandwidth contention — is:

configuration	serial slots	critical slots	max theoretical win
QWISP_FUSE_MOE2=1	654	~554	≤15.3%
stock HEAD defaults	694	~594	≤14.4%

#143's binding decision rule is "P3 bubble share > ~20% ⇒ priority lever; < ~10% ⇒ drop it". The unattainable upper bound is already below the 20% bar, and the realistic figure is lower still (the independent branches are simultaneously streaming large weights). The gate cannot be passed, so running P3 to decide this issue is not necessary.

4. One premise is factually wrong

The issue claims three independent chains in the MoE block (routed / shared / shared-gate qmm8). There are two: the shared-gate qmm8 is already folded into the S2 convergence kernel (SeedlessFusedVerify.swift:1820) and is not a separately schedulable dispatch. The GDN z example is also not an independent chain — z is a value waiting on coreOut, with no concurrent dispatch to overlap. The attention q-prep/k-prep/v-write fork is real, but it is three small kernels out of 694.

5. Risk side is unchanged

As the issue itself states: a missing barrier fails silently and schedule-dependently — it can pass RAWTESTS, pass 512/512 lossless, and corrupt one token in production later. Accepting that exposure on the bit-proven decode core, for a bounded ≤15% that the measured decomposition says is mostly unreachable, is the wrong trade.

Verdict

Closing not-planned, per this issue's own "Honest priority note". The cheap win (#144) is open and does not depend on this. If the bottleneck is ever re-measured as dispatch-bound rather than GEMV-occupancy-bound — or hardware gains a device-scope barrier — reopen with that evidence.

Refs: notes/10:14 · notes/09 §1, :14 · SeedlessFusedVerify.swift:1820, :3066, :3296, :3358, :3403 · #143 (corrected baseline + roofline)