gqa(adr): review fixes — unify KV placement to contiguous, ring/reduce cost model, opt3 wording, rank def

Per review:
- Placement unified to *contiguous* C×P blocks throughout (S0/S2.1/S2.2/S0.5);
  round-robin demoted to the rejected alternative. Adds driver SP-enable
  threshold fallback (smaller C / C=P=1 for short/early decode).
- Ring-vs-reduce cost model in S5.5: reduce ~ G*T_q*log(C*P) (O dominant;
  m,l scalars), ring ~ 2*S (total K+V bytes a CUBE injects over C-1
  rotations; recv_async pipelines so latency ~ max-step) -> ring wins when
  T_q > 2S/(G*log(C*P)).
- opt3 'removes the bubble' -> 'hides (subject to scheduler+engine balance)'
  everywhere; table 'hidden*' with footnote.
- 'rank' defined (SP participant = a PE in a CUBE, KV shard in its HBM->TCM).
- out-proj handoff contract (S0.5.4); S11 gate-type note (absolute latency
  deferred to ADR-0064; structural + relative-to-baseline gates now).
- greenlet-as-contrast tightened to 'primitive-op (tl.dot) path' (S1, SB).
KO mirror synced; DDD gets the SP-enable threshold fallback. Docs only.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

docs/adr-proposed/ADR-0060-algo-gqa-fused-attention-ahbm.md

@@ -82,7 +82,7 @@ def gqa_decode_v1(q_ptr, k_ptr, v_ptr, o_ptr, S_kv_local, d, C, P, scale, *, tl)
 ```
 
 ```python
-# OPTION 3 — software pipelining (current primitives; bubble removed) ← SHIP THIS
+# OPTION 3 — software pipelining (current primitives; bubble hidden, not provably removed) ← SHIP THIS
 def gqa_decode_v3(q_ptr, k_ptr, v_ptr, o_ptr, S_kv_local, d, C, P, scale, *, tl):
     cube_id = tl.program_id(axis=1); pe_id = tl.program_id(axis=0)
     kv  = head_of_group(cube_id); q_g = tl.load(q_base(kv), (G * 1, d))
@@ -118,8 +118,12 @@ def gqa_decode_v2(q_ptr, k_ptr, v_ptr, o_ptr, S_kv_local, d, C, P, scale, *, tl)
 | | new HW cmd | GEMM bubble | CPU intra-tile wait | issues | now |
 |---|---|---|---|---|---|
 | **opt1 current** | no | **yes** | yes | `O(tiles·ops)` | ✓ |
