dasLLAMA gpt-oss arc: idot dot migrations, Metal fusion lane, fused MoE expert chain, fleet boards - #3579
Merged
Merged
Conversation
…the gate The babysit loop was re-running full preflight (~20 min) before every push, even for a 6-line review-fix commit. CI runs the whole matrix on every push for free, so per round only the fast local gates that CI is unforgiving about (compile/format/lint on the changed files, + docs/JIT smoke only if that class was touched) are worth running — everything else is left to CI. Full preflight stays only for a genuinely large change. Documents `git push --no-verify` as the sanctioned bypass of the per-round preflight token (deliberate + CI-covered, not the silent skip the make_pr rule bans), and how to tell an infra-canceled runner red from a real one. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…mm seed, act ghost-tail cut, prefix folds into bucket Three structural cuts on the gptoss pp512 critical path (Wave 1 fusion lane): - MetalMoeMulMmMx4 takes wb/hasb and seeds its accumulator tiles with stride-0 simdgroup_loads of the expert's bias row — the three per-layer MetalMoeBiasRows dispatches (the 196K-tg early-exit launch storm, 15.1ms of window) fold away; the moe_bias chase arm now knocks the fold out (hasb=0) for attribution. MetalMoeBiasRows deleted. - swiglu_oai prefill twin (MetalSwigluOaiPf) exits ghost rows past the last bucket's padded end on a uniform-broadcast liveness bound (~24% of mpad at 512; -1.6ms). The 3-binding MetalSwigluOai stays as-is — the decode rail's enc_ew2 dispatches it with slots 0-2 (the fam-gptoss forced-step cell caught the shared source the first time around). - MetalMoePrefix (1-tg serial walk) folds into MetalMoeBucket: each bucket tg computes its expert's padded-prefix base in-tg and publishes basep[e] — one dispatch and one serial barrier level per layer gone, bit-identical CSR. gptoss metal pp512: GPU window 512.2 -> 495.2ms, 998 -> 1030.9 t/s best-window; adjacent lcpp_bench pairs read 1032.6 vs 1033.8 (0.999, statistical tie, was 0.966) with tg128 1.045-1.064 held. Prefill dispatches 650 -> 626. 30B near-tie re-paired clean (pp 1.018, tg 1.138). Gates: kernels suite, prefill base/kq/qkv, all five MoE fam parity arms (PARITY_FULL) green. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…512 tie das metal 1031.5 +/- 2.1 vs lcpp stock 1033.8 +/- 2.4 => pp512 0.998 (board 1.00), tg128 88.1 vs 84.4 => 1.044 held; official rail, adjacent warm pair, tripwire silent (cv 0.2%). Also re-merges site/files/dasllama/bench_records.json, which was stale since before #3564 (still carried the 945.4 / 0.92 card). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Q8bZnWkVWeXy8jfrxSojso
Anton's M4 Pro mini (10P+4E, 64GB, macOS 15.2), full pub_catalog sweep with adjacent llama-bench refs at the pinned sha (ebd048f), fresh full tune (sdot kstep2 family holds; smmla measured live and declined — 147.5 vs 165.9 GMAC/s, still one MMA pipe on M4). Reds, all pp: gpt-oss metal 0.93, E4B metal 0.98, E4B accel 0.99. Standouts: 35B metal tg 1.27, 30B accel pp 1.78, gpt-oss cpu pp 1.76. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…4B/gpt-oss) Irina's M3 Air (4P+4E, 16GB, macOS 26.4.1), native build + native pinned refs. CPU legs all green (gpt-oss accel pp 1.71); metal slightly red on the small GPU (12B 0.97/0.94 — M1-tuned tiles + passive thermal); gpt-oss metal legs honestly absent (12.1GB exceeds the 16GB box's Metal working set). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…in-das idot4 — hand-NEON sdot4x4 leg retired The P1 assumption "backends pattern-match the generic idot4 lowering to native dots" fails on AArch64: the stride-4-shuffle widen-multiply form compiles to a zip/uzp/smull expansion, zero SDOTs (probe-disassembled), while the plain autovectorized loop DOES get sdot from the vectorizer. So the ledgered per-ISA strengthening lands now: - llvm_jit_intrin: idot/idot4 signed x signed emit ONE @llvm.aarch64.neon.sdot (idot adds a v4i32 reduce -> addv) on the host-features rail — g_target_host_features is exactly when +dotprod is force-appended, so the generic/cross rail (kernel-free exes) keeps the portable IR and can never hit "Cannot select". unsigned x signed needs USDOT (i8mm, absent on M1) — rides the i8mm un-gate follow-up. LLVM_JIT_CODEGEN_VERSION 0x48 -> 0x49. - dasllama_math_default: new dot_q8q8_idot4x4/_f16s — the 4-accumulator GEMV dot in plain das (byte16 loads + idot4), bit-identical fold to the hand-NEON sdot4x4 leg (exact int sums, same float order). - dasllama_math_aarch64_neon: the hand-NEON dot_q8q8_sdot4x4/_f16s are DELETED; the arm64-sdot backend's q8q8 kernels dispatch the plain-das dots. NEON intrinsics remain only for the mx4 + laneq tiles. dasllama_math_gen row tails switch with it. Interleaved A/B (M1 Max, 8 P-cores, DRAM-rotated GEMV, warm adjacent cells): single-thread idot4x4 50.9-51.8 GB/s vs hand-NEON 50.2-50.9 vs tuned-autovec vec16 35.9-38.5 (n=2048 d=8192 and n=2816 d=1408); threaded all cells tie at the bandwidth ceiling (~102/112 GB/s) — board-neutral, kernel-parity proven. 1-acc idot rewrites of the template measured NO gain (per-block hsum drains the dot chain) — the measured reason the autovec template stays for x64. Gates: tests/type_lattice/test_lattice_idot interp+jit; dasLLAMA test_matmul, test_matmul_batch, test_kernel_backend, test_groupn, test_kquant, test_fused_decode, test_mxfp4, test_batch_decode, test_forward, test_prefill, test_parity (45/45 with DASLLAMA_CPU_PREFILL=1; the 4 tripwire errors without it reproduce identically on master) — all -jit. Riders: two pre-existing LINT017s in llvm_jit_run.das (uint64(length) -> long_length). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…n das (shuffle+idot4)
Audit: the 18 integer lattice vector types (byte2-16, ubyte2-16, short2-8,
ushort2-8) carried ctors/converts/equality/swizzles and NO shift or bitwise
ops (fp16 family has closed arithmetic; workhorse int/uint vectors have the
full bit surface). This adds exactly `>>` — shift by a scalar int, count
masked to the LANE width, signed lanes arithmetic / unsigned logical — the
nibble-unpack shape (`shuffle(lut, nib >> 4)`). No <<, no >>=, no & | ^
(deliberate minimal scope).
- vectypes.h das_sv_shr template + 18 addExterns (the addEquNeqVal rail:
extern for interp/AOT, JIT overrides natively).
- JIT: is_lattice_op2_native routes int-lattice ">>" to visitExprOp2_NonMonade
(else the builtin-address gate would drop the enclosing function to interp);
the shift preamble masks/truncs/splats the count at lane width; AShr/LShr
arms accept tInt8/tInt16/tUInt8/tUInt16 vector bases.
- shuffle: aarch64 now emits AND + one @llvm.aarch64.neon.tbl1 (baseline NEON,
same gate as the tbl16 intrinsics) instead of 16 dynamic extract/inserts.
LLVM_JIT_CODEGEN_VERSION 0x49 -> 0x4a.
- dot_mx4q8_template rewritten in plain das: ubyte16 nibble load, shuffle LUT
lookups (low via the &15 mask, high via the new per-byte >> 4), idot4
accumulate — replaces tbl16_lo/hi + sdot4_w. Bit-exact vs the scalar oracle;
disassembly shows the IDENTICAL per-block instruction mix (2 tbl + 2 sdot +
1 and + 1 ushr); isolated GEMV at gptoss shape (n=d=2880) 16.1 vs 16.3 GB/s
(-1.5%, fallback-u2 hint tuned for the old body — per-box tune re-crowns).
The [tuned] grid and sidecar keys are unchanged.
- conformance: gen_type_conformance gains has_shr + a "shift right" section
(arithmetic vs logical, count masking, shift-by-zero) on the 18 lattice
types + int2/uint4 controls; 20 files regenerated, 438/438 interp + jit.
- datatypes.rst: the lattice int families are now storage + converts + >>.
Measured reason NOT to migrate dot_q8q8kv: its vec16 template already
compiles to sdot via the LoopVectorizer and hits ~39-40 GB/s single-thread at
KV row lengths (n=64/128/256); the 4-acc idot4x4 shape loses 2x there
(per-call fixed cost at 3.5ns/row rows), and the 1-acc idot form measured
exactly at vec16 parity in the GEMV race - no headroom. KV dots stay on the
tuned templates.
Riders: pre-existing LINT017 sweep (uint64(length) -> long_length) in
llvm_jit.das + llvm_boost.das, and the llvm_boost describe() LINT016
(':=' -> '= clone_string').
Gates: tests/type_lattice 438/438 interp+jit; dasLLAMA test_mxfp4,
test_kernel_backend, test_groupn -jit green.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
… + compound assigns) Completes the >> commit to full addFunctionVecBit parity on the 18 integer lattice vector types: per-lane << (lanes wrap like the C cast), bitwise & | ^, and the five compound assigns (<<= >>= &= |= ^=). Same rail as >>: das_sv_* templates in vectypes.h + addExterns (interp/AOT), native JIT lowering (is_lattice_op2_native now routes the whole 10-op surface; isBinary accepts int-lattice vectors — the ~ Op1 arm stays unreachable for lattice, matching the workhorse families where ~ is also absent). Shift counts mask to the lane width on every form. Conformance: the shift-right section grows into shift-right/shift-left/ bitwise sections (arithmetic vs logical >>, << lane wrap pinned via two's-complement expected values computed in the generator, count masking, all five compound assigns) on the 18 lattice types + int2/uint4 controls — 478/478 interp + jit. datatypes.rst updated. Real-.dlim A/B of the migration commits (branch vs master worktree, same binary, lcpp_bench, tune/warm discard + timed adjacent pairs, M1 Max t8): Llama-3.2-1B Q8 pp512 887/861 vs 874/844 tok/s (+1.4/+2.0% lean), tg128 82.4/82.4 vs 82.4/82.3 (tie); gpt-oss-20b mx4 pp512 ~200 both sides, tg128 within +-1.5% with the sign following RUN ORDER (a reversed pair flipped it) — statistical parity, no regression on either model. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…dot-gptoss # Conflicts: # skills/babysit.md
…dot-gptoss # Conflicts: # site/files/dasllama/bench_records.json
…d metal files Post-merge lint of the gptoss-fusion surface under the post-#3576 rules: 14 int(...) memcpy-size casts dropped (int64 overloads exist) and 3 int64(length(...)) -> long_length in dasllama_metal_prefill.das — all pre-existing on master, swept because the fusion lane owns these files in this PR. bench_metal_moe_lab was already clean. Lint 0/0, format verified, test_kernel_backend -jit green in the integration worktree. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ector) One measurement window of -n single-token forwards at --depth, engine prof_add buckets reset after prefill+warmup, pace directly comparable to a lcpp_bench tg cell. --accel mirrors lcpp_bench's flag for profiling the accel flavor's decode. The previous incarnation was a session tool that got cleaned up — only its tune sidecar survived in the main tree; this one is a tracked benchmark so the detector stops evaporating between decode chases. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…uant -> down+reduce The routed-expert groupn sequence ran 3 fork/joins per layer with act/requant/reduce single-lane between them; the chain runs both stages inside ONE team_parallel_stages rendezvous — stage 0 = (expert, 32-row span) gate+up rows + bias + act + requant, stage 1 = dim spans computing every expert's down rows + bias + the weighted sum in slot order. Bit-identical to the groupn path (same rows cores, same dot->bias->act element order, same 32-block requant) — proven by the new test_fused_decode_moe_chain gpt-oss arm (ids + logits bit-exact, chain vs per-op). New backend slot: mx4_rows (MatmulMx4RowsFn) — arm64/x64-gen register mx4q8_gemv_gen directly (the [tune] GEMV is the rows core, the mm_rows precedent), arm64-sdot gets a row-major dot_mx4q8 walk. The chain gates on kernel_backend_has_mx4_rows() + team mode + 32-divisible spans, so every other backend keeps the per-op dispatch path. decode_prof gains --no-fused (the chain-vs-per-op A/B lever). Rider: LINT017/018 sweep in dasllama_common.das (43 long_length + 4 memcpy casts, the touched-file sweep discipline). Measured M1 note: gpt-oss cpu decode wall is FLAT under the fusion (blk_ffn 12493 vs 12468 us/token, 3 joins/layer -> 1) — a real datum against the join-cost theory of the remaining lcpp gap; the chain still wins the structure (fewer syncs, threaded act/requant) and is the substrate for further expert-chain work. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The --no-fused A/B is only trustworthy when the log carries proof of which path ran (the matrix that motivated this silently benched chain-vs-chain through a zsh quoting bug). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…d-tool parity) Lifted from history/dasLLAMA/benchmarks/decode_prof.das (the archived predecessor — which this file should have credited instead of assuming deletion). First use settled the lane question: at -t N all N lanes (N-1 workers + the computing caller) run 95-97% busy in the decode window; t7 and t8 do equal wall work, so the t7==t8 pace plateau is shared-resource saturation, not a parked or half-productive lane. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Contributor
There was a problem hiding this comment.
Pull request overview
This PR bundles several dasLLAMA “gpt-oss arc” lanes into one integrated change set: it extends the integer lattice vector bit-operation surface end-to-end (runtime/AOT + LLVM JIT + conformance tests/docs), migrates key dasLLAMA dot paths onto the idot4/shuffle builtins (retiring hand-NEON dot bodies), and introduces a fused routed-MoE expert decode chain plus a permanent decode profiler and new fleet performance records.
Changes:
- Add full bitwise/shift operator support for 8/16-bit integer lattice vectors, including LLVM JIT lowering and expanded type-lattice tests/docs.
- Migrate dasLLAMA Q8·Q8 and MXFP4·Q8 dot implementations to
idot4/shuffle-based paths and update arm64 NEON backend registration accordingly. - Add fused MoE expert-chain decode path gating on a new backend
mx4_rowscore, plus a newdecode_profbenchmark and updated performance records.
Reviewed changes
Copilot reviewed 41 out of 42 changed files in this pull request and generated no comments.
Show a summary per file
| File | Description |
|---|---|
| utils/dasgen/gen_type_conformance.das | Extends generated conformance tests with an optional bit-ops section for vector types. |
| tests/type_lattice/test_int2.das | Adds shift/bitwise operator conformance checks for int2. |
| tests/type_lattice/test_uint4.das | Adds shift/bitwise operator conformance checks for uint4. |
| tests/type_lattice/test_short2.das | Adds bit-ops conformance coverage for short2. |
| tests/type_lattice/test_short3.das | Adds bit-ops conformance coverage for short3. |
| tests/type_lattice/test_short4.das | Adds bit-ops conformance coverage for short4. |
| tests/type_lattice/test_short8.das | Adds bit-ops conformance coverage for short8. |
| tests/type_lattice/test_ushort2.das | Adds bit-ops conformance coverage for ushort2. |
| tests/type_lattice/test_ushort3.das | Adds bit-ops conformance coverage for ushort3. |
| tests/type_lattice/test_ushort4.das | Adds bit-ops conformance coverage for ushort4. |
| tests/type_lattice/test_ushort8.das | Adds bit-ops conformance coverage for ushort8. |
| tests/type_lattice/test_byte2.das | Adds bit-ops conformance coverage for byte2. |
| tests/type_lattice/test_byte3.das | Adds bit-ops conformance coverage for byte3. |
| tests/type_lattice/test_byte4.das | Adds bit-ops conformance coverage for byte4. |
| tests/type_lattice/test_byte8.das | Adds bit-ops conformance coverage for byte8. |
| tests/type_lattice/test_byte16.das | Adds bit-ops conformance coverage for byte16. |
| tests/type_lattice/test_ubyte2.das | Adds bit-ops conformance coverage for ubyte2. |
| tests/type_lattice/test_ubyte3.das | Adds bit-ops conformance coverage for ubyte3. |
| tests/type_lattice/test_ubyte4.das | Adds bit-ops conformance coverage for ubyte4. |
| tests/type_lattice/test_ubyte8.das | Adds bit-ops conformance coverage for ubyte8. |
| tests/type_lattice/test_ubyte16.das | Adds bit-ops conformance coverage for ubyte16. |
| include/daScript/misc/vectypes.h | Implements small-vector shift/bitwise helpers used by interpreter/AOT and extern bindings. |
| src/builtin/module_builtin_vector_ctor.cpp | Registers the new small-vector bitwise/shift operators as builtins/externs. |
| modules/dasLLVM/daslib/llvm_boost.das | Fixes several LLVM C-API wrapper calls to use long_length, and adjusts basic-block name handling. |
| modules/dasLLVM/daslib/llvm_jit.das | Updates debug-info string-length plumbing and adds JIT lowering support for lattice int-vector shifts/bitwise ops. |
| modules/dasLLVM/daslib/llvm_jit_run.das | Bumps LLVM_JIT_CODEGEN_VERSION and updates LLVM module-flag string length usage. |
| modules/dasLLVM/daslib/llvm_jit_intrin.das | Adds aarch64-native sdot lowering for idot/idot4 (where gated) and aarch64-native tbl1 lowering for shuffle. |
| modules/dasLLAMA/dasllama/dasllama_math.das | Adds mx4_rows backend slot and public accessors/gates for fused MoE expert-chain use. |
| modules/dasLLAMA/dasllama/dasllama_math_default.das | Introduces dot_q8q8_idot4x4 (+ f16-scale twin) plain-das dot implementations. |
| modules/dasLLAMA/dasllama/dasllama_math_gen.das | Switches generated kernels’ row-major tails to the new idot4x4 dot implementation and wires mx4_rows. |
| modules/dasLLAMA/dasllama/dasllama_math_aarch64_neon.das | Migrates arm64 backend dots to idot4x4/shuffle-based forms and wires mx4_rows registration. |
| modules/dasLLAMA/dasllama/dasllama_common.das | Adds fused routed-MoE expert-chain decode path gated by mx4_rows, and performs long_length/memcpy-size sweeps in touched areas. |
| modules/dasLLAMA/tests/test_fused_decode.das | Adds a parity test asserting fused MoE expert-chain decode is bit-exact vs per-op sequence (fixture-gated). |
| modules/dasLLAMA/benchmarks/matmul/bench_metal_moe_lab.das | Adds an additional Metal lab binding/control buffer for a bias-disabled prod arm. |
| modules/dasLLAMA/benchmarks/decode_prof.das | Adds a permanent decode profiling benchmark that reports per-bucket prof_add totals over a decode window. |
| modules/dasLLAMA/performance/records/m3air.json | Adds new M3 Air fleet performance record set. |
| modules/dasLLAMA/performance/records/m1.json | Updates select M1 performance record entries (date/sha/results). |
| doc/source/reference/language/datatypes.rst | Documents the new integer lattice vector bit-operation surface and shift semantics. |
| skills/babysit.md | Adds CI triage guidance to distinguish infra cancellations from code regressions. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
borisbat
added a commit
that referenced
this pull request
Jul 27, 2026
Two conflicts in dasllama_common.das, both resolved to this branch:
kv_ensure — the growth loop is split into a [cold_path]
kv_grow_block_table here; master still had it inline.
Both sides had already moved to long_length, so the 64-bit
sweep merged silently — we ran the same fix independently.
scratch_resize — the @scratch parameter. #3579 arrived with a new comment on
the same function reading "grow-only no-init: contents are
stage output (fully overwritten)", which is independently
the exact property @scratch declares.
Their new moe_experts_chain needed one annotation: the fourth
team_parallel_stages lambda capture, same dispatch-API cost as the three
already marked. Everything else in it passes untouched, because its only
sizing op is s.chain_stages, already @scratch on this branch.
Verified on the merged tree with a daslang rebuilt for #3579's lattice
operators (ubyte16 >> int needs its C++): engine compiles clean, 65 changed
files lint clean, ENVIRONMENT.md shows no drift, PERF fixtures hold, and
test_fused_decode is 13/13 under -jit with the gpt-oss arm skipped for want of
the gguf.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
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
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.
The gpt-oss arc bundle: the idot dot-kernel migrations, the gpt-oss Metal fusion lane, the fused MoE expert chain, the M4/M3-Air fleet boards, and the permanent decode profiler. Four previously-unmerged lanes plus this session's work, integrated and measured together.
What's in
idot dot migrations + lattice bit surface (
bbatkin/dasllama-idot-migrationlane)dot_q8q8anddot_mx4q8GEMV families rewritten in plain das over theidot4/shufflebuiltins; the hand-NEON dot bodies are deleted. Native JIT arms: ss-formidot4/idotemit onellvm.aarch64.neon.sdot(gated exactly where+dotprodis force-appended; the generic/cross rail keeps portable IR),shufflelowers to a singletbl1.LLVM_JIT_CODEGEN_VERSION0x48 → 0x4a.<<>>&|^plus the five compound assigns on all 18 int lattice vector types, extern + native-JIT rail, conformance 478/478 interp+JIT (shift-left wrap pinned; without a native arm a lattice builtin op drops the enclosing function to interp — every op got one).dot_q8q8iso ST 50.9-51.8 GB/s vs hand-NEON 50.2-50.9. KV dots stay on tuned templates (measured: vec16 already sdot-optimal via the LoopVectorizer).gpt-oss Metal fusion lane (
bbatkin/dasllama-gptoss-fusionlane)Fused MoE expert chain (new this session)
team_parallel_stagesrendezvous: stage 0 = (expert, 32-row span) gate+up rows + bias + act + requant; stage 1 = dim spans computing every expert's down rows + bias + the weighted sum in slot order.mx4_rows(the mx4 twin ofmm_rows): arm64/x64-gen registermx4q8_gemv_gendirectly, arm64-sdot gets a row-majordot_mx4q8walk. The chain gates on team mode + the slot + 32-divisible spans; every other backend keeps the per-op path.test_fused_decode_moe_chainarm (gpt-oss,DASLLAMA_PARITY_FULL-gated) asserts ids + logits identical chain-vs-per-op; the full fused-decode suite is 14/14.decode_prof — the position-0 detector, now permanent (
modules/dasLLAMA/benchmarks/decode_prof.das)prof_addbuckets;--no-fused(chain A/B),--accel,--trace(jobque lane trace of the window). Successor of the archivedhistory/dasLLAMA/benchmarks/decode_prof.das.Fleet boards (
bbatkin/dasllama-m4-boardlane)records/m4.json: first M4 Pro board, 45/48 cells green.records/m3air.json: M3 Air subset. Plus the babysit-skill canceled-runner triage note, and LINT017/018 sweeps in every touched file.Current gpt-oss measurements on this branch (adjacent pairs, alternating order)
The remaining cpu tg gap is characterized (per-lane kernel parity with llama.cpp; an M1-specific memory-saturation wall from 7 lanes on, trace-proven with all lanes 95%+ busy) — the one open CPU lever (single-stream fused-scale mx4 plane) and the M4 post-CLT re-baseline are follow-ups.
Gates
test_aot_subsetbuild. tests-cpp / tests-aot / sequence are deferred to CI by agreement; pushed with--no-verifyfor the same reason.test_fused_decode.das14/14 with the new gpt-oss MoE-chain arm underDASLLAMA_PARITY_FULL=1(ids + logits bit-exact, chain vs per-op)🤖 Generated with Claude Code