Source layout plan — the runtime, launch and kernel layers

Status: every step has landed (2026-10-05). Steps −1, 0, 1, 2, 3, 4 and 6 are on master with the merge SHAs in the table below; Step 5's Linux half (the MINFER_OP_TIMING module-load re-measure and the stale-prose sweep) landed with this document's own row, and its Mac half (#53's Metal DeviceMemory answer, the second implementation the common decision waits on) landed on macbook (macOS 27.0.1, Apple M4 Pro) — the decision itself is recorded in §1.3 rule 3. Step 4 (Metal, #265) landed from a Mac in three increments. This is the plan of record for splitting the four long backend files and for the naming convention the crate follows afterwards. Tickets: #261 (umbrella) + #262 cuda.rs · #263 cuda/kernels/ · #264 CPU · #265 Metal (Mac) · #266 anchor checker · #267 test files. Step −1 landed #138 and #225 first, as decided.

Status

Each step appends its dated record here when it lands (gates run, counts, box label).

1. Device is the first axis, the layer is the second axis inside each device. The crate keeps its per-device modules (cuda.rs, metal.rs, the CPU trio kernel.rs/quants.rs/vec_ops.rs) and each of them is split into its own inner axis. There is no top-level L1/, L2/, L3/ directory tree: the layers only become directories inside a device.

2. The cross-device interface stays flat and singular. src/graph/backend.rs (Backend, KvProvider) plus src/graph/registry.rs remain the one device seam — exactly as llama.cpp keeps ggml-backend.cpp + ggml-backend-impl.h + ggml-backend-reg.cpp flat beside the per-device directories. No new trait is introduced by this plan.

3. A common is only allowed to exist when it has a second real implementation. Interface eligibility = at least two real implementations and at least two callers using it with the same semantics. The one candidate is allocplan::DeviceMemory: CUDA answered it first and, since #53, Metal answers it too. This rule goes into docs/ARCHITECTURE.md.

   Step 5 pre-analysis (2026-10-05, on 4900298) — no common module yet, and the Linux half owes none. The candidate's three halves: the type (allocplan::DeviceMemory, src/graph/allocplan.rs:67) and the pure policy (budget_decision/weight_budget, allocplan.rs:124 + graph/offload.rs:310, with callers in graph/alloc.rs:889/908 and both loaders) are already device-agnostic and CI-tested; the device answer has exactly one implementation (CudaState::device_memory(), src/cuda/methods/accounting.rs:31, reached through the CUDA-only resolver models::device_memory() at src/models/mod.rs:74, 3 call sites), and no Backend trait hook for it exists. #53 adds the second answer behind that existing resolver, so nothing has to move to make it possible; whether the device-answer code then belongs in a common module stays open until it exists, and "no new module" is a live answer (the type and the policy are already in allocplan). Creating the abstraction now would be the single-real-implementation case this rule forbids, so it is deliberately not created here.

   Post-analysis (2026-10-05, on the #53 branch) — the second implementation landed, and the answer is still "no new module". #53 added MpsState::device_memory() (Metal's recommendedMaxWorkingSetSize, src/metal/runtime.rs), so allocplan::DeviceMemory now has the two real implementations the rule asks for. The eligible interface is not the enum alone, though: it is the whole device-answer path, and that path has three parts whose homes are already right. The type and the pure policy (budget_decision, weight_budget) stay in graph/allocplan.rs / graph/offload.rs — device-agnostic and CI-tested. Each device answer stays in its own device module, next to the state it queries (CudaState::device_memory(), src/cuda/methods/accounting.rs; MpsState::device_memory(), src/metal/runtime.rs). The routing is the four-line models::device_memory() (Metal > CUDA > CPU), and it is the two-caller seam the rule is about: graph/alloc.rs::memory_budget (E4's feasibility gate) and both loaders' auto fit (models/qwen2/loader.rs, models/qwen3/loader.rs) read it with the same semantics. A common module would therefore have to hold either the two device methods — moving them out of the modules that own the device state, for no caller's benefit — or the routing, a four-line match: a module with one function is the abstraction-for-one-caller case this rule exists to refuse. Counts at the decision: 2 implementations (CUDA, Metal) and 3 same-semantics callers (the E4 gate + the two loaders) of the resolver; 0 new modules.

4. Each backend picks its own inner axis (this is what llama.cpp actually does — it is not uniform): CUDA = kernel family, Metal = layer (ggml-metal-device.* → ggml-metal-ops.cpp → kernels/), CPU = ISA (ggml-cpu/arch/{x86,arm,…}).

5. Existing file paths are preserved. src/cuda.rs stays the module file and gains src/cuda/<part>.rs children (the layout already used by src/cuda/tests.rs); the same for metal.rs, quants.rs, vec_ops.rs, kernel.rs. This keeps all crate::… paths, the check_dead_code_annotations.py grandfather keys, the dead-code baseline file = fields and the documentation's file references valid.

6. Policy is expressed as pure predicates next to the family they gate (llama.cpp's ggml_cuda_should_use_mmq / _mmvq / _mmf convention), not as a separate policy module and not by relocating the decision point. The predicates become pure functions with unit tests that need no device.

7. src/cuda.rs keeps its name and its type; the family impl blocks become descendants of one intermediate parent (src/cuda/methods.rs — not impl.rs: impl is a Rust keyword, so mod impl; is a syntax error — expected identifier, found keyword 'impl', verified with a two-file rustc probe on 2026-10-04), so privacy does the work: private fields of CudaState (defined in cuda) and private helper methods (defined in cuda::methods) are visible in every family file. The split is therefore a pure move — 0 field-visibility edits, 0 pub(super) — with exactly one mechanical edit: the 86 extern "C" launch declarations get pub(crate) so the two launch test files keep resolving them through use super::*;.

Confirmed execution decisions (2026-10-04): land #138 and #225 first (Step −1); this document lands as a docs-only PR (Step 0); the seven tickets in §7.4 are filed now.

Non-goals: no behaviour change, no renaming of public items, no new abstraction layer, no Metal work on a non-Mac box.

2. The four layers today

L2 is not a layer that can be moved away from L1: the CUDA launchers are inherent methods of CudaState, and the Metal ones are methods of MpsCommandBuffer. Splitting L1/L2 apart by directory would be a type refactor, not a file move — that is why the layer axis stays inside each device.

3. What each backend's second axis is

The one rule both device backends share: kernel sources live in <backend>/kernels/. llama.cpp is not uniform here (its CUDA keeps *.cu/*.cuh flat in ggml-cuda/, with only vendors/ and template-instances/ as subdirectories, while its Metal puts shaders in ggml-metal/kernels/); this plan chooses the consistent form the request asked for, and states the rule once.

Two honest wrinkles:

• A CUDA .cu holds the kernels and their host-side launchers (the C++ half of L2), because a launcher must live in the TU that instantiates its kernel (§5). src/cuda/kernels/ therefore means "the CUDA translation units", not "device code only"; the Rust half of L2 is src/cuda/methods/.
• CPU is the exception: quants.rs and vec_ops.rs are not device-private layers — they are the crate's numeric kernel library (graph/kvformat.rs uses quants::quantize_row_q8_0_into, graph/cuda_backend.rs uses vec_ops::RopeStyle), so they stay where they are and are split by ISA, not into a kernels/ directory.

4. Steps

Step −1 — land the two colliding tickets first (decided 2026-10-04)

• #138 (F5 late cross-backend wait). It edits copy_to_host and consumes stream_wait_event / cudaStreamWaitEvent — both in cuda.rs family G, and both are the only two grandfathered bare allow(dead_code) sites. Landing it first removes code the split would otherwise move and deletes two GRANDFATHERED_BARE keys plus one docs/dead-code-baseline.toml entry.
• #225 (pre-warm cost table). It corrects the same measurement the .cu split will change (the fatbin's per-module one-time load), so the corrected record must exist before Step 2 re-measures it.
• Acceptance: each ticket's own gates; master hard-synced afterwards.

Step 0 — documentation and hygiene (docs-only PR)

• Land this document as docs/SOURCE-LAYOUT-PLAN.md; add it to the AGENTS.md docs index and update the AGENTS.md Layout block.
• Fold the stale facts found while measuring into #219 (it already owns two of them): AGENTS.md:3 ~4400 LOC (production code is 55,529 lines), inference_e2e_walkthrough/15-cuda-backend.md:4/29/36 line counts, CUDA-BACKEND-DESIGN.md §"the gates" — 120 sites / 120 / 120 (the count at that revision; 130 on 6b6d94f). The banner naming the deleted CudaCommandBuffer was rewritten to src/cuda/methods/dispatch.rs:124 by Step 1 (#262), the step that moved it. If #219 is not widened, the rest become ticket 7 in §7.4.
• Add the interface-eligibility rule (§1.3) and the layer definition (§2) to docs/ARCHITECTURE.md.
• Acceptance: check-docs green (check_docs_links.py, check_status.py --check, book build). No counter in docs/status.toml changes (the suite counts do not move in this step).

Step 1 — src/cuda.rs → src/cuda/*.rs (landed)

• src/cuda.rs (6 594 → 1 014 lines) keeps: the module doc, AttnWindow, pub struct CudaState (all 29 fields stay private), the free items (cuda_error_name, cstr_owned, CudaPtr, CudaDevicePropBuf, the memcpy/attribute consts, StreamScratch/StreamBinding/bind_stream, ModelLoadGuard, PinnedPool/PinnedBuf, CaptureStaging, MmqCache, layout_of/format_of, concat_rows, the KV-layout constants, gemm_prewarm_disabled, …), the ten #[cfg(test)] mod declarations, and the mod/use/#[cfg(test)] pub(crate) use lines that re-export the two moved FFI surfaces.
• src/cuda/methods.rs (323 lines) — the 15 non-pub helpers whose callers land in a second family file, computed as the transitive closure of the cross-file call graph: context_stream, get_or_grow (called from seven families), mmq_quantize_transposed, mmq_quantize_native, decode_quantize_native, record_mmq_cache_native, prefill_gemm_f16, the four no_*_mmvq predicates, no_q80_p32, fused_b_on, no_w16cache, no_prefill_gemm, no_fa_prefill; plus the 18 mod <family>; declarations and the #[cfg(test)] pub(crate) use <family>::*; re-exports the two launch test files resolve through use super::*. plane_budget_ok stays in methods/weights.rs (every caller is there), so the split has 0 field-visibility edits and 0 pub(super).
• 18 family files under src/cuda/methods/ (28–728 lines each): each holds its impl CudaState block and its own pub(crate) extern "C" launch declarations (the 86 declarations move with their family; no symbol is used by two families). The declarations are re-exported two levels (methods.rs → cuda.rs) under #[cfg(test)], because a single-level glob does not reach cuda::tests and an ungated one is an unused_import under deny(warnings); the families whose declarations a methods.rs helper calls are re-exported ungated.
• src/cuda/ffi_runtime.rs (145 lines) — the cudart/driver FFI block (its declarations become pub(crate), the second mechanical widening) and the test-only extern block.
• src/cuda/methods/policy.rs (113 lines) holds the MMQ gate family J (mmq_gate_on, mmq_enabled, mmq_active, cc, mmq_a_fuse_mode). The no_*/fused_b_on predicates are cross-family (dispatch, prefill_f16, attention), so they live in methods.rs with the other shared helpers: parking them in a sibling policy.rs is exactly what would have forced the pub(super) edits this step avoids.
• Acceptance: CUDA unit 567 / 0 / 42; CPU 481 / 0 / 36 on dgxspark (aarch64, GB10 sm_121) (479 on the CI runner) + integration 10 / 0 / 6 — the rows #138 moved when it landed (565 → 567, 480 → 481 / 478 → 479); cargo fmt --all --check; check_source_layout.py; check_dead_code_annotations.py (its two src/cuda.rs: keys re-pointed to src/cuda/ffi_runtime.rs:cudaStreamWaitEvent and src/cuda/methods/events.rs:stream_wait_event); check_dead_code_oracle.py --config {cpu,cuda} with only the two moved file = lines in docs/dead-code-baseline.toml; real-model gates FEATURES=cuda scripts/real_model_gates.sh 42 / 0 ×2 with bitwise-identical greedy output; check_doc_line_anchors.py green (the split's per-line map is /home/yusiwen/minfer-split/step1/line-map.tsv, and the anchors whose old line was only a locator were re-anchored to the symbol).

Step 2 — src/cuda_kernels.cu → src/cuda/kernels/ (landed: 1 header + 17 TUs)

Route (a): launchers move with the kernels they launch (llama.cpp's CUDA shape). Two of the 77 launchers are the exception (blueprint: /home/yusiwen/minfer-split/step2/README.md §2) — they launch kernels that land in two different target files, so "the launcher moves to its kernel's file" needs the qualification: launch_gqa_attn_split_f16kv launches both gqa_attn_split_partial (decode) and gqa_attn_split_partial_hybrid (hybrid), resolved by merging attention_hybrid.cu into attention_decode.cu; launch_mmq_raw_nb_bt_nt launches mmq_raw_nb_bt_kernel (nb) and mmq_ksplit_reduce_kernel (bt_q6k), resolved by moving mmq_ksplit_reduce_kernel next to the NB kernels. No -rdc=true, no new nvcc flag (route (b), -static-global-template-stub=false, is the recorded fallback; see §5). Prerequisite: ticket 6 in §7.4 (the documentation-anchor checker) lands first or in parallel, because this step moves 163 line anchors in 23 documents.

Placement (decided 2026-10-04): src/cuda/kernels/, so that both device backends obey one rule — kernel sources live in <backend>/kernels/ (Metal gets src/metal/kernels/). The path src/cuda_kernels.cu disappears, so the 342 documentation mentions of it are swept in this step (they are being swept for anchors anyway).

Granularity: every file at or below ~800 lines, no kernel split across files. The previous draft stopped at "ten family TUs", which left attention (~1,600 lines) and mmq_prefill (~2,270) too large. The section inventory measured on 6b6d94f regroups into:

(the extern "C" launcher block 4052–5135, 1,086 lines, contributes ~50–150 lines to each file above; that is why matmul_f32act and attention_* look slightly over their section size.)

• minfer_prewarm_kernels (9188–9237) is decomposed into one extern "C" minfer_prewarm_<family>_kernels() per file plus a dispatcher that keeps the symbol name the Rust side declares at src/cuda/methods/prefill_mmq.rs:48.
• Grouping into PRs (revised): 5–6 groups, not one family each — the file count grew from 10 to 20, so the natural batches are: (1) common.cuh + guard.cu (the infrastructure), (2) attention (3 files), (3) the MMQ prefill family (4 files), (4) the MMVQ decode family (4 files), (5) ops/KV/gemm (5 files), (6) matmul_f32act + leftovers. Each group is independently verifiable on the device.
• Tooling in the same PRs:
  • check_cuda_launch_returns.py discovers src/cuda/kernels/*.cu as a list and audits one file per audit() call (_RESOLVE_LINES is a module global); it must also assert that no <<<>>> lives in a .cuh (the invariant that keeps the audit complete).
  • tests/fixtures/cuda_launch_sites.tsv keeps its four columns and is regenerated in the build list's order. The identity key is the ordered (owner, site, kernel-fragment) list, and only 2 of the 130 triples repeat, so four columns still identify every site; the --fixture writer emits no file column and src/cuda/issue162_tests.rs's assert_eq!(f.len(), 4) stays as it is. The one tooling change is --check-fixture's filter, len(w) == 4 → len(w) >= 4, so a future column cannot silently empty the fixture (a 5-column fixture under the old filter audited 0 of 130 sites — the pre-verified mutation in the blueprint's §5).
  • build.rs compiles the list, emits one .o per file, keeps libcuda_kernels.a, and adds one rerun-if-changed per .cu and per .cuh (a header edit that does not trigger a rebuild is the silent-stale hazard of this step). It also gains a new-file guard: every .cu/.cuh found in src/cuda/kernels/ must appear in the explicit list and every listed file must exist, so a new kernel file cannot silently not compile.
• Acceptance per group: 130-site audit + --check-fixture; MINFER_TEST_ISSUE162=1 device gate; CUDA unit 565 / 0 / 42; real-model gates 42 / 0 ×2 with bitwise-identical greedy output; cold-start timing recorded (the fatbin module count changes — see §7 and #225).

Stage record (2026-10-04, dgxspark (aarch64, GB10 sm_121)). One PR per group. Every stage regenerates tests/fixtures/cuda_launch_sites.tsv in build.rs's order (still 130 rows, still four columns) and states the three numbers — the moved files, the remaining src/cuda_kernels.cu, the audit's site count — so "nothing lost" is checkable in each PR rather than only at the end. The non-mutating gates are the same at every stage: audit 130 / --check-fixture exit 0, CUDA unit 567/0/42 + integration 10/0/6, CPU 481/0/36 + 10/0/6, real-model 42/0 on both models, 0 nvcc warnings.

Three measured corrections to the tables above, applied as the stages land:

• guard.cu is 305 lines, not 402. The §4 range 5521–5922 also covered launch_fa_prefill_kv (102 lines), but that launcher launches fa_prefill_kv, whose instantiations live in the FA-prefill section — route (a) puts it in attention_prefill.cu, so G1 takes only the #147/#162 state and its one-owner minfer_launch_* definitions.
• minfer_prewarm_kernels needs five per-family registration functions, not six: mmq_nb, mmq_raw, mmq_bt_q6k, attention_prefill, attention_decode are the only translation units that own template __global__ instantiations the pre-warm address-takes (the other pre-warm entries are plain kernels and stay in the dispatcher, which is where the symbol src/cuda.rs declares it lives). The dispatcher itself moves with mmq_bt_q6k.cu in G3.
• gemm_smem.cu merges into gemm_wmma.cu (937 lines, not 568 + 369): the MINFER_GEMM_OPTIN_SET table, gemm_f16_fn_for and both GEMM launchers address-take and launch gemm_f16_nt_kernel_t instantiations that only gemm_wmma.cu defines. Two files would be the cross-TU template shape that fails to link, so route (a) merges them — 18 family TUs, not the plan's 19.
• The mmq_ksplit_reduce_kernel site appears twice. launch_mmq_raw_nb_bt_nt and launch_mmq_raw_nb_bt_q6k_nt both launch it (fixture rows 108 and 111), so moving the reducer next to the NB kernels leaves the q6_K BT launcher with a cross-TU launch. That is legal where route (a) forbids a cross-TU reference: the reducer is a plain __global__, and only a template instantiation fails without -rdc (blueprint §6 fact 2 vs §7.3).

Step 3 — CPU files

In-place split along the ISA axis that is already there (src/quants.rs mod neon_kernels / mod neon_q8k, src/vec_ops.rs mod neon_f16 / mod neon_vec, plus the inline #[cfg(target_arch = "x86_64")] *_avx2 bodies). No path changes: the three file names stay the module deciders. Three stages, one PR each, quants → vec_ops → kernel:

• 3A src/quants.rs — dot_q4_0 / dot_q4_1 / dot_q5 / dot_q8_0 / kquant / quantize_q8_0 / quantize_q8_k / avx2 / neon (the two inline NEON modules, flattened into the one file), and the inline #[cfg(all(test, target_arch = "aarch64"))] mod neon_correctness promoted to src/quants/neon_correctness.rs — which is also the extraction that lets scripts/check_source_layout.py's first rule be widened to read the cfg predicate instead of the literal #[cfg(test)] (#274).
• 3B src/vec_ops.rs — vec / rms_norm / rope / softmax / silu / f16 (with the inline mod neon_f16) / bf16 / neon (the promoted mod neon_vec).
• 3C src/kernel.rs — dispatch / pool / embed.

Acceptance for each stage: CPU 481 / 0 / 36 + 10 / 0 / 6 (read from docs/status.toml; the plan's older text said 480, which predates #138's two tests), the dead-code (name, kind) set identical on aarch64 and x86_64, and check_source_layout.py / check_dead_code_annotations.py / cargo fmt --all --check green.

Stage A landed (PR #290): src/quants.rs 1,340 → 61 lines + dot_q4_0 71 · dot_q4_1 44 · dot_q5 94 · dot_q8_0 64 · kquant 185 · quantize_q8_0 125 · quantize_q8_k 114 · avx2 81 · neon 407, plus the extracted neon_correctness.rs 137. neon_kernels and neon_q8k are flattened into the one neon.rs (no item name collides), their cross-references lose the super::neon_kernels:: prefix. pub(super) replaces "private to quants" on the items a sibling reaches, so the reachable set is unchanged; the parent's pub use list is what keeps crate::quants::… (and graph/kvformat.rs's two calls) resolving. The old file had 48 fn definitions and the new files have the same 48 (excluding the pre-existing tests.rs); all 29 cpu-map.tsv items resolve in their mapped targets.

Stage B landed (PR #291): src/vec_ops.rs 1,344 → 52 lines + vec 394 · rms_norm 154 · rope 17 · softmax 91 · silu 102 · f16 318 (the inline mod neon_f16 stays nested, its super::F16_SIMD_PATH_CALLS unchanged) · bf16 95 · neon 165 (the promoted mod neon_vec — its items keep pub(super), the same pub(in vec_ops) reach they had as a nested module). 47 fn definitions on each side of the move. Three f16 re-exports (dot_f16_f32, dot_f16_f32_scalar, f16_dot_path, F16DotPath) are #[cfg(test)]: their only consumer is vec_ops::tests, and a non-test pub use of an unused name is an unused_imports error under #![deny(warnings)]. RopeStyle::Interleaved's docs/dead-code-baseline.toml file = field moves to src/vec_ops/rope.rs in the same PR.

Stage C landed (PR #292): src/kernel.rs 667 → 22 lines + dispatch 72 · pool 313 · embed 278. 10 fn definitions on each side of the move. Pool's fields and the MmJob/PoolJob/ParForJob types become pub(super) because dispatch.rs submits through them (the same pub(in kernel) reach they had as private items of kernel), and Pool's gate field carries its hazard comment into pool.rs. One re-export is deliberately not carried: cpu_quant_matmul keeps its pub in dispatch.rs but is not re-exported at the kernel root, because its only caller is cpu_quant_matmul_f32 in the same file and #![deny(warnings)] rejects a pub use of an unused name; no crate::kernel::cpu_quant_matmul path exists anywhere in the tree.

Step 4 — Metal (Mac round, after #255)

src/metal.rs and src/metal.metal are not compiled on Linux (src/main.rs:31-32 gates the module; a failed .metal compile only warns and writes an empty metallib marker), so this step is part of #260 and is verified by the Mac-local gates.

Shape: src/metal/{runtime,encode,ops,policy}.rs + src/metal/kernels/*.metal + *.h — the same rule as CUDA (<backend>/kernels/ = the shader sources), which is also llama.cpp's Metal shape (ggml-metal-device → ggml-metal-ops → kernels/, 22 .metal + common.h/dequantize.h/quantize.h).

Granularity: every .metal file at or below ~800 lines. The 5,151-line metal.metal regroups as:

build.rs compiles the parts into one metallib, and the runtime newLibraryWithSource fallback (src/metal/runtime.rs include_str!) needs the parts joined (concat!) or a thin umbrella source; both entry points must see the same set. Then #53 (reserve/assign + the DeviceMemory report) lands on top of the new layout.

Landed (2026-10-05, macbook (macOS 27.0.1, Apple M4 Pro)) — three increments as Addendum 2 asked. Increment 1 made build.rs's shader set an explicit SHADER_SOURCES list plus a check_shader_file_list() guard, concatenated it into $OUT_DIR/minfer.metal, and pointed the runtime include_str! at the same file (shader set unchanged; the metallib stayed byte-identical, 13af518e…). Increment 2 moved one family per commit (16 commits) and deleted src/metal.metal; the generated source is byte-for-byte the old one (212 923 B, 5 151 lines, 61 kernel void names, same set), and the metallib hash moved to 7a4a7cd4…. Increment 3 split src/metal.rs with 0 visibility edits — the type definitions and the private dispatch primitives stay in the parent module, so the children reach them without pub(super) (the CudaState-stays-in-cuda.rs shape) — and moved the test-only matmul_on_gpu_buf with them so src/metal/tests.rs still resolves it.

Three measured corrections to the table above, applied as the increments landed:

• quantize.h is not created. The plan's dequantize.h + quantize.h row assumed GPU-side quantize helpers; the tree has none (the only quantize strings are comments — activations are quantized on the CPU). The pragmatic source of truth is the tree, so src/metal/kernels/ holds common.h + dequantize.h only; a quantize.h with no helper would be a file the guard must list and nothing reads.
• get_rows.metal includes the warm-up kernel (2 340–2 523), so it is 184 lines, not 169.
• rope.metal is 37 lines (the 2 743–2 746 banner is not adjacent to kernel_rope_f32, which is at 2 876–2 908 after the P1 parallel-attention section); fa_parallel.metal is 129 lines (2 747–2 875). The plan's "≈60 / 162" rows mixed the two.

A fourth compile entry point the ticket did not name: the four tests/*_isolation.rs integration tests include_str! the shader source and compile it themselves (9 sites). They now include_str! the same $OUT_DIR/minfer.metal, so all four consumers — build.rs, src/metal.rs's fallback and the integration tests — see one file set by construction.

Step 5 — close the loop

Re-measure the N-module cold start and update the §2.4 pre-warm table (#225); implement allocplan::DeviceMemory for Metal (#53, CUDA already answers it) so that "the device memory report" becomes the first interface with two real implementations; add the mechanical documentation-anchor check (§6) if it is not already landed.

Landed. The MINFER_OP_TIMING re-measure is §5.1 (Linux half, PR #293); the anchor checker is #266. The Metal DeviceMemory half landed on macbook (macOS 27.0.1, Apple M4 Pro) (PR #296) — the measured record (before/after counts, the red-baseline note, the mutation transcript) is the E4/E5 "Metal half" record in docs/ARCHITECTURE-EXECUTION-PLAN.md, and the common decision it unblocks is §1.3 rule 3 above: two implementations, three same-semantics callers, no new module.

Step 6 — the long test files (in scope, decided 2026-10-04)

The largest files in the crate are tests: src/graph/cuda_backend/tests.rs (8,384 lines, 143 tests), src/models/qwen2/graph/tests.rs (3,374), src/server/batch/tests.rs (2,629), src/cuda/issue162_tests.rs (1,186), src/graph/alloc/tests.rs (1,773), src/graph/kvcache/tests.rs (1,134), src/conversation/tests.rs (1,156), src/tooling/tests.rs (1,669), src/sampler/tests.rs (1,232), src/graph/cuda_backend/tests.rs … — this step splits them by op family / topic into <module>/tests/<topic>.rs (the same rule: every file named by a mod declaration, checked by scripts/check_source_layout.py).

• Why last: src/graph/cuda_backend/tests.rs is the evidence base for Steps 1–2 (its 143 tests are what proves the moves), and every later PR's line references would churn if it moved first.
• Order inside the step: graph/cuda_backend/tests.rs first (the largest), then the other >1,000-line test files, one PR each.
• Acceptance: the test counts are identical (nothing added or removed — the same tests run from new files, which check_source_layout.py's rule 2 is precisely there to guarantee), plus the step-appropriate gates (CUDA unit 565/0/42 for the executor tests, CPU 480/0/36 + 10/0/6 for the rest).
• Note: no size ratchet is added (decided 2026-10-04) — the ~800-line target in this document is guidance, enforced by review, not by a script.

5. Why no -rdc=true, and why route (a)

Measured on dgxspark (aarch64, GB10 sm_121), CUDA 13.0, 2026-10-03 (two-file probe, four cross-TU patterns, each compiled and run):

So the default toolchain forces "the launcher lives in the TU that instantiates the kernel" — which is route (a) and is also llama.cpp's CUDA shape. The evidence base for the split's other costs:

• one nvcc invocation today, 13 -gencode targets (12 SASS + compute_121 PTX); serial compile 117.3 s / 114.1 s (two runs), 675 MB RSS, 35.8 MB object; nvcc --threads 0 16.4 s / 15.6 s;
• the CUDA build is not bit-reproducible today (two identical serial runs differ by 16 bytes in the .text of two cubins) — so the split's evidence is runtime gates, not binary identity;
• the recorded per-module fatbin load was ~2.2 ms, set-size independent (docs/CUDA-BACKEND-DESIGN.md §2.4's cost table), so ten modules were extrapolated at ~13–22 ms one-time. Step 5 measured it (2026-10-05) and the extrapolation was wrong — see below.

5.1 The re-measured module load (Step 5, 2026-10-05)

Box dgxspark (aarch64, GB10 sm_121), CUDA 13.0, driver 580.178.04, nvcc 13 -gencode targets (12 SASS + compute_121 PTX) × the 17 TUs; nvidia-smi before the run: SM clock 2 411 MHz idle (warm) / 208 MHz (cold), 0 % util, no other compute process. Two binaries, both built in this repository's worktrees: pre-split = bc30152 (930e1e2^, src/cuda_kernels.cu = one 10,215-line TU = 1 fatbin module) and post-split = 4900298 (17 TUs = 17 modules). Command of record (<0.5B> = the cached qwen2.5-0.5b-instruct-q4_0.gguf):

MINFER_OP_TIMING=1 target/release/minfer <0.5B> "hello"

which prints the prefill-GEMM smem pre-warm loop's own duration — the point at which the fatbin's module is finalized. fresh process = a new minfer invocation; cold = the binary's and the model's pages evicted with posix_fadvise(POSIX_FADV_DONTNEED) first (an agent shell cannot drop_caches), warm = back-to-back fresh processes with the pages resident.

Per-module figures (post-split, warm, one live cudaFuncGetAttributes per TU, temporary probe inside minfer_prewarm_kernels — reverted before the PR): mmvq_q6k 57 · mmvq_aquant 57 · mmq_raw 70 · mmq_nb 73 · matmul_f32act 85 · mmq_int8 93 · mmvq_skipwrite 100 · kv_store 102 · mmvq_multi 111 · mmq_bt_q6k 118 · ops_misc 123 · ops_elementwise 245 · attention_decode 223 · gemm_fused_dequant 40 · attention_prefill 447 · gemm_wmma ≈ 350–590 µs (its line also carries the 12 attribute queries), i.e. 40–450 µs per module, ~150 µs median — not the recorded 2.2 ms, which was the whole pre-split module's cost, not a per-module constant.

Verdict: the split's cost is acceptable. In steady state the 17-module fatbin loads in the same ~2.3 ms as the pre-split single module, because the load tracks the code a module contains, not the module count; the MINFER_OP_TIMING line moves from 2.3 ms to 0.4 ms only because it now times one seventeenth of the work. The honest extra is cold: a page-cache-cold start pays ≈ +25 ms (≈ 19 ms → ≈ 45 ms), because every one of the 16 module registrations faults the fatbin's pages again. The cold per-module figures (same probe, evicted pages, 975–5 214 µs) are ~10× their warm values across the whole size range — including 975 µs for the 284-line gemm_fused_dequant.cu — so a per-registration overhead rides on top of the size-proportional part, and 16 registrations pay it 16 times. That is ≈ 1.8 % of the ~1.4 s cold-start wall time on this 0.5B model, once per process, and the mitigation the plan names (minfer_prewarm_kernels trimmed to the modules a run needs) would only move it into the first forward's lazy loads — which is why no trimming is applied. The correction also applies to §2.4's old explanation of the 14.5 ms cold row: it is page-cache-cold, not the GPU clock (a 40 s idle cooldown at a 208 MHz SM clock reads the warm 2.3 ms; the same binary with evicted pages reads 18.5–20.7 ms).

The full per-module transcripts and the two worktrees' build logs are in the Step 5 record in docs/ARCHITECTURE-EXECUTION-PLAN.md.

6. Documentation plan

6.1 What the split invalidates

6.2 Policy: live documents are edited, historical records are frozen

• Live documents (the ones a maintainer reads to find code): edited in the step that moves the code, with anchors converted to symbol anchors (`prefill_mmq` (`src/cuda/kernels/mmq_*`)) wherever the line number was only a locator.
• Historical records (docs/cuda_optimization_steps/*.md, docs/QWEN2.5-*.md, docs/DEBUGGING-*.md, docs/KNOWN-CPU-ISSUES-*.md, docs/PARAMETER_AUDIT.md's older tables, docs/ARCHITECTURE-EXECUTION-PLAN.md's per-ticket entries, and experiments/cuda/*.md — the probe run records, which quote the nvcc … ../../src/cuda_kernels.cu command as it was run) keep their text — they record a measurement taken against a revision, and rewriting them would falsify the record. They are resolved through the path mapping table this document keeps (§6.4).
• The machine ledgers the step records name are of their day. They moved beside their checkers on 2026-10-09/10 (ADR-0023, ADR-0024) and now live at scripts/status.toml, scripts/test-baselines.toml, scripts/dead-code-baseline.toml and tests/fixtures/f6-fixtures.json. §6.1's "machine-checked today" row names the current reader; the step text above keeps the paths it was written with.
• The checker must know about the freeze: scripts/check_doc_line_anchors.py (ticket 6) carries a frozen-file set (the GRANDFATHERED_BARE pattern), so a frozen record does not fail CI, and the set can only shrink. This is the one design constraint the frozen policy puts on ticket 6.

6.3 Per-step update table

docs/status.toml is not edited by Steps 0–6: the suite counts do not move (code moves, tests move, no test is added or deleted). If a step ever changes a count, scripts/check_status.py --check must be updated in the same PR, and this document says so in that step's record.

6.4 The path mapping table (lives here; grows per step)

The frozen records resolve old paths through this table, and the live sweeps are generated from it:

6.5 The macOS hand-off (decided 2026-10-04: Step 4 is Mac-local)

Step 4 cannot be executed or verified on the Linux box, so this document must be sufficient alone for a macOS agent: §3's rule, §4 Step 4's file table, §5's cross-TU constraint, §6.3's doc sweep, and §10's verification row. The ticket (T5) and #260 both link here, and Step 4's record names the Mac box explicitly (gate-contract rule 5: an absolute box label, e.g. macbook (macOS 15.x, Apple M4)), never "this box".

7. Interaction with the open issues (as of 2026-10-03, 38 open)

Symbol-level scan of all 38 issue bodies against the identifiers defined in the files to be split (601 distinctive symbols; plus a direct src/<file>:NNN path scan). 18 issues reference affected code or files, or target code that this plan moves. (#150's worker_loop hit is the server's worker_loop_serial, not kernel.rs — counted as unaffected.)

7.1 Must be sequenced against this plan

7.2 Needs a body/anchor update only

#137 (async staging, copy_cross/await_cross + Metal), #135 (walkthrough/architecture stale Backend enum — same docs), #54 (re-run Metal gap measurements), #52 (mixed-quant QKV epilogue), #39 (debug_assert! in release on Metal), #231 (five macOS-only attention call sites, 8 src/…:NNN references).

7.3 Unaffected

#215, #209, #205, #204, #203, #198, #195, #179, #157, #150 (server-side worker_loop_serial, not kernel.rs), #133, #132, #126, #125, #118, #103, #62, #40, #38, and #229 (allocator dead-code bookkeeping only).

7.4 Tickets this plan files (decided 2026-10-04: filed now, labelled)

That is 7 tickets, all filed 2026-10-04; every sub-ticket carries Part of #261, and #261 carries the Step −1…6 checklist, the interaction table of §7.1, the target tree of §8 and the documentation plan of §6. The former "stale size/claim sweep" ticket was folded into #219 as a comment (decision 3, §9).

8. Resulting tree

src/kernel/, src/quants/, src/vec_ops/, src/metal/ and src/cuda/ already exist today — they hold only tests.rs (plus metal/mmap_align_test.rs and cuda/'s ten issue probes). The plan therefore does not create a new convention: the production parts simply join the directories that are already there. [S1]…[S4] name the step that produces each entry; a parenthesised line count is the file's length after the step that created it, and the .cu/.metal ranges refer to the pre-split file.

src/
├── main.rs                                    (unchanged)
├── cuda.rs                          [S1]  module cuda: doc + `AttnWindow` + `pub struct CudaState`
│                                             (fields stay private) + free items + `mod methods;`
│                                             `mod ffi_runtime;` + the `#[cfg(test)] pub(crate) use`
│                                             re-exports + the ten `#[cfg(test)] mod` declarations (1 014)
├── cuda/                                    (exists: 10 test files today, unchanged)
│   ├── methods.rs                   [S1]  the 15 cross-family helpers + the 18 `mod` declarations
│   │                                        + the `#[cfg(test)] pub(crate) use` re-exports (323)
│   ├── methods/
│   │   ├── accounting.rs            [S1]  B   42   weights_bytes / device_memory
│   │   ├── attention.rs             [S1]  P  516   gqa / split / batched / prefill
│   │   ├── buffers.rs               [S1]  E   28   cuda_malloc / cuda_free
│   │   ├── capture.rs               [S1]  H  137   CUDA-graph capture / replay
│   │   ├── copy.rs                  [S1]  F  182   H2D / async / D2H / pinned / D2D
│   │   ├── dispatch.rs              [S1]  I  440   matmul_f32_ptr* + the MMQ dispatch tree
│   │   ├── elementwise.rs           [S1]  O  195   norm / add / mul / silu / swiglu / rope
│   │   ├── events.rs                [S1]  G  183   events, async staging, sync, latch
│   │   ├── gpu_act.rs               [S1]  N  233   on-GPU quantize / gather / embed
│   │   ├── init.rs                  [S1]  A  331   device probe / tier / singleton
│   │   ├── kvstore.rs               [S1]  R  350   KV store + fused QKV epilogue
│   │   ├── mmq_quant.rs             [S1]  K  306   A-quantize + MmqCache
│   │   ├── mmvq.rs                  [S1]  Q  728   decode MMVQ + q8_0 p32 planes
│   │   ├── policy.rs                [S1]  J  113   MMQ gate predicates
│   │   ├── prefill_f16.rs           [S1]  M  307   f16 GEMM + w16 cache
│   │   ├── prefill_mmq.rs           [S1]  L  488   auto_ksplit, prefill_mmq
│   │   ├── stream.rs                [S1]  D   67   bound/context stream, create/destroy
│   │   └── weights.rs               [S1]  C  726   register_weight + q6k/q4k expansion
│   ├── ffi_runtime.rs               [S1]  cudart/driver FFI (`pub(crate)`) + the test-only extern
│   │                                        block (145)
│   └── kernels/                     [S2]  the CUDA translation units (kernels + their host
│       │                                   launchers; `<backend>/kernels/` is the one rule both
│       │                                   device backends share)
│       ├── common.cuh               [S2]  ~450: defines + device helpers + KV load idiom +
│       │                                   declarations of the #147/#162 helpers
│       ├── guard.cu                 [S2]  5521–5922  single owner of the #147/#162 state and of
│       │                                   the minfer_launch_* definitions (external linkage)
│       ├── matmul_f32act.cu         [S2]  63–693            ≈750 with launchers
│       ├── mmvq_aquant.cu           [S2]  694–1094          fused-producer A-quantize prepass
│       ├── mmvq_skipwrite.cu        [S2]  1095–1651         mode-2 skip-write variants
│       ├── mmvq_q6k.cu              [S2]  1652–1942         pipelined + dense split-plane q6_K
│       ├── ops_misc.cu              [S2]  1943–2375         padded Q6_K, gather/embed, f32×f32, f16×f32
│       ├── ops_elementwise.cu       [S2]  2376–2654         quantize f32→Q8_0, norm, bias, add/mul,
│       │                                                     SiLU/SwiGLU, i32 decode, RoPE
│       ├── kv_store.cu              [S2]  2794–3084 + 10173–10215  KV store, fused QKV epilogue,
│       │                                                     arena row move
│       ├── attention_decode.cu      [S2]  3085–3677         GQA, E1 window, kv_map, split-K, batched
│       ├── attention_hybrid.cu      [S2]  3678–4047         hybrid rpw (hd 128, f16 KV)
│       ├── attention_prefill.cu     [S2]  5136–5520         FA-style prefill (staged KV)
│       ├── gemm_wmma.cu             [S2]  5923–6490         dequant-to-f16 + wmma HGEMM
│       ├── gemm_smem.cu             [S2]  6491–6859         dynamic-smem formula + checked opt-ins
│       ├── gemm_fused_dequant.cu    [S2]  6860–7135         8p fused dequant-in-GEMM
│       ├── mmq_int8.cu              [S2]  7136–7540         R1 int8 MMQ prefill GEMM
│       ├── mmq_raw.cu               [S2]  7541–8074         P6 raw-byte MMQ
│       ├── mmq_nb.cu                [S2]  8075–8667         raw-nibble NB + A-layout transform
│       ├── mmq_bt_q6k.cu            [S2]  8668–9407         r38 q6_K BT
│       └── mmvq_multi.cu            [S2]  9408–10172        multi-token MMVQ + doc103/104
├── metal.rs                         [S4]  module metal: doc + free items + `pub use`
├── metal/                                   (exists: mmap_align_test.rs + tests.rs)
│   ├── runtime.rs                   [S4]  L1: MpsState / MetalDevice / library + pipeline cache
│   ├── encode.rs                    [S4]  L2: MpsCommandBuffer encoding
│   ├── ops.rs                       [S4]  L2: the op → encoding table
│   ├── policy.rs                    [S4]  pure predicates (MINFER_METAL_* / MINFER_* knobs)
│   └── kernels/                     [S4]  L3, ≤~800 lines each:
│       ├── common.h · dequantize.h          (no `quantize.h` — the plan's row
│       │                                    assumed GPU-side quantize helpers
│       │                                    and the tree has none; §4 Step 4)
│       ├── mul_q4_0_q8_0.metal · mul_f32act_q4q5.metal · mul_f32act_kquant.metal
│       ├── mul_mm.metal · mul_mm_kq.metal · get_rows.metal · norm_elementwise.metal · rope.metal
│       ├── kv.metal · qkv_fused.metal
│       ├── fa_parallel.metal · fa_split.metal · fa_decode.metal · fa_prefill.metal
│       ├── f16.metal (#164) · f32.metal (#317) · bf16.metal (#208) · attn_window.metal (#44a)
│                                            (the four added by the Metal round after S4)
│                                            (runtime fallback joins them with `concat!`)
├── kernel.rs                        [S3]  module kernel: `mod` + `pub use`
├── kernel/
│   ├── dispatch.rs                  [S3C] cpu_quant_matmul / cpu_quant_matmul_f32 (12–44, 322–390)
│   ├── pool.rs                      [S3C] Pool / par_for / set_cpu_threads (45–321, 364–390)
│   ├── embed.rs                     [S3C] embed_tokens (391–)
│   └── tests.rs                            (exists)
├── quants.rs                        [S3]  module quants: `pub use`
├── quants/
│   ├── dot_q4_0.rs · dot_q4_1.rs · dot_q5.rs · dot_q8_0.rs    [S3A]
│   ├── kquant.rs                    [S3]  Q4_K/Q5_K/Q6_K dots
│   ├── quantize_q8_0.rs · quantize_q8_k.rs                    [S3A]
│   ├── neon.rs                      [S3A] was `mod neon_kernels` / `mod neon_q8k`, flattened
│   ├── avx2.rs                      [S3A] was the inline `*_avx2` bodies
│   ├── neon_correctness.rs          [S3A] was the inline `#[cfg(all(test, aarch64))] mod
│   └── tests.rs                            (exists)
├── vec_ops.rs                       [S3]  module vec_ops: `pub use`
├── vec_ops/
│   ├── vec.rs · rms_norm.rs · rope.rs · softmax.rs · silu.rs   [S3B]
│   ├── f16.rs                       [S3B] the f16 dot/matmul + the nested `mod neon_f16`
│   ├── bf16.rs                      [S3B] bf16 row decode + matmul
│   ├── neon.rs                      [S3B] was `mod neon_vec`
│   └── tests.rs                            (exists)
├── graph/                                   (unchanged: `backend.rs` + `registry.rs` stay the one
│                                             device seam; `*_backend.rs` stay the executors)
└── … (all other modules unchanged)

Non-src/ changes that ride along: build.rs (CUDA/Metal file lists + one rerun-if-changed per .cu/.cuh/.metal/.h + one .o per .cu, still one libcuda_kernels.a; Metal parts → one metallib; plus the new-file guard that a file in kernels/ cannot silently be absent from the list), scripts/check_cuda_launch_returns.py (directory discovery + the "no <<<>>> in a .cuh" assertion) + tests/fixtures/cuda_launch_sites.tsv (regenerated in build-list order, still 130 rows, still four columns), scripts/check_doc_line_anchors.py (new, ticket 6), docs/SOURCE-LAYOUT-PLAN.md (this file) + AGENTS.md Layout block + docs/ARCHITECTURE.md (layer definition + interface-eligibility rule) + the path/anchor sweeps in docs/BACKENDS.md, docs/CUDA-BACKEND-DESIGN.md, docs/inference_e2e_walkthrough/{07,15}*.md, docs/GPU_SAFETY.md, docs/BUILD.md, docs/ARCHITECTURE-ROADMAP.md, docs/CUDA_OPTIMIZATION.md, docs/LLAMA-CPP-MMQ-ANALYSIS.md, docs/ARCHITECTURE-EXECUTION-PLAN.md.

9. Decisions (all taken — nothing open)

10. Verification matrix

Suite counts are read, not remembered. The numbers in the table above are the pre-#138 rows; the current rows are docs/status.toml (after #138: CUDA 567 / 0 / 42, CPU 481 / 0 / 36 on dgxspark (aarch64, GB10 sm_121) and 479 on the CI runner, integration 10 / 0 / 6). Every step must read the rows from that file rather than quote this document, and a step that moves a row updates AGENTS.md and docs/status.toml in its own PR.

Standing rule for every step: a layout PR moves code and nothing else; anything else it notices is filed, not fixed inside it.

The layering convention and the test-module rule (moved from AGENTS.md)

The backend layers: each device backend is organised as L1 runtime, L2 launch/dispatch, L3 kernel sources and L4 graph executor, and only L4 is polymorphic — graph/backend.rs + graph/registry.rs stay the single device seam, so the directory tree is device-first and the layer is the second axis inside each device (<backend>/kernels/ is where kernel sources live). A shared common is added only when two backends implement it and two callers use it (allocplan::DeviceMemory is the one candidate today). CPU's quants.rs/vec_ops.rs are deliberately not a device-private layer: they are the crate's numeric kernel library, shared with graph/kvformat.rs and graph/cuda_backend.rs. The four long backend files are split along this convention — the three CUDA files, the CPU trio and the two Metal files are done (src/cuda/, src/cuda/kernels/, src/quants/, src/vec_ops/, src/kernel/, src/metal/, src/metal/kernels/). Plan and target tree: docs/SOURCE-LAYOUT-PLAN.md (#261); the file list above is updated by each step as the code moves.

src/graph/: mod.rs ComputeGraph/CNode · ops.rs Op + NodeMeta · builder.rs GraphBuilder · scheduler.rs assign → split → execute · backend.rs + cpu_backend.rs/metal_backend.rs/cuda_backend.rs executors · registry.rs backend registry (F4; F5's copy_cross/await_cross) · alloc.rs liveness allocator + persistent KV regions (E4) · allocplan.rs size-class ladder + pure plan + DeviceMemory/budget_decision · offload.rs layer offload plan (E5) · kvcache.rs cell store, removal/shift/compaction, span list, prefix sharing (C1–C3, C8b) · kvformat.rs KV format + MINFER_CACHE_TYPE gate (C4) · kvsession.rs versioned KV session container (C5) · cache.rs/params.rs params-only graph reuse · fusion.rs SwiGLU fusion · copystats.rs split-boundary counters · batch.rs batch composition · dot.rs/json.rs exporters.

Unit tests live beside their module as <module>/tests.rs, declared #[cfg(test)] mod tests;, so a non-test build does not parse them (e.g. src/graph/alloc/tests.rs for src/graph/alloc.rs). src/graph/op_matrix.rs was already this pattern. Note this buys build hygiene, not speed: a warm cargo check --release --features cuda measured 1.43–1.45 s with the tests inline vs 1.52–1.57 s extracted. Both halves are enforced: scripts/check_source_layout.py (CI check-docs) rejects an inline #[cfg(…test…)] mod … { (the cfg predicate is read, so a compound #[cfg(all(test, …))] is caught too), and rejects a src/**.rs that no mod declaration names — the latter is the quiet one, because an undeclared file is never compiled and the tests inside it would silently not run. A long test module is split further into <module>/tests/<topic>.rs, each topic declared by a mod <topic>; in its tests.rs (src/graph/cuda_backend/tests/ is the first, #267); the same declaration walk reaches those files, so the orphan rule guards the split as well.

Decisions governing this document

This page is the current contract; the decisions behind it are frozen in the ADR corpus:

• ADR-0012 — Device is the first axis, the layer the second — and no premature common
• ADR-0005 — Metal becomes a first-class backend
• ADR-0023 — Each machine ledger lives beside its checker, one per prose target