# Feature — Gemma-4 A4B Dual Chunk Attention (DCA), training-free 1M context

> **Status (2026-06-10): DESIGN LANDED, patch `0078` not yet captured.** Feature home for the
> forthcoming `vendors/patches/llamafile/0078-gemma4-dca.patch`. The patch capture + build + RULER
> gate run in an **out-of-tree clean-room** (a separate llamafile checkout) — see "Capture without
> reset" — because the live `vendors/sources/llamafile` working tree holds the only applied source of
> patches `0001–0077` and MUST NOT be `reset`/`clean`ed (buglog 2026-06-04: `build:llamafile:reset`
> = `git reset --hard` + `git clean -fdx` wipes uncommitted patch source). T87.pD.

## Goal
Extend Gemma-4 26B-A4B (and 31B dense) to ~1M context **training-free**, matching Qwen2.5-14B-1M's
RULER class (92.5 @ 1M), by porting **Dual Chunk Attention** (DCA, the Qwen 1M mechanism) into the
llamafile graph. Runs on consumer GPUs (RTX 3090) — the path our users actually have — instead of
the parked vLLM FA4-sm120 route (which needs an unfinished CuTe Blackwell kernel; see backup_models
task #592). DCA composes with YaRN mscale for the final 2×–4× past the chunk grid.

## Why llamafile is sufficient (feasibility verified 2026-06-10, read-only)
Gemma-4 26B-A4B attention is **mixed-head** (config + tensor-shape confirmed):
- **25 sliding layers @ head_dim=256** (window 1024, θ=1e4) — bounded by their window, need NO
  long-range handling.
- **5 global/full layers @ head_dim=512** (idx **5, 11, 17, 23, 29**; `global_head_dim=512`,
  `num_global_key_value_heads=2`, `attention_k_eq_v=true` → no `v_proj`, K=V; `partial_rotary_factor
  0.25` → n_rot=128 of 512; θ=1e6, `rope_type=proportional`). **These carry all long-range
  information → DCA touches ONLY these 5 layers.**
- No `attn_logit_softcapping` (Gemma 3/4 dropped it; only `final_logit_softcapping=30` on the LM
  head) → the chunk merge is the plain no-softcap online-softmax.

Every primitive DCA needs already exists in `vendors/sources/llamafile/llama.cpp` (RULER-validated):
- **Proportional rope** — `src/models/gemma4-iswa.cpp:56-59` feeds `model.layers[il].rope_freqs` as
  `freq_factors` to `ggml_rope_ext` for `!is_swa` layers; `rope_freqs` loads at
  `src/llama-model.cpp:2955`. Per-layer `freq_base_l/freq_scale_l/n_rot_l` (`gemma4-iswa.cpp:47-49`).
- **YaRN** — `ggml_rope_ext(..., ext_factor, attn_factor, beta_fast, beta_slow)` (`gemma4-iswa.cpp:73,97`)
  + hparams `n_ctx_orig_yarn`/`yarn_*` + `rope_scaling_type ∈ {none,linear,yarn,longrope}`.
- **hd=512 flash-attention** — patch `0075-d512-turbo-vec` (D-generic FA-VEC for the global layers).
- **LSE-emitting FA + exact online-softmax merge** — patch `0070-rolling-kv`: `launch_fattn` per-row
  `dst_lse = max + logf(rowsum)` via the `opencoti_fattn_dst_lse` consume-once channel +
  `GGML_OP_STREAMING_FLASH_ATTN` + `streaming_combine_kernel`. **This is DCA's chunk-partial merge**;
  generalizes from 2 regions (window⊕tail) to N chunk-groups.
- **Training-free position-remap precedent** — `grp_attn_n/grp_attn_w` (LongLM Self-Extend),
  `tools/completion/completion.cpp:511-517`. Same CLASS as DCA. **Caveat: completion-tool-only**, so
  DCA hooks at the **graph** `build_attn` seam (serves the server opencoti drives).
- Clean per-layer seam: `build_attn(inp_attn, wo, …, Qcur, Kcur, Vcur, …, f_attention_scale, il)`
  (`gemma4-iswa.cpp:102`).

**No blocking limitation.** Nuances handled by design: extension config (YaRN factor / `n_ctx_orig`)
via GGUF metadata or CLI; Self-Extend is completion-only so we hook the graph; our A4B GGUFs already
carry `rope_freqs` (they serve at 256k today).

## DCA algorithm (global layers only)
Split positions into chunks of size `c` (`c` < trained length; e.g. 65536). For each global layer,
attention is computed in position regimes merged by LSE:
- **intra-chunk** — queries attend keys in the SAME chunk at LOCAL positions `[0,c)` (preserves the
  trained short-range rope behavior).
- **inter-chunk** — queries attend EARLIER chunks at a clamped inter-chunk (chunk-index) position so
  relative positions stay inside the trained range.
- **successive-chunk** — the immediately-preceding chunk uses a continuous bridging position so the
  intra/inter seam doesn't jump.
On Gemma-4 only the **rotary 128 dims** (partial_rotary 0.25) are remapped; the other 384 pass
through. Each regime = a separate rope'd-Q × KV attention emitting LSE; the partials merge via the
existing online-softmax combine. YaRN mscale (attn_factor) optionally rides on top for the final
2×→4×.

## Patch `0078-gemma4-dca` structure
Base for capture: live tree = `0001-0042 + 0070..0077` applied. Number `0078` (applies AFTER `0077`).

### A. Flags (additive, default-OFF → byte-identical when off)
`common/arg.cpp` + `common/common.h` + `src/llama-cparams.h`: `--dca {off,on}` (default off),
`--dca-chunk-size N` (default = n_ctx_train ⇒ inert: `c ≥ seqlen` collapses to a single intra
regime = plain attention), optional `--dca-yarn-factor F`. `LLAMA_ARG_DCA*`, `LLAMA_EXAMPLE_SERVER`.
New `cparams`: `bool dca_enabled; uint32_t dca_chunk; float dca_yarn_factor;`. OFF path grep-clean
byte-identical to pre-0078.

### B. Graph hook (one surgical marker) — `src/models/gemma4-iswa.cpp`
For `!hparams.is_swa(il) && cparams.dca_enabled` ONLY, route attention through the new additive
builder instead of the single `build_attn(...)`. Marker `// opencoti F5 gemma4-dca`. Sliding layers
and `dca_enabled==false` keep the unchanged `build_attn` → OFF byte-identical.

### C. Additive builder — NEW `src/models/gemma4-dca.cpp` (+ decl in `src/models/models.h`)
`build_attn_dca_gemma4(...)`: builds intra/successive/inter rope'd-Q variants (3× `ggml_rope_ext`
on `inp_pos`-derived remapped position tensors), runs 3 `build_attn`-style FA calls over the global
KV emitting `dst_lse` (reuse `0070`'s channel) with chunk masks, merges the 3 partials with a
`streaming_combine`-family reduce (extend the 2-region form to N-region). Keeps the in-tree hook to a
single marker (NEW file per UPSTREAM_SYNC §Additive).

### D. Position inputs — additive `llm_graph_input_dca`
Mirrors `llm_graph_input_mtp` (0074) / position-window inputs (0070): holds intra/inter/succ position
vectors derived from `inp_pos` + `dca_chunk`. Registered in graph-input assembly; no upstream-source
edit beyond registration.

### E. Acceptance gate — NEW `.opencoti/dca-ruler-gate.sh`
- **OFF byte-identity**: `--dca off` greedy decode byte-identical to the pre-0078 binary.
- **Below-one-chunk no-op**: at ctx ≤ 256k with `--dca-chunk-size ≥ ctx`, `--dca on` RULER vt/niah
  == `--dca off` (DCA must be a no-op when `c ≥ seqlen`). Uses the RULER native runner.
- **Extension win**: RULER vt/niah at 512k & 1M with `--dca on` vs YaRN-only degradation — target
  the T87 1M class (≥ Qwen2.5-14B-1M's 92.5@1M behavior on the global layers). Uses the
  merge→serve→RULER-ladder driver.
- Acceptance = logit-distribution / RULER-score equivalence, NOT greedy byte-equality past one chunk
  (online-softmax is fp-non-associative — same standard as `0070`).

### F. README row + UPSTREAM_SYNC
Add the `0078-gemma4-dca` row to `vendors/patches/llamafile/README.md` (## Current series) and
register the single `// opencoti F5 gemma4-dca` marker in `docs/protocols/UPSTREAM_SYNC.md` (same
mechanism as `0076`/`0077`). Provenance: opencoti-original; algorithm from the DCA paper
(Qwen2.5-1M) + the deleted v0 vLLM `dual_chunk_flash_attn.py` reference (read, not transcribed).

## Capture without reset (the live tree holds the only patch source)
The standard `bun run build:llamafile:{check,apply,build}` flow assumes a pristine submodule
(`git apply --check`/end-to-end apply), which the live tree is NOT (0001-0077 already applied) and we
must NOT `reset` it. Therefore:
1. **Clean-room build/gate**: in a SEPARATE out-of-tree llamafile checkout (not the held submodule),
   `reset → bootstrap → apply 0001-0077 → apply 0078 edits → build (cosmocc + ggml-cuda DSO) → run
   gate E`. The held `vendors/sources/llamafile` is never touched.
2. **Capture the diff** against an out-of-tree base **copy** of the touched files (NOT `git apply`
   from inside an in-tree baseline — that is a silent no-op, buglog bug-431; use `diff -u base live`
   to produce the patch and `patch -p1` to re-apply into a copy when verifying).
3. **Verify round-trip**: `0078-gemma4-dca.patch` reverse-applies clean vs the post-0078 tree and
   forward-reproduces it byte-for-byte (0 diffs), both directions.
4. **Land** only the artifacts in the opencoti repo: `0078-gemma4-dca.patch`, the README row, the
   UPSTREAM_SYNC entry, `.opencoti/dca-ruler-gate.sh`, this doc. The held submodule working tree
   stays exactly as the other session left it.

## Cheaper baseline to A/B first
`grp_attn` (Self-Extend) already exists; A/B its RULER@512k for A4B as the bar DCA must beat. Blocker:
completion-tool-only → needs a tiny server-path `grp_attn` hook OR running the completion tool for
the probe. Document the gap vs full DCA before committing to the 0078 build.

## Section-C merge: implementation decision (grounded 2026-06-10, read-only)

The graph seam (`gemma4-iswa.cpp` global block, `!is_swa(il)`): `rope_ext(Qcur/Kcur, inp_pos, freq_factors=rope_freqs, n_rot_l, freq_base_l=1e6, ...)` → `build_attn(inp_attn, wo, …, Qcur, Kcur, Vcur, …, f_attention_scale, il)`. Hook here for `dca_enabled && !is_swa`.

0070 primitives available host-side: `ggml_streaming_flash_attn(q,k,v,mask,scale,max_bias,softcap)` (emits O+LSE), `ggml_streaming_flash_attn_window(q,k_win,v_win,k_tail,v_tail,mask,…)` (2-region FUSED FA+combine), `ggml_flash_attn_ext_tail_partial(q,k,v,…)` → un-normalized partial `[2+DV, n_head, n_q, n_b]` = (m, s, VKQ). **There is NO standalone host-exposed N-way (O,lse) combine** — the online-softmax merge is fused inside `GGML_OP_STREAMING_FLASH_ATTN` (`fattn.cu streaming_combine_kernel`).

DCA needs to merge **3** regime partials (intra/successive/inter). Two paths:
- **C1 — graph-math merge (host-only, PREFERRED, no DSO rebuild):** per regime, emit the un-normalized partial (m_i, s_i, O_i·s_i); combine N partials with a small subgraph of EXISTING ggml ops: `m=max_i(m_i)`, `w_i=exp(m_i−m)`, `S=Σ w_i s_i`, `O=Σ w_i O_i s_i / S`. All elementwise/reduce ops already in ggml → compiles into the host binary, executes on the existing fat `ggml-cuda.so 0.10.1` (sm_86+sm_120). Build = cosmocc `make` only. **Risk:** must get each regime's FA to emit the GPU partial to a graph tensor (verify `ggml_streaming_flash_attn` partial-output mode is graph-addressable, not just the thread-local channel; the `_tail_partial` op is CPU-only — too slow at 512k, so the GPU streaming op's partial path is required).
- **C2 — new N-way combine kernel (DSO rebuild):** add `ggml_streaming_combine_n(partials[])` + a `streaming_combine_n_kernel` generalizing the 2-region kernel → rebuild `ggml-cuda.so` via `cuda.sh`. Cleaner/faster but heavier build + a CUDA-source edit (per UPSTREAM_SYNC, a NEW kernel file to keep the in-tree hook minimal).

**Decision: implement C1 first** (host-only build, fastest to a gated result); fall back to C2 only if the GPU partial isn't graph-addressable or numerics demand a fused kernel. Either way the 3 regime Qs come from 3 `ggml_rope_ext` on remapped `inp_pos` views (intra=local [0,c); inter=chunk-index clamp; succ=continuous bridge), remapping only the rotary 128 dims (partial_rotary 0.25); the other 384 pass through.

**Build path is cosmocc, NOT cmake** (confirmed: patch sources `ggml-neo-pipeline/iqk/turbo/d512-vec` are NOT in any `CMakeLists.txt`; native `cmake --build` link-fails on those symbols). Clean-room = out-of-tree copy of the held `vendors/sources/llamafile` tree (already has 0001-0077 applied to source); edit C1 there; `bun run build:llamafile:make` (host binary, reuses fat DSO); gate E with `--dca on` + `--override-kv gemma4.context_length` serve; capture 0078 by `diff` vs a frozen copy of the held tree. Held submodule never reset.

### Section-C VERDICT (2026-06-10, op contract verified read-only) — C1 FALSIFIED → C2 required

Read the actual 0070 op definitions (`vendors/patches/llamafile/0070-rolling-kv.patch`). The "implement
C1 first" decision above is **falsified by the op contract**:

1. The GPU op `ggml_streaming_flash_attn` returns **normalized O only** — `ne[4]={DV, n_head, n_q,
   n_batch}`, F32 (patch L539). Its per-tile LSE lives in **internal pool allocs** `O_tiles`/`lse_tiles`
   (L1327-8) consumed by a **`static __global__ streaming_combine_kernel`** *inside* fattn.cu (L1113,
   invoked L1859). **LSE is never a graph tensor.** → cannot call it 3× and merge (O,LSE) host-side.
2. The only op that emits the addressable un-normalized partial `[2+DV,n_head,n_q,n_b]`=(M,S,VKQ),
   `ggml_flash_attn_ext_tail_partial`, is **CPU-only** (dispatched solely in ggml-cpu/ops.cpp L694/811;
   **no CUDA case** — CUDA dispatch L1895/1912 is `STREAMING_FLASH_ATTN` only) **and decode-only `n_q==1`**
   (its own header, L618-620). Useless for the 512k-token **prefill** where DCA matters; CPU scalar at
   512k is intractable regardless.
3. Eager/materialized DCA (3 masked QK^T blocks combined by pure graph math, single softmax — which
   *would* be genuine C1, no DSO rebuild) needs the full O(n²) score matrix: ≈1 TB/head @512k. Infeasible.

**Therefore: the performant DCA path requires a CUDA-source edit → DSO rebuild (C2).** There is no
host-only route to per-regime (O,LSE) at prefill scale. Build flips from cheap `build:llamafile:make`
(reuse fat DSO) to **`cuda.sh` ggml-cuda.so rebuild** (the multi-hour cosmocc CUDA path), done in the
out-of-tree clean-room (held submodule still never touched).

**C2, minimal-footprint shape (reuses machinery that already exists internally):** expose the existing
per-tile partial + combine as graph-addressable —
- (a) a GPU op variant that runs FA over one regime and writes the **un-normalized partial** (M,S,VKQ)
  to a graph tensor (the streaming op already computes exactly this per tile before its internal
  combine — emit it instead of/in addition to normalized O), and
- (b) a host-exposed N-way `streaming_combine_n(partials[3])` generalizing `streaming_combine_kernel`
  (which already combines N tile-partials) so the 3 regime partials merge on-device.
Per UPSTREAM_SYNC, new CUDA code lands in a NEW file; the in-tree hook stays a single marker. The 3
regime Qs still come from 3 `ggml_rope_ext` on remapped `inp_pos` views (intra=local [0,c);
inter=chunk-index clamp; succ=continuous bridge), rotary-128 dims only (partial_rotary 0.25); other
384 pass through. Numerics validated against the existing op's internal combine (same LSE math).

## MULTI-ARCH generalization (2026-06-10, user directive: "don't limit DCA to Gemma4 — same 1M extension on Qwen")

DCA is now an **architecture-general** training-free 1M extension, not a Gemma4 one-off. Patch renamed
`0078-dca` (was `0078-gemma4-dca`). Verified the Qwen seam is identical to Gemma4's:
`qwen2.cpp`: `ggml_rope_ext(Q/K, inp_pos, freq_factors=nullptr, n_rot=n_embd_head /*full rotary*/,
rope_type, n_ctx_orig, freq_base, freq_scale, …)` → `build_attn(inp_attn, wo, …, Q,K,V, …, scale, il)`
— same shape, differing only in `freq_factors` (Gemma4 global = `rope_freqs`; Qwen = null) and
`n_rot` (Gemma4 partial-128 of 512; Qwen full head_dim). So **one generic helper serves all**:

- **Section C (core, arch-agnostic)** — new `dca.{cpp,h}`: `llm_graph_input_dca` (the 3-regime remapped
  position vectors, built once like `inp_pos`) + `llm_graph_context::build_attn_dca(inp_attn, inp_dca,
  wo, wo_b, Qcur_PREROPE, Kcur_PREROPE, Vcur, …, dca_params, rope_params, kq_scale, il)`. It applies the
  3 regime ropes (intra/successive/inter) via `ggml_rope_ext` on remapped positions, runs the
  partial-emit FA per regime, merges with `streaming_combine_n`, then the `wo` projection (mirrors
  `build_attn`'s tail). `rope_params` carries exactly the args each builder already passes to
  `ggml_rope_ext` (n_rot, rope_type, n_ctx_orig, freq_base/scale, ext/attn_factor, beta_fast/slow,
  freq_factors) → handles Gemma4 partial-rotary AND Qwen full-rotary with no special-casing.
- **Section B (per-arch hooks, small)** — in each builder, gate the full-attention layers through
  `build_attn_dca` instead of `rope_ext+build_attn`: Gemma4 `is_dca_layer = !is_swa(il)` (5 global only);
  Qwen `is_dca_layer = true` (all layers, no SWA). First-landing hooks: `gemma4-iswa.cpp`, `qwen2.cpp`
  (Qwen2.5-1M, the canonical DCA target), `qwen3.cpp`, `qwen3moe.cpp`; `qwen35/qwen35moe/qwen2moe` are
  one-marker each, added as needed. Marker `// opencoti F5 dca`.
- **Section D (CUDA, arch-agnostic)** — partial-emitter op + `streaming_combine_n`; operates on tensors,
  no arch knowledge. This is the C2 DSO-rebuild payload.
- **Gate E** — RULER-VT @512k on BOTH A4B Q4_K_M AND a Qwen 1M model, vs the naive-ext bar (VT 0.36).

### Clean-room standup (confirmed mechanics)
Build pipeline: `bun --cwd packages/opencoti-llamafile script/build-pipeline.ts {reset,apply,make,cuda}`.
`apply` = `reset`(`git clean -fdx`+checkout pin `6490e16`) → `bootstrap`(download cosmocc + Mozilla-Ocho
`make setup`) → `git apply` 0001-0077. The held submodule has 8 uncommitted WIP files (`cuda.sh`,
`llamafile.c/h`, build mk, nested `llama.cpp`) — `reset` would WIPE them (buglog bug-431). So the
clean-room is a **separate `git clone /shared/dev/opencoti <dca-clone>`** (committed state only; held WIP
excluded by design), where `apply`→edit 0078→`cuda` runs freely. Capture 0078 by diffing the clone's
post-0078 source vs its post-0077 (pre-edit) source. Held tree never reset, never touched.

## Section-C ENGINEERING DESIGN (2026-06-10, grounded in clone applied source) — the cache-rope problem

Reading `llm_graph_context::build_attn` (kv variant) exposed the crux: **K is stored to the KV cache
already rope'd** (model builder ropes Kcur → `build_attn` → `mctx_cur->cpy_k(Kcur)`), and attention
reads it back rope'd. DCA needs the SAME keys re-rope'd with DIFFERENT positions per regime
(intra=native pos, inter=chunk-clamped, succ=bridge) — a single cached rope can't provide that.

**Decision: DCA layers cache K UN-ROPE'd and apply all rope at attention time, per regime, inside
`build_attn_dca`.** Consequences:
- The per-arch hook (Section B) passes **pre-rope Q AND pre-rope K** to `build_attn_dca` (not just
  re-routes the call). On Gemma4 global, pre-rope K = `attn_k_norm(K_proj)` (the tensor right before
  `ggml_rope_ext`); V = `rms_norm(K_proj)` is already un-rope'd (k_eq_v) and unchanged. On Qwen,
  pre-rope K = the reshaped projection before its `ggml_rope_ext`.
- `build_attn_dca` stores pre-rope K + V to cache (`cpy_k`/`cpy_v`), then per regime r∈{intra,succ,
  inter}: `Q_r=rope(Qcur_pre, pos_q_r)`, `K_r=rope(get_k_raw, pos_k_r)`, `partial_r =
  streaming_flash_attn_PARTIAL(Q_r, K_r, V, mask_r, kq_scale)`; then `O =
  streaming_combine_n([partial_intra, partial_succ, partial_inter])`; then `wo` projection (mirrors
  build_attn tail: `build_lora_mm(wo,·)` + optional `wo_b`).
- Rope-on-cached-raw-K is an ESTABLISHED pattern here: `llama_kv_cache::build_rope_shift`
  (llama-kv-cache.cpp:3893, used by context-shift) re-ropes cached K via a position/`k_rot` tensor —
  the structural template (DCA ropes a READ COPY with absolute regime positions, does NOT mutate the
  cache).
- **`llm_graph_input_dca`** (built once like `inp_pos`, host-filled from batch positions + KV-cache
  cell positions) supplies, per regime: `pos_q_r` (len n_tokens), `pos_k_r` (len n_kv), and `mask_r`
  ([n_kv × n_tokens] additive). The pos remaps + masks encode the DCA chunk math
  (chunk_size c = `--dca-chunk-size`, default ← n_ctx_orig). intra = same-chunk causal, native pos;
  inter = q-chunk > k-chunk, k pos within-chunk + q pos clamped ≤ pretrain window; succ = adjacent
  chunk bridge. rope uses `rope_params` (n_rot, rope_type, freq_base/scale, ext/attn_factor,
  beta_fast/slow, freq_factors) — Gemma4 partial-128 / Qwen full-rotary handled by the n_rot arg.

**Toggle is launch-time only** (cache contents differ: raw-K for DCA vs rope'd-K for non-DCA) — fine,
it's a serve flag. Wiring caveat: Gemma4 0076 shared-KV means some global layers `reuse KV cache of
earlier layers` (`has_kv(il)` false) — DCA must rope-on-read consistently for the layer-group that
owns the cache; verify against the 0076 shared-kv mechanics when wiring Section B.

## Implementation progress (2026-06-10) + Section-D simplification

Patch number **0078 confirmed** (user, keep 0078). Build host = solidpc RTX 3090 (sm_86, CUDA 13.2).

**Section-D shrinks to ONE new op + graph-math combine** (no new combine kernel). Insight: if the new
FA op emits NORMALIZED O packed with per-row LSE as `[DV+1]` (reusing 0070's existing `dst_lse`
channel — the `lse = m + logf(sum exp)` math already lives in `launch_fattn`), the 3-regime merge is
pure existing-ggml graph math: `M = max_r(lse_r)` (binary max via `a + relu(b−a)`), `w_r =
exp(lse_r−M)`, `O = Σ_r w_r·O_r / Σ_r w_r` (`ggml_exp/mul/add/div`, broadcasting w over DV). So the
DSO rebuild only adds a CUDA *forward* for the new op (FA + write lse), not a bespoke combine kernel.

**DONE this session (clean-room `/shared/dev/opencoti-dca`, baseline nested `43ccc8a56`):**
- **Section A (flags)** — `--dca`/`--dca-chunk-size`/`--dca-yarn-factor` through all 6 opencoti sites
  (arg.cpp, common.h, common.cpp, llama-context.cpp ×2, llama-cparams.h, llama.h). Verified.
  Script: `backup_models/scripts/dca/apply_section_a_flags.py`.
- **Section D-host** — `GGML_OP_FLASH_ATTN_EXT_LSE` + `ggml_flash_attn_ext_lse(q,k,v,mask,scale,
  max_bias,logit_softcap)` → dst `[DV+1, n_head, n_q, n_batch]` ([0,DV)=O, [DV]=lse). Registered in
  ggml.h enum+decl, ggml.c name/symbol arrays, both static_asserts (100→101), constructor.
  Script: `backup_models/scripts/dca/apply_section_d_host.py`.

**REMAINING:** D-cuda (CUDA forward for the new op: `ggml_cuda_flash_attn_ext` path + dst_lse→dst[DV]
slot; supports_op + scheduler-pin like the bug-259 STREAMING_FLASH_ATTN hook) → Section C
(dca.cpp/h: `llm_graph_input_dca` + `build_attn_dca`) → Section B (per-arch hooks) → cuda.sh DSO
rebuild → RULER-VT gate (A4B + Qwen) → capture `0078-dca.patch`.

## Section D-cuda DONE (2026-06-10)

`ggml_cuda_flash_attn_ext_lse` authored + wired (8 sites, `apply_section_d_cuda.py`):
- **fattn.cu** — the forward: O via an O-typed view of the packed dst → existing
  `ggml_cuda_flash_attn_ext` (proven streaming-op `ggml_tensor o=*dst; o.op=FLASH_ATTN_EXT` trick);
  per-row lse via the existing static `streaming_lse_kernel` into the trailing `n_rows` (softcap-aware,
  works for any n_q incl. prefill — avoids the `decode_lse`/bug-263 finalize fragility). Asserts Q f32 / K f16.
- **fattn.cuh** decl; **ggml-cuda.cu** compute_forward dispatch + supports_op (validates the O FA at
  ne[0]=DV, not the packed DV+1); **ggml-backend.cpp** bug-259 scheduler CUDA-pin; **ggml-cpu** ×3
  (compute_forward abort, n_tasks single-task, supports_op=false) — CUDA-only mirror of STREAMING_FLASH_ATTN.

Op fully complete (host + CUDA). Compile-verified later in the single cuda.sh DSO build (after C+B).
**REMAINING:** Section C (dca.cpp/h: `llm_graph_input_dca` + `build_attn_dca` with DCA position-remap +
graph-math combine — O view ne=[DV,H,nq,nb] @offset0, lse view ne=[1,H,nq,nb] @offset DV*n_rows) →
Section B (per-arch hooks) → cuda.sh DSO rebuild → RULER-VT gate → capture 0078-dca.patch.

## Section C design (2026-06-10) — grounded against the clean-room seams

Read of the clean-room (`43ccc8a56`) settles the algorithm and the integration points.

**Seams (verified):**
- kv-variant `build_attn` (llama-graph.cpp:2571) and **iswa**-variant `build_attn`
  (llama-graph.cpp:2841). Both: `cpy_k/cpy_v(k_cur,…)` store **whatever K they're handed** →
  `get_k/get_v(ctx0,il)` read-back → `build_attn_mha` → `build_lora_mm(wo)` + `wo_b`. Gemma-4
  globals ride the **iswa** path (`build_attn_inp_kv_iswa`, `is_swa(il)==false` → base mctx);
  Qwen rides the **plain-kv** path. ⇒ `build_attn_dca` needs **both** input-type overloads.
- Host position source: `set_input_kq_mask` (llama-kv-cache.cpp:3773) / `set_input_pos_bucket`
  (:3812) read `v_cells[strm].pos_get(j)` / `cells.is_empty(j)` per cached cell. A
  `llm_graph_input_i::set_input` only gets the ubatch → the per-cell fill **must** be a new
  `llama_kv_cache[_context]::set_input_dca` (mirrors `set_input_kq_mask`), delegated to from
  `llm_graph_input_dca::set_input`.
- Rope contract (gemma4-iswa.cpp:73/97): `ggml_rope_ext(x, inp_pos, freq_factors, n_rot_l,
  rope_type, n_ctx_orig, freq_base_l, freq_scale_l, ext, attn, beta_fast, beta_slow)` — Gemma-4
  globals use `n_rot_l=hparams.n_rot(il)` (partial-128 on the hd512 global), `freq_base_l=1e6`.
  Qwen full-rotary uses `n_rot = n_embd_head`, `freq_base=1e6`/yarn. `build_attn_dca` reproduces
  this call three (Q) + one (K) times with regime position vectors instead of `inp_pos`.

**Algorithm — rope-K-ONCE, rope-Q-THRICE (ChunkLlama scheme).** All three regimes share one key
assignment `pos_k[j] = cellpos(j) mod c`; only the query position differs:

| regime | `pos_q[i]` | keep-mask (causal `j≤i`, chunk `=⌊pos/c⌋`) |
|---|---|---|
| **intra** | `i mod c` | `chunk(j) == chunk(i)` |
| **succ**  | `(i mod c) + c` | `chunk(j) == chunk(i) − 1` |
| **inter** | `c` (const) | `chunk(j) <  chunk(i) − 1` |

The three keep-masks **partition the causal lower-triangle exactly** → the LSE-merge of the three
softmaxes equals full attention under DCA-remapped positions (no double-count, no gap). No q-k
distance exceeds the pretrain window: intra `<c`, succ `∈(0,2c)`, inter `=c−ok∈[1,c]`. `c` =
`dca_chunk_size` or (auto) `n_ctx_orig_yarn`. YaRN mscale composes via the existing rope `ext/attn`
factors (`dca_yarn_factor` feeds `cparams.yarn_*`; the rope kernel itself is unchanged).

**`build_attn_dca(inp, wo, wo_b, Qpre, Kpre, Vcur, …, kq_scale, il)` graph body:**
1. `cpy_k(Kpre)` / `cpy_v(Vcur)` to cache (Section B hands **pre-rope** K + un-roped V).
2. `Kraw = get_k(ctx0,il)`; `V = get_v(ctx0,il)`; `Kdca = ggml_rope_ext(Kraw, inp_dca->pos_k, …)`
   — **one** rope of the large cached K (raw cells, regime-invariant key positions).
3. For `r ∈ {intra,succ,inter}`: `Qr = ggml_rope_ext(Qpre, inp_dca->pos_q[r], …)`;
   `Pr = ggml_flash_attn_ext_lse(Qr, Kdca, V, inp_dca->mask[r], kq_scale, max_bias, softcap)`
   → packed `[DV+1, H, nq, nb]`. O view `ggml_view_4d(Pr, DV,H,nq,nb,…,0)`; lse view
   `ggml_view_4d(Pr, 1,H,nq,nb,…, DV*nb0)`.
4. Combine (pure ggml): `M = a+relu(b−a)` over the 3 lse; `wr = ggml_exp(lse_r − M)` (broadcast over
   DV); `O = (Σ wr·Or) / (Σ wr)`; reshape `[n_embd, n_tokens]`; `build_lora_mm(wo)` + `wo_b`.

**`llm_graph_input_dca`** (new, dca.h): `pos_k` (I32 [n_kv]), `pos_q[3]` (I32 [n_tokens]), `mask[3]`
(F32→F16 cnv, `[n_kv, GGML_PAD(n_tps,GGML_KQ_MASK_PAD), 1, n_stream]` like `build_attn_inp_kq_mask`),
`const llama_kv_cache_context * mctx`, `uint32_t c`. `set_input(ubatch)` →
`mctx->set_input_dca(pos_k, pos_q, mask, ubatch, c)`.

**Files touched by Section C:** new `src/dca.cpp` + `src/dca.h` (the two builders) **and**
`llama-kv-cache.{cpp,h}` (the `set_input_dca` host-fill — folded into C) **and** `llama-graph.h`
(the `build_attn_dca` decls + `llm_graph_input_dca`). dca.cpp compiles against the already-present
host `ggml_flash_attn_ext_lse` (D-host); the CUDA forward (D-cuda) need not exist until the DSO
build. **Sub-blocks:** C1 = dca.h + `llm_graph_input_dca` + `set_input_dca` host-fill (this block);
C2 = `build_attn_dca` graph bodies (both overloads).

## Section C2a (host-fill) DONE (2026-06-10)

`dca.h` (C1) + the host-fill half of C2 landed in the clean-room. Builds independently of the
`build_attn_dca` graph bodies (next block) and of D-cuda (host `ggml_flash_attn_ext_lse` already
present from D-host).
- **`src/dca.cpp`** (new) — `dca_resolve_chunk_size` (cparams.dca_chunk_size or auto =
  n_ctx_orig_yarn / n_ctx_train), `llm_graph_input_dca::set_input` (delegates to
  `mctx->set_input_dca`), `::can_reuse` (false — pos/masks are position-specific), and
  `llm_graph_context::build_attn_inp_dca(mctx_kv)` (allocates pos_k I32[n_kv], 3× pos_q I32[n_tokens],
  3× mask F32[n_kv,n_tokens,1,1] + F16 cnv; single-stream, mirrors `build_attn_inp_kq_mask`).
- **`src/llama-kv-cache.cpp`** — `llama_kv_cache::set_input_dca` (reads `v_cells[0]` per-cell
  positions → `pos_k=cellpos%c`, 3× pos_q, and the 3 partition masks; `n_stream==1` assert like
  `set_input_pos_bucket`; causal + same-seq + nonempty band test) + the `llama_kv_cache_context`
  per-batch wrapper.
- **`src/llama-kv-cache.h`** — `set_input_dca` decl in both classes (×2).
- **`src/llama-graph.h`** — `build_attn_inp_dca` method decl + `class llm_graph_input_dca;` fwd-decl.

Apply script `backup_models/scripts/dca/apply_section_c_hostfill.py` (5 anchored edits, all OK).
Verified: kv-cache.cpp ×2, kv-cache.h ×2, graph.h ×2, dca.cpp def ×1, braces balanced.

**⚠ Build-wiring TODO for the cuda.sh/make block (#598):** confirm the cosmocc Makefile globs
`src/*.cpp` (so the new `src/dca.cpp` compiles) — else add it to the file list explicitly, or
`build_attn_inp_dca` link-errors.

**REMAINING:** Section C2b (`build_attn_dca` two overloads — rope cached raw-K once with pos_k +
rope Q thrice + 3× `ggml_flash_attn_ext_lse` under the partition masks + LSE graph-math combine +
wo/wo_b; declared in llama-graph.h alongside this block's decls). **Before authoring C2b, read
`get_k`'s returned cache-K layout (llama-kv-cache.cpp:2204) + `build_rope_shift`'s caller
(~:4000-4030)** to nail the rope-on-cached-K axis (positions vary on the cell axis). → Section B
(per-arch hooks) → cuda.sh DSO rebuild → RULER-VT gate → capture 0078-dca.patch.

## Section C2b (build_attn_dca bodies) DONE (2026-06-10) — Section C COMPLETE

The rope-on-cached-K axis is settled by `get_k` (llama-kv-cache.cpp:2305): it returns
`[n_embd_head, n_head_kv, n_kv, 1]` — **cells on ne[2]** — so `ggml_rope_ext(get_k, pos_k, …)` with
`pos_k` length `n_kv` indexes the cell axis exactly as `inp_pos` indexes live tokens. K is roped
ONCE; Q thrice; both pre-permute (matching the model rope), then permuted into the flash-attn layout.

- **`src/dca.cpp`** — `build_attn_dca_core` (rope cached raw-K once; read+permute V once; per regime
  rope Q + `ggml_flash_attn_ext_lse` under `mask_cnv[r]`; extract O-view `[DV,H,NQ,1]@0` + lse-view
  `[1,H,NQ,1]@DV·nb0`, both `ggml_cont`; combine `M=a+relu(b−a)`, `w=ggml_exp(lse−M)`, `O=Σw·O/Σw`;
  `reshape_2d`→`build_lora_mm(wo)`+`wo_b`) + the two public overloads (iswa: store to
  `inp->mctx->get_base()`, asserts `!is_swa`; plain-kv: `inp->mctx`; both guard `k_cur`/`v_cur` for
  Gemma-4 shared-KV layers).
- **`src/dca.h`** — `dca_rope` struct (per-layer freq_base/freq_scale/n_rot/freq_factors).
- **`src/llama-graph.h`** — 2 `build_attn_dca` overloads + `build_attn_dca_core` decls + `dca_rope`
  fwd-decl. Apply script `apply_section_c2b_decls.py`.

**Empty-regime safety (verified, no code needed):** a chunk-0 query's succ/inter regimes have an
all-`-inf` mask row → FA O=0 and `streaming_lse_kernel` (fattn.cu:693 `lse = (l>0)?…:-INFINITY`)
emits `-inf` (the `l<=0`/NaN guard already there) → `w=exp(-inf−M)=0`, dropping the regime cleanly.
Intra is never empty (diagonal), so `M` is always finite.

**Two runtime caveats for the gate (#598):**
1. ~~**f16 KV required**~~ — **LIFTED by #444 (DCA all-KV C1, 2026-06-18).** `build_attn_dca_core` no
   longer asserts f16; `dca_lift_to_f16` dequant-on-lifts a quantized/bf16/turbo cache to a transient
   f16 for the f16-only fused kernel (the stored cache stays quantized — the long-ctx memory win). Scalar
   quants (q8_0/q4_0/q4_1/q5_0/q5_1) hop `q8→f32→f16` because CUDA CPY has no direct scalar-quant→f16
   kernel (that pair CPU-spills); f16 no-op, bf16/turbo cast direct. Host-only (no `.cu` change). Gated
   `real_frac=0` vs f16-DCA on Qwen3-8B-Q8 at multi-chunk ctx (`.opencoti/m444-dca-allkv-gate.sh`). **#445
   (C2/C3) extends the gate to the full matrix — q8_0, q5_1, q5_0, q4_0, bf16 + asymmetric q8_0-K/q4_0-V
   ALL PASS** (`.opencoti/m445-dca-scalar-matrix-gate.sh`): per type `real_frac(<t>-DCA vs f16-DCA) ==
   real_frac(<t>-noDCA) == 0.0`, so DCA adds zero argmax error beyond the type's own quant loss. (Prefill
   note: noDCA-quant FA prefill is type-specific — q4_0/bf16 have a fast batched tile-FA kernel ~3.5k
   tok/s, q8_0/q5_0/q5_1 fall to the decode-VEC path ~12–42 tok/s; DCA's f16-lift prefills uniformly
   ~1.5k tok/s — a win for the slow types.) q8-DCA
   decode ≈0.47× f16-DCA, and that 2× is characterized (not a regression to chase away): it is
   GPU-resident (CPU-spill would be <5 tps) and NOT cuda-graph loss (`GGML_CUDA_DISABLE_GRAPHS=1`
   changes neither path — the fused DCA op uses no graph). The cost is the dequant-on-lift itself,
   DOMINATED by the `quant→f32` CPY kernel's `<<<ne, 1>>>` one-thread-per-CUDA-block launch geometry
   (cpy.cu, ~1/32 warp utilization vs the 256-thread scalar cpy), with the f32 two-pass hop secondary;
   raw cast bandwidth alone (~3 ms/tok) does not explain it. The single-pass direct `quant→f16` CPY
   kernel (proper thread geometry, no f32 intermediate) recovers it — C-series perf follow-on at #448.
   So `-ctk q8_0 -ctv q8_0` is now valid with `--dca on`. The 0078 patch re-capture is batched at #448
   (umbrella C6); the live source is banked `.opencoti/dca/m444-dca.cpp.banked`.
2. **FA precision defaults** — unlike `build_attn_mha` (which forces `GGML_PREC_F32` via
   `ggml_flash_attn_ext_set_prec`, an op-specific call that would assert on `_LSE`), the lse op runs
   at default prec. If RULER shows precision drift at 1M, force F32 on the O-view in D-cuda's
   `ggml_cuda_flash_attn_ext_lse` (set the prec slot before dispatching the inner FA).

Verified: dca.cpp 3 defs + braces 26/26; graph.h 3 decls + `dca_rope` fwd; ruff clean.

**Section C COMPLETE** (C1 header + C2a host-fill + C2b bodies). Compiles against host
`ggml_flash_attn_ext_lse` (D-host); CUDA forward (D-cuda) needed only at the DSO build.
**REMAINING:** Section B (per-arch hooks: gemma4-iswa global layers + qwen2/3/3moe all layers — gate
on `cparams.dca_enabled`, pass PRE-rope Q+K, build `build_attn_inp_dca(base)` once, fill `dca_rope`)
→ cuda.sh DSO rebuild (confirm `src/dca.cpp` globbed) → RULER-VT gate → capture 0078-dca.patch.

## Section B (gemma4) DONE (2026-06-10)

`src/models/gemma4-iswa.cpp` now routes the 5 GLOBAL (`!is_swa`) layers through `build_attn_dca`
when `cparams.dca_enabled`; SWA layers + the dca-disabled path are byte-unchanged. 7 anchored edits
(`apply_section_b_gemma.py`):
- `#include "dca.h"`.
- before the layer loop: `llm_graph_input_dca * inp_dca = cparams.dca_enabled ?
  build_attn_inp_dca(inp_attn->mctx->get_base()) : nullptr;` (built once on the base/global context).
- per layer: `const bool use_dca = cparams.dca_enabled && !hparams.is_swa(il);`.
- Q-rope + K-rope wrapped in `if (!use_dca) {…}` — DCA layers cache **PRE-rope** Q+K (rope happens
  per regime inside build_attn_dca). Vcur stays the rms-normed un-roped projection.
- both `build_attn(…)` call sites (has_kv + shared-KV reuse-KV branch) gain an `if (use_dca)`
  arm calling `build_attn_dca(inp_attn, inp_dca, wo, nullptr, Qcur, [Kcur|null], [Vcur|null],
  {freq_base_l, freq_scale_l, n_rot_l, freq_factors}, f_attention_scale, il)`.

DCA reduces to identity for n_ctx ≤ c (all positions in chunk 0 → intra-only, succ/inter empty,
distances < c → native rope), so enabling `--dca` never regresses short-context serving. Verified:
2 build_attn_dca calls, 1 inp builder, use_dca 5×, braces 27/27, ruff clean.

**REMAINING:** Section B (qwen2/qwen3/qwen3moe — all layers full-attention, `use_dca =
cparams.dca_enabled`, same PRE-rope Q+K + `build_attn_inp_dca(static_cast<…kv_cache_context*>(mctx))`
on the plain-kv overload) → cuda.sh DSO rebuild (confirm `src/dca.cpp` globbed) → RULER-VT gate
(Gemma-4 A4B first, f16 KV) → capture 0078-dca.patch.

## Section B (qwen) DONE (2026-06-10) — Section B COMPLETE → ALL SOURCE AUTHORED

`qwen2.cpp` / `qwen3.cpp` / `qwen3moe.cpp` route **all** layers through `build_attn_dca` when
`cparams.dca_enabled` (Qwen = the canonical training-free 1M arch; every layer is full attention).
The Q/K-rope + `build_attn` blocks are byte-identical across the three files, so one shared anchor
set (`apply_section_b_qwen.py`, 5 edits ×3) applies to each: `#include "dca.h"`; `use_dca =
cparams.dca_enabled` + `inp_dca = build_attn_inp_dca(inp_attn->mctx)` (plain-kv context, no
get_base); Q/K rope wrapped in `if (!use_dca)`; the `build_attn` site gains an `if (use_dca)` arm
calling `build_attn_dca(..., model.layers[il].wo, model.layers[il].bo, Qcur, Kcur, Vcur,
{freq_base, freq_scale, n_rot, nullptr}, 1/√head, il)` (wo_b = bo carried; qwen3/3moe `wo_s`
post-scale untouched). Verified: each file 1 dca-call, 1 inp builder, use_dca 5×, braces balanced,
ruff clean.

**ALL 0078-dca SOURCE AUTHORED** — A (flags) + D (LSE op host+CUDA) + C (input + build_attn_dca) +
B (gemma4 globals + qwen2/3/3moe all layers). Nothing compiled yet — the single `cuda.sh` DSO build
(#598) compiles A+C+D+B together. New files: `src/dca.{h,cpp}`. Apply scripts (8) in
`backup_models/scripts/dca/`.

**REMAINING (#598, the build+gate+land block):**
1. **Build wiring** — confirm the cosmocc Makefile globs `src/*.cpp` (so `src/dca.cpp` compiles);
   add it explicitly if not. Then `bun --cwd packages/opencoti-llamafile script/build-pipeline.ts
   make` + `… cuda` (the multi-hour `ggml-cuda.so` DSO rebuild, sm_86). Fix any compile errors
   surfaced (first real compile of A/C/D/B).
2. **RULER-VT gate** — serve A4B Q4_K_M (or bf16) with `--dca on --override-kv
   gemma4.context_length=int:1048576`, **f16 KV** (NOT `-ctk q8_0` — DCA core asserts f16), default
   prec. RULER-VT @ ≥256k vs the naive-extension bar (VT=0.36 from T87.ext). Then a Qwen 1M model.
3. **Capture** — `git diff 43ccc8a56` in the clean-room → `0078-dca.patch`; write patch + README row
   + UPSTREAM_SYNC entry + `.opencoti/dca-ruler-gate.sh` into the HELD opencoti repo (top-level
   tracked only; submodule never touched); round-trip `apply` verify.