docs/features/gemma4_dca.md

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/cleaned (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.