-| **opt3 sw-pipe** | no | no | yes (reordered) | `O(tiles·ops)` | ✓ |
-| **opt2 ex_composite** | **#2 only** | no | **no** | `O(tiles)` | ✗ (build #2) |
+| **opt3 sw-pipe** | no | hidden\* | yes (reordered) | `O(tiles·ops)` | ✓ |
+| **opt2 ex_composite** | **#2 only** | hidden\* | **no** | `O(tiles)` | ✗ (build #2) |
+
+\* "hidden" not "removed": filling the GEMM engine during softmax depends on
+the scheduler/engine balance (softmax-vs-GEMM time, K DMA readiness, queue
+depth) — it *reduces/hides* the bubble, not provably eliminates it.
 
 (`#1` = the *existing* GEMM composite + `scale`; only `#2` = softmax + P·V +
 the stateful online-softmax accumulator is new. MATH engine already has
@@ -257,9 +261,12 @@ two levels** of sequence-parallelism (SP):
 - **Level-2 (intra-CUBE, across PEs):** each CUBE's slice is split again
   across its `P` PEs.
 
-So one KV head maps to **`C × P` ranks**, all within one SIP. **How the
-`G` query heads map onto those ranks differs by case** (the two kernels,
-TL;DR / §5):
+So one KV head maps to **`C × P` ranks**, all within one SIP. A **rank** is
+not memory — it is one **SP participant** = a specific **PE in a specific
+CUBE** (`rank = cube_local·P + pe_local`); its KV shard *resides in* that
+PE's owned HBM region (`K_base[rank]`) and is DMA'd to the PE's TCM for
+compute. **How the `G` query heads map onto those ranks differs by case**
+(the two kernels, TL;DR / §5):
 
 - **Decode** (§4): **Q is replicated** — every rank holds all `G` query
   heads, **M-folded** (stacked into the matmul M / row dimension: `Q` for
@@ -290,10 +297,14 @@ drives its 2 CUBE Groups (2 KV heads) — the head is selected spatially by
 CUBE coordinate, **not** an in-kernel `for kv` loop. The CLI runs the
 4 SIP-devices logically in parallel to cover all 8 heads.
 
-**Token → rank placement** within a CUBE Group (round-robin SP, balanced
-to ≤1 token): KV token `i` lands on Level-1 rank `(start + i) mod C`, then
-Level-2 rank `((i // C) + start_pe) mod P` — a hierarchical round-robin so
-each of the `C × P` ranks owns a dense, contiguous local slice.
+**Token → rank placement** within a CUBE Group is **contiguous block
+sharding**: rank `r` owns the position block `[r·B, (r+1)·B)`,
+`B = ⌈S/(C·P)⌉` (Level-1 picks the CUBE, Level-2 the PE within it). A
+contiguous block is **required** so prefill can causal-skip whole future
+blocks **and** so prefill + decode share one KV cache (§2.1). (Round-robin —
+interleaved `token i → rank (start+i) mod (C·P)` — gives better
+short-context decode balance but cannot causal-skip and cannot be shared
+with prefill; it is the **rejected** alternative, §2.1.)
 
 Reduction is **hierarchical and intra-SIP**: a KV head **does** span the
 `C` CUBEs of its group (Level-1), reduced over the CUBE NOC; this
@@ -326,8 +337,8 @@ existed only because the `H_kv=1` baseline never needed it. Crossing
   launches.** ⇒ this kernel is **pure read** on the KV cache.
 - **P4.** V is not rotated; V-cache holds raw projected V.
 - **P5.** RoPE position upstream is the token's **absolute global
-  position** `global_idx = local_slot·(C·P) + rank` where `rank =
-  cube_local·P + pe_local` (§2.1). Round-robin placement ≠ RoPE position.
+  position** `global_idx = rank·B + local_slot` (`B = ⌈S/(C·P)⌉`, contiguous
+  block, `rank = cube_local·P + pe_local`, §2.1). Placement ≠ RoPE position.
 
 ### 0.5.2 Shape symbols
 `T_q` = query length this launch (decode: 1; prefill/chunk: chunk width).
@@ -341,7 +352,7 @@ matters.
 | Input | Shape (per KV head) | Location | Notes |
 |---|---|---|---|
 | `Q` | `[G, T_q, d]` | per-rank HBM (loaded to TCM) | post-RoPE (P1). `G` query rows of the group batched. |
-| `K_cache` | `[S/(C·P), d]` | per-rank HBM, base `K_base[rank]`, contiguous | post-RoPE (P2). Read-only. Dense locally, strided by `C·P` in global index (§2.1). |
+| `K_cache` | `[S/(C·P), d]` | per-rank HBM, base `K_base[rank]`, contiguous | post-RoPE (P2). Read-only. A **contiguous** global position block `[r·B,(r+1)·B)` (§2.1). |
 | `V_cache` | `[S/(C·P), d]` | per-rank HBM, base `V_base[rank]` | raw V (P4). Read-only. |
 | `global_token_counter` | scalar | launch arg | kernel derives local len, slot↔global, causal bounds. |
 | `start_cube`, `start_pe` (= `f(request_id)`) | scalars | launch arg | Level-1/Level-2 rotation. |
@@ -364,6 +375,14 @@ out-projection must consume them accordingly:
 | **Decode** (§4 reduce) | `O = [G, 1, d]` (all heads) | `O_base` at the **CUBE-Group root** | head-replicated → all `G` heads end on one rank after the 2-level reduce. |
 | **Prefill** (§5.5 ring) | `O_i = [1, T_q, d]` (one head each) | per-CUBE `O_base[i]`, **distributed** | head-parallel → CUBE `i` writes head `i`'s rows in place; no reduce. |
 
+**Handoff contract to out-projection (downstream, separate kernel).** This
+ADR *guarantees* the two layouts above; designing out-proj is out of scope,
+but the handoff must match: **decode** delivers all `G` heads co-located at
+the Group root (out-proj reads them locally, or scatters per its own SP);
+**prefill** leaves head `i` on CUBE `i` (out-proj is itself head-parallel,
+or must **gather** the `G` heads first). Pin this per-kernel in the
+`out_proj` launch contract (§B decode-split item 2).
+
 **Decode intermediate** (non-root rank, on IPCQ — not kernel-visible):
 `(m_i, ℓ_i, O_i)` pushed up the 2-level tree to the parent (Level-2 PE
 parent, then Level-1 CUBE parent) via `tl.send` (§4). `O_i` is the heavy
@@ -433,8 +452,10 @@ arithmetic) lives in the **kernel** (plain Python `if`/arithmetic in the
 greenlet body). This is exactly what kernbench's greenlet model already
 permits (`kernel_runner.py`, ADR-0020 D3).
 
-> **What this supersedes.** An earlier iteration proposed a pure greenlet
-> primitive path (all `tl.dot`, no composite) on the argument that
+> **What this supersedes.** An earlier iteration proposed a pure
+> primitive-op path (all `tl.dot`, no composite — both are issued from the
+> greenlet, so the contrast is *scheduler-managed composite* vs
+> *kernel-issued ops*, not "greenlet vs composite") on the argument that
 > "composite yields no latency benefit." That holds **only because** the
 > simulator currently charges **zero** per-op CPU issue cost
 > (`dispatch_cycles=0`, `pe_cpu.py`) — it models away exactly the CPU
@@ -472,16 +493,15 @@ One KV head's sequence is sharded across `C × P` ranks (Level-1 inter-CUBE
 
 - Per-rank KV buffers sized to `⌈max_context / (C·P)⌉ × d × dtype`, K and V
   separately.
-- Within a rank, assigned tokens are **appended contiguously** (slot
-  0,1,2,…); the per-rank buffer is dense ⇒ DMA stays contiguous, even
-  though global indices are strided by `C·P`.
-- Slot → global (hierarchical round-robin): `rank = cube_local·P + pe_local`
+- Each rank's **contiguous position block** is stored dense (`B =
+  ⌈max_context/(C·P)⌉` slots) ⇒ DMA stays contiguous.
+- Slot → global (**contiguous block**): `rank = cube_local·P + pe_local`
   where `cube_local = (cube_id − start_cube) mod C`,
   `pe_local = (pe_id − start_pe) mod P`; then
-  `global_idx = local_slot·(C·P) + rank`.
+  `global_idx = rank·B + local_slot`.
 
-> **Shared prefill/decode layout — use *contiguous* blocks, not
-> round-robin.** Because prefill *writes* the KV cache (`qkv_rope`,
+> **Why the placement is *contiguous* blocks (rationale).** Because
+> prefill *writes* the KV cache (`qkv_rope`,
 > upstream) and decode *reads + extends* the same cache, both kernels share
 > one physical layout — no reshard at the prefill→decode boundary. Prefill's
 > **causal skip needs contiguous position blocks** (a late contiguous block
@@ -492,8 +512,14 @@ One KV head's sequence is sharded across `C × P` ranks (Level-1 inter-CUBE
 > the `C` CUBE-level blocks. *Caveat:* contiguous under-uses ranks for
 > **short** context (only frontier ranks hold data) — acceptable given the
 > long-context target; short-context balance is a separate study (§B).
-> (The round-robin formula above is the alternative for decode-only balance;
-> the contiguous block form is the one that serves both kernels.)
+> **Fallback:** for short/early decode where contiguous-block sharding
+> under-utilizes ranks, the driver may select a smaller `C` (or `C=P=1`); SP
+> is enabled only past the context-length threshold where KV-sweep time
+> dominates reduction and placement imbalance.
+> **Rejected: round-robin** (interleaved `token i → rank (start+i) mod (C·P)`)
+> — better short-context decode balance *alone*, but a round-robin block
+> spans all positions so it cannot causal-skip and cannot be shared with
+> prefill.
 
 ### 2.2 Driver per-launch duties (minimal)
 The driver supplies, per launch, the bases + counter + rotation; the
@@ -508,19 +534,21 @@ kernel derives the rest:
 | `cube_id`, `pe_id` (`tl.program_id` 1/0), `C`, `P` | CUBE-Group geometry |
 | `q_block_meta` | prefill/SP query block start + length |
 
-Let `R = C·P`, `cube_local = (cube_id − start_cube) mod C`,
-`pe_local = (pe_id − start_pe) mod P`, and this rank's index
-`rank = cube_local·P + pe_local`. Kernel-derived (plain arithmetic):
-- **my turn to write this step:** `(start_rank + counter) mod R == rank`
-- **my valid length:** `base = counter // R; rem = counter % R;
-  my_len = base + (1 if rank < rem else 0)`
+Let `R = C·P`, `B = ⌈max_context/R⌉`, `cube_local = (cube_id − start_cube)
+mod C`, `pe_local = (pe_id − start_pe) mod P`, and this rank's index
+`rank = cube_local·P + pe_local`. Kernel-derived (plain arithmetic, with
+`counter` = # filled tokens):
+
+- **my valid length** (contiguous block `[rank·B, (rank+1)·B)`):
+  `my_len = clamp(counter − rank·B, 0, B)`.
+- **my turn to write this step:** the new decode token (position `counter`)
+  lands on rank `counter // B` (its contiguous block).
 - **read range:** tiles `0 .. ⌈my_len / TILE⌉`.
 - **causal bounds / per-tile skip:** from global positions of the query
   block vs each tile.
 
-Driver = bases + counter + rotation. The single formula
-`(request_id + token_idx) mod (C·P)` over the two SP axes is the entire
-placement policy.
+Driver = bases + counter + rotation. The placement policy is the single
+contiguous-block map `token t → rank ⌊t/B⌋` (rotated by `start_*`).
 
 ---
 
@@ -728,10 +756,10 @@ def gqa_decode_sp(q_ptr, k_ptr, v_ptr, o_ptr, counter, start_pe, start_cube,
 
 ### 5.3 DECODE, with SP / KV-parallel (`C × P` ranks)
 
-- One request's KV is hierarchical-round-robin across the `C·P` ranks of
+- One request's KV is **contiguous-block sharded** across the `C·P` ranks of
   the CUBE Group (Level-1 over `C` CUBEs, Level-2 over `P` PEs); each rank
-  owns ≈`my_len` tokens and GQA-batches its `G=8` query rows over its local
-  tiles.
+  owns its `B`-token block and GQA-batches its `G=8` query rows over its
+  local tiles.
 - `T_q=1` ⇒ short per-rank sweep ⇒ **reduction dominates** ⇒ the §4 2-level
   reduce (`⌈log₂ P⌉` + center-mesh over `C`) is the structurally important
   part. Reduction latency is hidden by other concurrent decode tokens when
@@ -757,6 +785,23 @@ head is **big**, so reducing it across ranks would move `[T_q,d]` per rank;
 instead we **shard the heads** and **move the (also big) KV**, which each
 head needs in full anyway.
 
+> **When ring beats reduce (cost model — the choice is *not* just "O is
+> big").** Compare **total cross-rank bytes**, head dim `d` factored out:
+>
+> - **reduce** ≈ `G·T_q · ⌈log₂(C·P)⌉` — the payload is `(m,ℓ,O)` but **`O`
+>   dominates** (`m,ℓ` are `[G·T_q]` scalars vs `O` = `[G·T_q, d]`), so this
+>   is `O`-bytes × tree hops; head-replicated, so the `G` factor is present.
+> - **ring** ≈ `2·S` — the **total K+V bytes one CUBE injects** over the
+>   `C−1` rotations (`2·(S/C)·(C−1) ≈ 2·S`); `recv_async` **pipelines** the
+>   rotations so *latency* ≈ the slowest single step, but *traffic* is the
+>   sum. **No `G` factor, no log** — one KV slice is *shared* across all `C`
+>   heads by the rotation.
+>
+> So **ring < reduce (by total bytes) when `T_q > 2S / (G·⌈log₂(C·P)⌉)`** —
+> decode (`T_q=1`) → **reduce**, prefill (`T_q≈S`, RHS `≈ 2/(G·log) ≪ 1`) →
+> **ring**. KV is big, but the reduce's `O` carries the extra `G·log` factor,
+> so ring still wins for prefill. (Threshold sweep: §9.)
+
 - **Head-parallel placement:** within a CUBE Group, **CUBE `i` owns exactly
   one query head `i`** (the `G` query heads → the `C=G` CUBEs, one Q head per
   CUBE) and KV slice `i`. Each CUBE computes **its one Q head's** full
@@ -786,12 +831,14 @@ The three decode variants are shown **in full in the TL;DR** (Kernel 1):
 **opt1** current `CompositeCmd` (has a GEMM-engine bubble while the CPU
 auto-waits on `Sj`), **opt3** software pipelining (issue the next tile's
 `Q·Kᵀ` before this tile's softmax, `Sj` in a persistent double buffer —
-removes the bubble, no new command kind), **opt2** `ex_composite` (split
+*hides* the bubble subject to scheduler+engine balance, no new command
+kind), **opt2** `ex_composite` (split
 into `#1` Q·Kᵀ = existing composite reading K first, `#2` softmax+P·V+
 accumulator = the only new flash-epilogue machinery, reading V later).
 
-**Recommend:** ship **opt3** now (no new machinery, removes the bubble);
-revisit **opt2** once the cost model (ADR-0064) makes the fewer-CPU-issues
+**Recommend:** ship **opt3** now (no new machinery, hides the bubble subject
+to scheduler+engine balance); revisit **opt2** once the cost model (ADR-0064)
+makes the fewer-CPU-issues
 win measurable (§8 item 4). The MATH engine already has max/sum/exp — the
 only genuinely new part of opt2 is the composite's **stateful `(m,ℓ,O)`
 accumulator**, not the ops.
@@ -859,7 +906,7 @@ new primitive** — only the kernel restructuring in §5.2 (per KV head,
   pure kernel control flow over `tl.send`/`tl.recv`.
 - Causal tile skip + additive boundary mask (§3/§5.4) — kernel `if` +
   a mask tensor added with `+`.
-- 2-level round-robin KV placement / valid-length arithmetic (§0, §2) —
+- 2-level contiguous-block KV placement / valid-length arithmetic (§0, §2) —
   launch-arg arithmetic + `DPPolicy` `row_wise` over both `cube` and `pe`.
 
 **New primitives for efficiency / scale (each has a supporting ADR):**
@@ -964,6 +1011,14 @@ reshape-not-transpose, and `bf16` is modelled as `f16`. Verification is
 therefore **structural/timing-first**, with numeric parity as a bounded
 secondary check.
 
+**Gate type.** *Absolute* latency targets are deferred — they are only
+meaningful once ADR-0064's cost model lands (`dispatch_cycles=0` today). The
+gates below are (a) **structural** (op_log shape) and (b) **relative-to-
+baseline** ratios that *are* testable now and must pass: reduce rounds
+`⌈log₂P⌉+center-mesh` **vs baseline `C·P−1`**; KV bytes `H_kv·S·d` **vs
+`H_q·S·d` (G×=8× less)**; causal GEMM count = lower-triangular **vs full
+grid**; prefill ring traffic = `C` KV rotations **vs** a `[G·T_q,d]` reduce.
+
 1. **Runs in data mode (`enable_data=True`):** the GQA kernel
    (`h_q = G·h_kv`) completes without the byte-conservation error that the
    baseline head-packing hits — for all four cases. (This needs the §5.2
@@ -1005,8 +1060,8 @@ secondary check.
 
 ## B. Open design items from the hybrid pivot (review later)
 
-These arose when the decision moved from a pure greenlet primitive path to
-the **composite hybrid + lazy `tl.load`** (this revision). None blocks the
+These arose when the decision moved from a pure primitive-op (`tl.dot`) path
+to the **composite hybrid + lazy `tl.load`** (this revision). None blocks the
 design; each needs a verification pass during implementation. Recorded here
 (rather than asked) per the working agreement — the recommendation is my
 predicted default; revise on review.
@@ -1095,11 +1150,12 @@ predicted default; revise on review.
     none found in this kernel's scope, but the `configure_sfr_*` neighbour
     wiring must expose CUBE-NOC `N/S/E/W` for Level-1, not just PE IPCQ.
 
-5. **`valid_len_2level` / hierarchical round-robin correctness.** The
-    two-axis `(cube_local·P + pe_local)` placement (§2.1) must tile the
-    sequence with no gaps/overlaps and keep each rank's local buffer dense.
-    **Recommend:** unit-test the index math (every global token maps to
-    exactly one rank; per-rank slices are contiguous) before the kernel.
+5. **`valid_len_2level` / contiguous-block correctness.** The two-axis
+    `rank = cube_local·P + pe_local`, block `[rank·B, (rank+1)·B)` placement
+    (§2.1) must tile the sequence with no gaps/overlaps and keep each rank's
+    local buffer dense. **Recommend:** unit-test the index math (every global
+    token maps to exactly one rank; per-rank blocks are contiguous) before
+    the kernel.
 
 ### Items from the decode-reduce / prefill-ring split (this revision)
 
@@ -1147,7 +1203,8 @@ predicted default; revise on review.
 7. **Decode CPU-pipelining variant to ship (§5.6).** Three decode variants
     exist (opt1 current / opt3 software-pipelining / opt2 ex_composite).
     **Recommend:** implement **opt3** (software pipelining: issue next Q·Kᵀ
-    before this tile's softmax, `Sj` in a persistent double buffer) — removes
-    the GEMM-engine bubble with no new command kind. Defer **opt2** (the
+    before this tile's softmax, `Sj` in a persistent double buffer) — hides
+    the GEMM-engine bubble (subject to scheduler+engine balance) with no new
+    command kind. Defer **opt2** (the
     two-composite `ex_composite`, only `#2` is new) until ADR-0064's cost
     model makes its fewer-issues win measurable.