ccl: add missing configure_sfr_intracube_pe_ring (fix milestone-gqa import)

The single_user_prefill/single_user_decode panels of milestone-gqa-llama70b
(landed in e748a62) import configure_sfr_intracube_pe_ring from
kernbench.ccl.sfr_config, but the function definition was left in the
local working tree and never made it into that commit. Because the
bench registry eager-imports every bench module at CLI startup, the
missing symbol breaks `kernbench` entirely and cascades into 7 test
failures + 18 collection errors for anyone on a fresh checkout.

Installs an 8-PE logical ring inside every cube on every SIP (E/W
direction labels, no intercube or inter-SIP edges). Mutually exclusive
with configure_sfr_intercube_multisip on the same engine — the bench
picks one per panel-run (ADR-0058 D2).

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

.claude/commands/report.md

@@ -0,0 +1,327 @@
+---
+description: Generate a public-facing architecture design document from approved ADRs and SPEC.md, with gap analysis reported to chat only.
+---
+
+# `/report` — Architecture Design Document Generator
+
+Generates a **public-facing** architecture design document at
+`docs/report/architecture-{YYYY}-{1H|2H}.md` derived from the current ADR
+corpus, SPEC.md, CLAUDE.md, and the canonical component list.
+
+This command is **strictly read-only** on `docs/adr/`, `SPEC.md`,
+`CLAUDE.md`, and `src/`. The only write is the report file itself
+(a derived artifact under `docs/report/`).
+
+---
+
+## Invocation
+
+Two modes:
+
+- `/report` — **dry-run** (default). No file is written. The command
+  reads sources, performs classification, and reports the planned TOC
+  + gap analysis to chat only. Use this to validate ADR-to-section
+  mapping before committing.
+- `/report write` — **write mode**. Performs the same procedure and
+  writes `docs/report/architecture-{period}.md`. Use after a dry-run
+  whose classification looks correct.
+
+Period determination (both modes), from system date:
+
+- month 1–6 → `{YYYY}-1H`
+- month 7–12 → `{YYYY}-2H`
+
+In write mode, if `docs/report/architecture-{period}.md` already exists,
+overwrite it without asking (regeneration is the expected operation).
+
+---
+
+## Output Contract
+
+### Document body (`docs/report/architecture-{period}.md`)
+
+Public release form. Reader is an external developer/architect. They do
+**not** have access to SPEC.md or ADR files. Therefore:
+
+- **No `ADR-NNNN` identifiers** in visible prose.
+- **No `SPEC R/§` identifiers** in visible prose.
+- **No internal jargon** assumed without definition.
+- **No diagram embeds** — only `<!-- DIAGRAM: ... -->` placeholders.
+- **Attribution via HTML comments** — every prose paragraph that derives
+  from a source carries an inline comment immediately above it:
+  `<!-- src: ADR-NNNN <section-name> -->` (multiple sources allowed).
+
+### Chat-only report (not written to any file)
+
+After writing the document, report to the user in the chat response:
+
+- File path written.
+- Section counts (e.g., "Detailed Architecture: 8 components covered,
+  2 in `builtin/` have no ADR backing").
+- **G1 gaps** — SPEC requirements (R-numbers / §) with no ADR citing them.
+- **G2 gaps** — ADRs missing **Context** or **Decision**. Alternatives
+  and Consequences are optional; their absence is NOT a gap.
+- **G3 gaps** — ADR cross-references without a back-reference.
+  Only flag when the referencer's ADR number is **less than** the
+  referenced ADR's number (older → newer). Newer ADRs citing older
+  infrastructure ADRs (higher number → lower number) are expected to
+  be one-way and are NOT flagged.
+- **G4 suggestions** — areas where an ADR seems missing based on the
+  ADR corpus + SPEC reading. Phrase as suggestions, not findings. Each
+  G4 item must say *why* it's suggested and remain falsifiable.
+- **G5 consistency issues** — ADR-to-ADR inconsistencies:
+  - **G5a (supersession not reflected)** — ADR-A states it supersedes
+    ADR-B, but ADR-B's Status is not marked as Superseded.
+  - **G5b (merge candidates)** — two or more ADRs cover near-identical
+    scope (detected naturally during section assignment, not via
+    exhaustive pair-wise scan).
+  - **G5c (explicit contradictions)** — two ADRs whose Decisions
+    directly oppose each other. Must cite both quotations; do not
+    speculate contradictions from topical similarity alone.
+- **TOC rationale** — for each section, list contributing ADR IDs
+  (this is for the user's verification only, never written to the
+  document itself).
+
+G4 must never appear in the document body. G1–G3 are also chat-only.
+
+---
+
+## Procedure
+
+### Step 1 — Determine period
+
+Use current system date. Compute `{YYYY}-1H` or `{YYYY}-2H`.
+
+### Step 2 — Ingest ADRs
+
+For each `docs/adr/ADR-NNNN-*.md`:
+
+- If both `ADR-NNNN-*.md` (Korean) and `ADR-NNNN-*.en.md` (English)
+  exist for the same number, **prefer the Korean `.md`** version.
+- Parse for the four canonical sections: Context, Decision, Alternatives
+  (also accept "Alternatives Considered"), Consequences.
+- Record presence/absence of **Context** and **Decision** for G2.
+  Alternatives and Consequences presence is recorded for use during
+  authoring, but their absence is not a gap.
+- Record ADR-NNNN cross-references for G3, preserving the direction
+  (referencer → referenced). G3 evaluation uses ADR numbers to
+  distinguish older→newer (flagged when missing back-link) from
+  newer→older (not flagged; see *Output Contract* G3).
+- Record Status (e.g., Accepted, Superseded, Draft) and any "supersedes
+  ADR-NNNN" text in the body for G5a.
+
+Process ADRs in **numerical order** for determinism.
+
+### Step 3 — Read canonical component list
+
+List `src/kernbench/components/builtin/*.py`, excluding `__init__.py`,
+`pe_types.py`, and `__pycache__/`. Sort alphabetically. This is the
+canonical order for Detailed Architecture subsections.
+
+### Step 4 — Read SPEC.md and CLAUDE.md
+
+For G1 detection: extract every `R<N>` and `§<X.Y>` identifier mentioned
+in SPEC.md. For each ADR, check which of these it cites. SPEC IDs with
+zero citing ADRs → G1.
+
+### Step 5 — Section assignment
+
+Assign each ADR to exactly one of:
+
+- **Design Principles** — project-wide rationale, philosophy, mission
+  (e.g., "why source-level kernel execution", "why fast multi-device
+  scaling"). Includes ADRs that describe foundational invariants
+  (e.g., latency model assumptions, verification strategy).
+- **High-level Architecture** — Tray / SIP / CUBE / PE hierarchy and
+  cross-layer boundaries (e.g., runtime API ↔ sim_engine ↔ components).
+- **Detailed Architecture** — single-component internal designs. One
+  subsection per file in the canonical component list. ADRs whose
+  primary topic is the internal structure of one component go here.
+- **Implementation Decisions** — **cross-cutting** algorithms / policies
+  / schemes / models that don't belong to a single component:
+  collective algorithms, parallelization policies, address schemes,
+  routing algorithms, model assumptions.
+
+Boundary rule between Detailed Architecture and Implementation Decisions:
+
+> Detailed Architecture = component-internal.
+> Implementation Decisions = spans multiple components OR is an
+> algorithm/policy/scheme/assumption rather than a structural choice.
+
+If an ADR fits two sections plausibly, prefer the one that minimizes
+duplication and pick the more specific bucket (Detailed if it primarily
+concerns one component, else Implementation Decisions).
+
+During classification, opportunistically detect ADR consistency issues:
+
+- **G5b (merge candidate)** — if two or more ADRs land in the same
+  Detailed Architecture subsection or the same Implementation Decisions
+  topic AND their primary scope is near-identical, record as a merge
+  candidate. Topical adjacency is not enough; the scopes must be
+  effectively the same question.
+- **G5c (explicit contradiction)** — if while reading you encounter two
+  ADRs whose Decisions directly oppose each other on the same question,
+  record both quotations verbatim with their ADR IDs. Do NOT speculate
+  contradictions from similarity, vocabulary, or domain overlap — only
+  explicit, citable opposition.
+
+Do NOT perform an exhaustive pair-wise scan of all ADRs. G5b/G5c are
+byproducts of normal reading; if not encountered, the chat report
+shows "(none)".
+
+### Step 6 — Write the document (write mode only)
+
+In **dry-run mode**, skip this step entirely. Proceed directly to Step 7.
+
+```markdown
+# KernBench — Architecture Design Document
+*{YYYY} {1H|2H}*
+
+## Design Principles
+<prose>
+
+## High-level Architecture
+<intro prose>
+
+### Tray
+### SIP
+### CUBE
+### PE
+
+## Detailed Architecture
+### <component-1>
+### <component-2>
+...
+
+## Implementation Decisions
+### <topic-1>
+### <topic-2>
+...
+```
+
+#### Authoring rules (apply to every section)
+
+- **Stay grounded.** Every claim must trace to an ADR's stated content
+  (Context / Decision / Alternatives / Consequences). No invented
+  motivation, no invented alternatives, no invented trade-offs.
+- **4-part discipline, naturally.** Each subsection should naturally
+  cover: the problem the design addresses, the decision made, the
+  alternatives considered, the consequences. Do **not** label these
+  with rigid headers like "**Problem.**" — weave them into prose. But
+  ensure all four are present *if the source ADR documents them*.
+- **Missing → omit, not fabricate.** If a source ADR has no
+  "Alternatives" section, do **not** invent alternatives for the
+  report. Simply write the remaining parts and record G2 in chat.
+- **Attribution.** Every paragraph derived from one or more ADRs
+  carries an HTML comment immediately above:
+  `<!-- src: ADR-NNNN <section> [, ADR-MMMM <section>] -->`.
+- **Diagram placeholders.** Where a diagram would help, insert
+  `<!-- DIAGRAM: <short description of what the diagram should show> -->`
+  on its own line. **Never** embed an image (`![...](...)`).
+- **Public tone.** Self-contained. Define internal terms (SIP, CUBE,
+  PE, Tray, NOC, IPCQ, TCM, etc.) on first use within the document.
+  Do not assume reader has read SPEC or ADRs.
+- **No internal references.** No `ADR-NNNN` in body text. No
+  `SPEC §X.Y` or `R<N>` in body text. These appear only inside HTML
+  attribution comments.
+- **Detailed Architecture component subsections.** Use the canonical
+  list from Step 3 in order. For each component file, write a
+  subsection drawing from any ADR that primarily concerns that
+  component. If no ADR covers a component, write a one-line stub
+  noting the component exists and flag it in chat report. If an ADR
+  covers a topic not in the canonical list, place it under
+  "Detailed Architecture → Other" (sub-subsection) and flag for
+  canonical-list extension in chat.
+- **Implementation Decisions topic naming.** Derive topic names from
+  ADR titles, made reader-friendly (no ADR number). Group related
+  ADRs under one topic when natural (e.g., multiple address-related
+  ADRs under "Address Scheme").
+
+### Step 7 — Generate chat report
+
+After Step 6 (write mode) or directly from Step 5 (dry-run mode),
+emit the following to chat. Do **not** write any of this to a file.
+
+In **dry-run mode**, replace the `Wrote:` line with:
+`**DRY-RUN — no file written.** Review TOC and gaps below. Run \`/report write\` to commit.`
+
+```
+## /report — Generation Summary
+
+**Wrote:** docs/report/architecture-{period}.md
+
+**Section coverage**
+- Design Principles: <N> ADRs
+- High-level Architecture: <N> ADRs
+- Detailed Architecture: <covered>/<total> components ; components without ADR: [...]
+- Implementation Decisions: <N> topics, <N> ADRs
+
+**TOC rationale (ADR → section mapping)**
+- Design Principles: ADR-NNNN, ADR-MMMM
+- High-level Architecture: ...
+- Detailed Architecture → <component>: ADR-NNNN
+- Implementation Decisions → <topic>: ADR-NNNN, ADR-MMMM
+
+**G1 — SPEC requirements without ADR support**
+- R<N> / §<X.Y>: not cited by any ADR
+- (or "none")
+
+**G2 — ADRs missing required sections (Context or Decision)**
+- ADR-NNNN: missing <Context|Decision>
+- (or "none")
+
+**G3 — Broken cross-references** (older → newer only)
+- ADR-NNNN cites ADR-MMMM (NNNN < MMMM); ADR-MMMM does not back-reference
+- (or "none")
+- Note: newer ADRs citing older infrastructure ADRs (NNNN > MMMM) are
+  not flagged here — one-way references are the expected pattern.
+
+**G4 — Suggested topics that may warrant a new ADR (verify before acting)**
+- <topic>: <why agent thinks it may be missing — must be falsifiable>
+- (or "none")
+
+**G5 — ADR consistency issues**
+- **G5a (supersession not reflected)**
+  - ADR-NNNN claims to supersede ADR-MMMM, but ADR-MMMM Status is "<status>"
+  - (or "none")
+- **G5b (merge candidates)**
+  - ADR-NNNN + ADR-MMMM: near-identical scope on <topic> — evaluate merge
+  - (or "none")
+- **G5c (explicit contradictions)**
+  - ADR-NNNN says "<quote>"; ADR-MMMM says "<quote>" — direct opposition on <question>
+  - (or "none")
+```
+
+---
+
+## Constraints (do not violate)
+
+1. **Read-only on source.** No writes to `docs/adr/`, `SPEC.md`,
+   `CLAUDE.md`, or `src/`. Only write is
+   `docs/report/architecture-{period}.md`.
+2. **No fabrication.** Every body paragraph traces to ADR content via
+   HTML attribution comment.
+3. **No diagram embeds.** Placeholders only.
+4. **No internal IDs in body.** ADR-NNNN and SPEC R/§ stay inside
+   HTML comments only.
+5. **Determinism.** ADRs processed in numerical order; components in
+   canonical (alphabetical) order. Same inputs → same output.
+6. **G4 stays in chat.** Never written to the document.
+7. **Korean bilingual preference.** When both `.md` and `.en.md`
+   exist for the same ADR number, use `.md`.
+8. **All ADRs included.** No exclusion list. ADRs about internal
+   tooling (CLI, diagram views, verification strategy) are still
+   included — usually under Design Principles or Implementation
+   Decisions, written in publishable form.
+
+---
+
+## Failure modes to avoid
+
+- **Padding** with general background not present in the source ADRs.
+- **Inferring alternatives** the ADR doesn't mention.
+- **Quietly skipping** an ADR because it seems internal. Include it,
+  rephrase for public audience.
+- **Inventing components** not in `src/kernbench/components/builtin/`.
+- **Auto-selecting diagrams** from `docs/diagrams/`. Only placeholders.
+- **Promoting G4 suggestions to the document.** They stay in chat.

.claude/settings.json

@@ -9,7 +9,59 @@
       "Bash(python -m kernbench.cli.main probe --topology topology.yaml)",
       "Bash(xargs grep -l \"class.*ComponentBase\\\\|class.*DefaultComponent\")",
       "Bash(python -m pytest tests/test_probe.py -v)",
-      "Bash(python -m pytest tests/test_probe.py tests/test_component_registry.py -v)"
+      "Bash(python -m pytest tests/test_probe.py tests/test_component_registry.py -v)",
+      "Bash(python -m pytest -o \"addopts=\" --no-header tests/test_intercube_root_center.py)",
+      "Bash(python -m pytest -o \"addopts=\" --no-header tests/test_tp_layers.py tests/test_tp_mlp.py)",
+      "Bash(git commit -m ' *)",
+      "Bash(git stash *)",
+      "Bash(python scripts/emit_overview_with_external_ref.py)",
+      "Bash(where inkscape *)",
+      "Bash(\"/c/Program Files \\(x86\\)/Microsoft/Edge/Application/msedge.exe\" --headless --disable-gpu --screenshot=\"$\\(pwd\\)/docs/diagrams/cube_mesh_view.png\" --window-size=1400,1300 \"file:///$\\(pwd)",
+      "Bash(python scripts/build_overview_slides.py)",
+      "Bash(git fetch *)",
+      "Bash(git pull *)",
+      "Bash(python -m pytest --no-header tests/test_allreduce_buffer_kind_sweep.py)",
+      "Bash(python -m pytest --no-header tests/test_pe_to_pe_latency.py)",
+      "Bash(python -m pytest --no-header tests/test_ipcq_buffer_kind_locations.py -v)",
+      "Bash(python -m pytest --no-header tests/test_ipcq_buffer_kind_locations.py tests/test_ipcq_buffer_kind_latency.py tests/test_allreduce_buffer_kind_sweep.py)",
+      "Bash(git checkout *)",
+      "Bash(python -m pytest --no-header tests/test_ipcq_buffer_kind_latency.py::test_slot_write_latency_orders_tcm_hbm_sram)",
+      "Bash(python scripts/emit_ipcq_send_recv_model_plots.py)",
+      "Bash(python -m pytest --no-header tests/test_pe_to_pe_latency.py -x)",
+      "Bash(python -m pytest --no-header tests/test_pe_to_pe_latency.py tests/test_ipcq_buffer_kind_locations.py tests/test_ipcq_buffer_kind_latency.py tests/test_allreduce_buffer_kind_sweep.py)",
+      "Bash(kill %1)",
+      "Bash(awk '{print $2}')",
+      "Bash(xargs -r kill)",
+      "Bash(python scripts/_debug_op_log.py)",
+      "Bash(SWEEP_SHAPES=\"16,32,64,128,256\" python scripts/gemm_sweep.py)",
+      "Bash(python scripts/plot_gemm_sweep.py)",
+      "Bash(python scripts/gemm_sweep.py)",
+      "Bash(python scripts/gen_pe_pipeline_diagram.py)",
+      "Bash(python scripts/gen_matmul_32x128x32_diagram.py)",
+      "Bash(python -m pytest tests/test_pe_pipeline.py -x --tb=short)",
+      "Bash(python -m pytest tests/test_pe_pipeline.py tests/test_e2e_pipeline.py tests/test_op_log.py -x --tb=short -q)",
+      "Bash(ls -la C:/Users/mukes/.claude/projects/c--Users-mukes-Mukesh-ywkang-git-kernbench2/ 2>&1 | head -20)",
+      "Read(//c/Users/mukes/.claude/projects/c--Users-mukes-Mukesh-ywkang-git-kernbench2/**)",
+      "Bash(awk 'NR==1812 || NR==1815' C:/Users/mukes/.claude/projects/c--Users-mukes-Mukesh-ywkang-git-kernbench2/e55237ed-5c1f-4a89-a3b9-9b74fec45366.jsonl)",
+      "Bash(awk 'NR==1058' C:/Users/mukes/.claude/projects/c--Users-mukes-Mukesh-ywkang-git-kernbench2/e55237ed-5c1f-4a89-a3b9-9b74fec45366.jsonl)",
+      "Bash(awk -F: '$1 > 1700 && $1 < 1815 {print $1}')",
+      "Bash(awk 'NR==1812' C:/Users/mukes/.claude/projects/c--Users-mukes-Mukesh-ywkang-git-kernbench2/e55237ed-5c1f-4a89-a3b9-9b74fec45366.jsonl)",
+      "Bash(awk 'NR>=1815 && NR<=1825' C:/Users/mukes/.claude/projects/c--Users-mukes-Mukesh-ywkang-git-kernbench2/e55237ed-5c1f-4a89-a3b9-9b74fec45366.jsonl)",
+      "Bash(awk 'NR>1815' C:/Users/mukes/.claude/projects/c--Users-mukes-Mukesh-ywkang-git-kernbench2/e55237ed-5c1f-4a89-a3b9-9b74fec45366.jsonl)",
+      "Bash(awk 'NR==1839' C:/Users/mukes/.claude/projects/c--Users-mukes-Mukesh-ywkang-git-kernbench2/e55237ed-5c1f-4a89-a3b9-9b74fec45366.jsonl)",
+      "Bash(git log *)",
+      "Bash(python -m pytest tests/test_op_log.py tests/test_pe_components.py tests/test_pe_pipeline.py -x --tb=short)",
+      "Bash(python -m pytest tests/test_pe_to_pe_latency.py tests/test_e2e_pipeline.py tests/test_e2e_data.py tests/test_data_executor.py tests/test_pe_dma_ipcq.py -x --tb=short)",
+      "Bash(python -m pytest tests/test_pe_pipeline.py::test_pe_dma_record_start_after_channel_acquire -x --tb=long)",
+      "Bash(python -m pytest tests/test_pe_pipeline.py::test_pe_dma_record_start_after_channel_acquire -x --tb=short)",
+      "Bash(python -m pytest tests/test_op_log.py tests/test_pe_components.py tests/test_pe_pipeline.py tests/test_pe_to_pe_latency.py tests/test_e2e_pipeline.py tests/test_e2e_data.py tests/test_data_executor.py tests/test_pe_dma_ipcq.py --tb=short)",
+      "Bash(python -m pytest tests/test_pe_pipeline.py -q)",
+      "Bash(python -m pytest tests/test_pe_pipeline.py tests/test_triton_emu.py -q)",
+      "Bash(python -m pytest tests/test_composite_epilogue.py -v)"
+    ],
+    "additionalDirectories": [
+      "c:\\Users\\mukes\\Mukesh\\ywkang_git\\kernbench2\\tests",
+      "C:\\Users\\mukes\\Mukesh\\ywkang_git\\kernbench2\\tests\\pe2pe_latency_plots"
     ]
   }
 }

.claude/settings.local.json

@@ -30,7 +30,10 @@
       "Bash(python -m pytest tests/test_pe_components.py -v)",
       "Bash(python -m pytest tests/test_triton_emu.py -v)",
       "Bash(python -m pytest tests/test_pe_components.py tests/test_triton_emu.py -v)",
-      "Bash(python -m pytest tests/test_pe_components.py::test_mcpu_multi_pe_kernel_launch tests/test_pe_components.py::test_qkv_gemm_bench_multi_pe_completes -v)"
+      "Bash(python -m pytest tests/test_pe_components.py::test_mcpu_multi_pe_kernel_launch tests/test_pe_components.py::test_qkv_gemm_bench_multi_pe_completes -v)",
+      "Bash(git add:*)",
+      "Bash(git commit:*)",
+      "Bash(git push:*)"
     ]
   }
 }

.gitignore

@@ -29,4 +29,6 @@ build/
 
 # Logs
 *.log
-.claude/
+.claude/*
+!.claude/commands/
+!.claude/commands/*.md

CLAUDE.md

@@ -5,27 +5,10 @@ SPEC.md and ADRs are the source of truth.
 
 ---
 
-## Terminology
+# Part 1 — General Behavior
 
-- runtime API:
-  Host-facing public API used by benchmarks and user code (e.g., tensor deployment, kernel launch).
-- simulation engine (sim_engine):
-  Discrete-event engine responsible for request injection, scheduling, and completion tracking.
-- components:
-  Device-side nodes modeling hardware behavior (IO_CPU, M_CPU, PE_CPU, routers, engines, etc.).
-
-## Authority & Scope
-
-- SPEC.md defines the architectural contract.
-- ADRs (docs/adr/ADR-*.md) define non-trivial architectural decisions.
-- If a change conflicts with SPEC.md or an ADR:
-  - STOP.
-  - Explain the conflict.
-  - Propose options (keep spec, update ADR, or narrow scope).
-- Do NOT silently change architecture.
-- The repository structure reflects architectural intent; Claude Code MUST respect existing module boundaries and file locations.
-
----
+> Reusable across repos. Describes *how* Claude Code interacts with the user
+> and constructs changes, independent of this project's domain.
 
 ## Design Questions
 
@@ -37,14 +20,21 @@ SPEC.md and ADRs are the source of truth.
   - ADRs
 - If a design question implies a change, default to Phase 1.
 
----
+## Surfacing Choices
+
+Applies to both design discussions and Phase 1 proposals.
+
+- If multiple valid interpretations of the request exist, present them.
+  Do NOT pick one silently.
+- If a simpler approach exists, say so. Push back when warranted —
+  do NOT just implement the more complex path the user proposed.
+- State required assumptions explicitly. If uncertain, ask before assuming.
 
 ## Change & Test Protocol (Mandatory)
 
 All non-trivial changes MUST follow a two-phase process.
-Design discussion is always allowed; code changes are not.
-
----
+Design discussion is always allowed.
+Production code changes require Phase 1 approval before Phase 2 applies them.
 
 ### Phase 1 — Proposal + Verification
 
@@ -63,20 +53,18 @@ Design discussion is always allowed; code changes are not.
 - Explain why the change is needed.
 - Explain consistency with SPEC.md and relevant ADRs.
 
-1) **Verification Plan**
+2) **Verification Plan**
 
-- SPEC requirement(s) / ADR(s) affected (e.g., R1/R2/R5, ADR-0002).
+- SPEC requirement(s) / ADR(s) affected.
 - Tests that validate the change:
   - existing tests to run, and/or
   - new tests to add.
-- Concrete input cases used by the tests:
-  - topology (SIP / CUBE / PE layout)
-  - request parameters (src, dst, size_bytes).
-- Expected observable assertions, such as:
-  - hop trace contains key waypoints,
-  - latency invariants (e.g., > 0, monotonic increase),
-  - deterministic route selection.
-  - **expected changes (or no changes) in generated diagrams**, if applicable.
+- Concrete input cases used by the tests.
+- Expected observable assertions.
+- Expected changes (or no changes) in generated artifacts, if applicable.
+
+(Project-specific expectations for what these inputs/assertions look like:
+see Part 2 → *Verification Plan — Project Expectations*.)
 
 If the Verification Plan is missing or vague, STOP.
 
@@ -89,7 +77,13 @@ If the Verification Plan is missing or vague, STOP.
 
 - Any production code changes
 - Any SPEC.md or ADR modifications
-- Any production diff output
+- Final, ready-to-apply unified diffs (Phase 2 only)
+
+#### Permitted for design discussion
+
+- Pseudocode, interface sketches, type signatures
+- Small illustrative snippets to clarify a design point
+- "Before / after" excerpts (not full diffs)
 
 #### Phase 1 Output
 
@@ -100,8 +94,6 @@ If the Verification Plan is missing or vague, STOP.
   - "No Phase 2 needed" OR
   - "Await approval for Phase 2"
 
----
-
 ### Phase 2 — Apply + Verify + Rollback
 
 #### Trigger
@@ -112,10 +104,10 @@ Phase 2 is triggered ONLY by the exact user approval phrase:
 
 #### Phase 2 Rules
 
-- Output **minimal unified diffs only**
-- Modify ONLY production files declared in Phase 1
-- Do NOT include explanations, comments, or unchanged code
-- Automatically apply the diff to the working tree
+- Keep changes minimal and scoped to the approved Phase 1 proposal.
+- Modify only production files declared in Phase 1.
+- Avoid unrelated edits, cleanup, or formatting churn.
+- Automatically apply approved changes to the working tree.
 
 #### Mandatory Verification
 
@@ -126,7 +118,7 @@ Phase 2 is triggered ONLY by the exact user approval phrase:
 If ALL tests PASS:
 
 - Keep the applied changes
-- Ensure generated diagrams (if affected) are consistent
+- Ensure generated artifacts (if affected) are consistent
 - Report success concisely
 
 #### Failure Path (Mandatory)
@@ -143,8 +135,210 @@ If ANY test FAILS:
 
 Tests must NEVER be weakened, removed, or altered to force Phase 2 to pass.
 
+Failing tests may indicate:
+- invalid assumptions,
+- architectural violations,
+- or incomplete modeling.
+
+Do not assume the test is wrong without explicit evidence.
+
+## Allowed Exceptions
+
+(Protocol Still Required)
+
+- comments or docstrings
+- formatting-only changes
+- type annotation changes with no runtime behavior change
+
+In exceptions, Phase 1 MUST explicitly state:
+**"No behavior change; tests unchanged."**
+
+## Coding Style
+
+Applies to all production code changes (Phase 2) and test code (Phase 1).
+The Phase 1/2 protocol decides *whether* and *what* to change;
+this section decides *how* the resulting diff should look.
+
+### Simplicity First
+
+**Minimum code that solves the problem. Nothing speculative.**
+
+- Write the minimum code that satisfies the Phase 1 proposal.
+- No abstractions for single-use code.
+- No "flexibility"/"configurability" not declared in Phase 1.
+- No error handling for impossible scenarios.
+
+Ask yourself: "Would a senior engineer say this is overcomplicated?" If yes, simplify.
+
+### Surgical Changes
+
+**Touch only what you must. Clean up only your own mess.**
+
+- Touch only files declared in the Phase 1 proposal.
+- Don't "improve" adjacent code, comments, or formatting.
+- Match existing style in the file, even if you'd do it differently.
+- If your changes orphan imports/variables/functions, remove them.
+- If you notice pre-existing dead code, do NOT delete it silently.
+  Mention it, and present options:
+    (a) delete (with approval),
+    (b) keep as-is,
+    (c) refactor to make it reachable / repurposed.
+  Let the user choose before acting.
+- Every changed line must trace to the Phase 1 proposal.
+
+## Enforcement Defaults
+
+General fallbacks. Apply to anything not explicitly covered above.
+
+- If unsure whether a change is non-trivial → treat it as non-trivial.
+- If unsure whether Phase 2 is allowed → STOP and ask.
+
 ---
 
+# Part 2 — Project-Specific (kernbench)
+
+> Specific to this repo's domain (SIP/CUBE/PE topology, runtime API, sim_engine).
+> Replace this entire Part when adapting the framework to another repo.
+>
+> Contains **foundations** (Authority & Scope → Terminology → Terminology
+> Discipline → Mental Model → Common Failure Modes) followed by **rules**
+> (Non-Trivial, Verification Plan, CLI, Derived Artifacts, ADR Translation
+> Discipline, runtime API / sim_engine Boundaries).
+
+## Authority & Scope
+
+- SPEC.md defines the architectural contract.
+- ADRs (docs/adr/ADR-*.md) define non-trivial architectural decisions.
+- If a change conflicts with SPEC.md or an ADR:
+  - STOP.
+  - Explain the conflict.
+  - Propose options (keep spec, update ADR, or narrow scope).
+- Do NOT silently change architecture.
+- The repository structure reflects architectural intent; Claude Code MUST respect existing module boundaries and file locations.
+
+### ADR Lifecycle
+
+ADRs live in one of four folders. Three carry **canonical English**
+content based on lifecycle state; the fourth holds Korean translations:
+
+- `docs/adr/` — **Accepted** (canonical English; current
+  implementation reflected).
+- `docs/adr-proposed/` — **Proposed**, **Stub**, or **Draft** (design
+  only / future-work exploration / retroactive documentation pending
+  verification). **Authoring language is free** (any language); the
+  promotion step (below) translates to English.
+- `docs/adr-history/` — **Superseded** or **Merged** (no longer the
+  authoritative source; kept as historical record). Frozen — language
+  policy not applied retroactively.
+- `docs/adr-ko/` — Korean translations of accepted ADRs (derived
+  artifact, 1:1 mirror of `docs/adr/`). English in `docs/adr/` is the
+  canonical source of truth; when KO and EN disagree, EN wins. See
+  *ADR Translation Discipline* below.
+
+Status field values:
+
+- `Accepted` — design is in current implementation.
+- `Proposed` — design is concrete but not yet implemented.
+- `Stub (Future Work)` — design space exploration; no commitment yet.
+- `Draft` — retroactive documentation drafted but not yet verified
+  against the implementation it describes.
+- `Superseded by ADR-NNNN` — replaced by another ADR.
+- `Merged into ADR-NNNN` — content absorbed by another ADR.
+
+Transitions:
+
+- **Proposed/Stub → Accepted**: when the ADR's decisions are
+  reflected in production code AND covered by tests. If the proposed
+  ADR is in Korean, translate to English and place the English in
+  `docs/adr/`; move the Korean original to `docs/adr-ko/`. If the
+  proposed ADR is in English, `git mv` it to `docs/adr/` and create
+  the Korean translation in `docs/adr-ko/`. Change Status to
+  `Accepted` in both files.
+- **Draft → Accepted**: when the ADR's text has been verified to
+  accurately describe the existing implementation. Same English /
+  Korean placement rule as above.
+- **Accepted → Superseded**: set Status to `Superseded by ADR-MMMM`
+  in both the EN and KO files and `git mv` both to their respective
+  history locations (`docs/adr-history/` for English; the KO copy
+  stays in `docs/adr-ko/` only if it was already mirrored — see *ADR
+  Translation Discipline* for the frozen-history exception).
+- **Accepted → Merged**: set Status to `Merged into ADR-MMMM`
+  (single-line stub) in both files and apply the same `git mv` rule
+  as the Superseded transition.
+
+Cross-references between ADRs use the `ADR-NNNN` ID and remain valid
+regardless of folder location. ADR numbers are **immutable**; never
+renumber. Numbering holes from moved ADRs are expected.
+
+## Terminology
+
+- runtime API:
+  Host-facing public API used by benchmarks and user code (e.g., tensor deployment, kernel launch).
+- simulation engine (sim_engine):
+  Discrete-event engine responsible for request injection, scheduling, and completion tracking.
+- components:
+  Device-side nodes modeling hardware behavior (IO_CPU, M_CPU, PE_CPU, routers, engines, etc.).
+
+## Terminology Discipline
+
+Use only terms established in SPEC.md, ADRs, existing notes, or code.
+Do not coin new terms (status labels, tiers, classifications, role names)
+without explicit user approval. When a needed term is missing or ambiguous,
+ask before introducing one. When proposing a rename, show the existing
+term and the proposed change side-by-side and wait for approval.
+
+## Mental Model
+
+The simulator is layered along **request flow**:
+
+  runtime API           (host-facing: tensor ops, kernel launch;
+                         topology-agnostic, no routing — ADR-0007)
+       ↓
+  sim_engine            (schedules events, routes requests,
+                         tracks completion via correlation IDs)
+       ↓
+  components            (device-side nodes: IO_CPU, M_CPU, PE_CPU,
+                         routers, engines — model HW behavior
+                         including interconnect)
+
+Configuration & decisions (orthogonal to request flow):
+- **topology**  — compiled at config time (ADR-0006); defines which
+  components exist and how they connect. Authoritative graph for sim_engine.
+- **policy** (routing / address / placement) — consulted by sim_engine
+  during request handling.
+
+Invariant: all latency arises from **explicit scheduled events on modeled
+components and links** (SPEC §0.1, R8). No implicit waits, no magic delays.
+
+Stay within layer boundaries; do not collapse or bypass for convenience.
+
+## Common Failure Modes
+
+Anti-patterns that violate the Mental Model or Golden Invariants (SPEC §0.1).
+If your change does any of these, STOP and reconsider.
+
+- **runtime topology mutation** — topology is compiled at config time; do not
+  add/remove nodes or edges during simulation (ADR-0006).
+- **nondeterministic iteration order** — never iterate sets, unordered dicts,
+  or anything else with implementation-defined order on the critical path.
+  Determinism is required (SPEC §0.1).
+- **routing policy inside runtime API** — runtime API is topology-agnostic;
+  routing/fan-out belongs in policy + sim_engine (ADR-0007).
+- **latency modeled outside sim_engine scheduling** — every delay must come
+  from an explicit scheduled event on a modeled component or link
+  (SPEC §0.1, R8). No magic sleeps, no hardcoded constants smuggled in.
+- **hidden cross-layer coupling** — do not skip layer interfaces.
+  e.g., runtime API must not call into components directly, bypassing sim_engine.
+- **silent ADR/SPEC reinterpretation** — surface conflicts; do not paper over them.
+  See *Authority & Scope* above.
+- **weakening tests to make Phase 2 pass** — fix the code, not the test.
+  See *Part 1 → Phase 2 → Failure Path*.
+- **asserting from memory without source check** — quantitative
+  architectural facts (topology counts, sizes, latencies, address widths,
+  port arities) must be sourced from SPEC.md or a specific ADR before
+  assertion. Memory is unreliable. If the source is silent, surface the
+  gap rather than guessing.
+
 ## What Counts as "Non-Trivial"
 
 (Protocol Required)
@@ -158,39 +352,82 @@ Any of the following:
 - changes affecting determinism or connectivity
 - changes touching two or more production files
 
----
+## Verification Plan — Project Expectations
 
-## Allowed Exceptions
+Concrete forms that Part 1's *Verification Plan* MUST take in this repo:
 
-(Protocol Still Required)
-
-- comments or docstrings
-- formatting-only changes
-- type annotation changes with no runtime behavior change
-
-In exceptions, Phase 1 MUST explicitly state:
-**"No behavior change; tests unchanged."**
-
----
+- SPEC requirement(s) / ADR(s) affected (e.g., R1/R2/R5, ADR-0002).
+- Concrete input cases:
+  - topology (SIP / CUBE / PE layout)
+  - request parameters (src, dst, size_bytes).
+- Expected observable assertions, such as:
+  - hop trace contains key waypoints,
+  - latency invariants (e.g., > 0, monotonic increase),
+  - deterministic route selection.
+  - **expected changes (or no changes) in generated diagrams**, if applicable.
 
 ## CLI Semantics
 
 - `kernbench run --device <id>` runs the benchmark on a single device.
 - Omitting `--device` runs the benchmark on all devices discovered in the topology (logically parallel).
 - Device enumeration is handled by the CLI only; benchmarks MUST remain single-device.
+- **Eval-bench exception (ADR-0054)**: a *milestone / eval bench*
+  (`milestone-1h-*`) may drive many configurations and build its own
+  per-config engines to regenerate a domain's full result + figure set; it
+  ignores `--device` and submits a sentinel tensor to satisfy the
+  "must submit ≥1 request" contract (ADR-0045 D4). This is the eval-harness
+  carve-out to the single-device rule, alongside the ADR-0024 multi-SIP CCL
+  exception.
 
 ## Derived Artifacts (Clarification)
 
 - Generated diagrams under `docs/diagrams/` are **derived artifacts**, not production code.
-- Creating or updating files in `docs/diagrams/`:
+- Korean ADR translations under `docs/adr-ko/` are **derived artifacts**
+  (mirror of the canonical English in `docs/adr/`); see *ADR Translation
+  Discipline*.
+- Creating or updating files in `docs/diagrams/` or `docs/adr-ko/`:
   - does NOT count as a production code change,
   - does NOT require Phase 2 approval,
   - MUST be consistent with SPEC.md and ADRs.
 
-## Enforcement Defaults
+## ADR Translation Discipline
+
+English in `docs/adr/` is the canonical source of truth. Korean in
+`docs/adr-ko/` mirrors it 1:1 as a derived artifact.
+
+**Bidirectional sync rule (MUST)**: any edit to a file in `docs/adr/`
+must be accompanied, in the same change, by a mirroring edit to
+`docs/adr-ko/<same-filename>.md`. The reverse also applies: edits to
+`docs/adr-ko/` must mirror back into `docs/adr/`. The two files must
+always describe the same architectural content.
+
+Mechanics:
+
+- When editing an EN ADR, propagate the change to its KO counterpart
+  by translating just the diff (preserve unaffected KO prose); do not
+  regenerate the whole KO file from scratch.
+- When editing a KO ADR, propagate to EN the same way.
+- Filename mirror: `docs/adr/X.md` ↔ `docs/adr-ko/X.md` (no language
+  suffix in either path).
+- The `## Status` *lifecycle keyword* (`Accepted`, `Proposed`,
+  `Stub (Future Work)`, `Draft`, `Superseded by ADR-NNNN`,
+  `Merged into ADR-NNNN`) must match between EN and KO. Parenthetical
+  commentary and any list items that follow the keyword may be
+  translated naturally (the verify tool ignores them when comparing).
+- Conflict policy: if the two diverge despite the rule, treat EN as
+  authoritative and overwrite KO. Surface the divergence to the user
+  before reconciling.
+- `docs/adr-proposed/` is exempt — single language only, no mirror
+  required until promotion.
+- `docs/adr-history/` is frozen — pre-existing mixed-language state
+  there is not migrated.
+
+Verification: `python tools/verify_adr_lang_pairs.py` checks that
+every EN ADR has a matching KO file, the title's ADR-NNNN matches the
+filename, and Status blocks are byte-equal. Run it on demand or wire
+it into CI. Exit code: 0 = OK, 1 = mismatch.
+
+## runtime API / sim_engine Boundaries
 
-- If unsure whether a change is non-trivial → treat it as non-trivial.
-- If unsure whether Phase 2 is allowed → STOP and ask.
-- SPEC.md and ADRs are the final authority.
 - runtime API MUST NOT hardcode topology/routing or internal hop sequences.
 - sim_engine MUST remain independent of runtime API semantics (no tensor/kernel policy logic).

README.md

@@ -155,5 +155,6 @@ kernbench/
 ## Documentation
 
 - [CHANGES.md](CHANGES.md) — changelog with detailed descriptions of each release
-- [docs/latency-model.md](docs/latency-model.md) — latency model explanation with worked examples
+- [docs/onboarding/latency-model.md](docs/onboarding/latency-model.md) — latency model explanation with worked examples
+- [docs/onboarding/](docs/onboarding/) — onboarding guides (architecture overview, latency model, CCL author guide, intro presentation)
 - [docs/adr/](docs/adr/) — Architecture Decision Records

SPEC.md

@@ -51,8 +51,8 @@ Major architectural decisions are documented in ADRs and referenced by number.
 - ADR-0007: runtime_api vs sim_engine responsibility boundaries
 - ADR-0008: Tensor deployment and allocation (Host allocator, PA-first)
 - ADR-0009: Kernel execution fan-out and completion semantics
-- ADR-0010: CLI device selection and multi-device execution semantics
-- ADR-0011: Memory addressing simplification (PA-first)
+- ADR-0010: Command line interface and execution semantics
+- ADR-0011: Memory Addressing — PA / VA / LA Address Models
 - ADR-0012: Host ↔ IO_CPU message schema (PA-first, PE-tagged shards)
 - ADR-0013: Verification strategy and Phase 1 test plan
 - ADR-0014: PE internal execution model (PE_CPU, PE_SCHEDULER, composite commands)
@@ -204,15 +204,23 @@ benchmark instances by default.
 
 ---
 
-## R10. Memory Addressing (Phase 0)
+## R10. Memory Addressing
 
-The simulator uses a **VA/PA memory model** (ADR-0011):
+The simulator defines three address models in ADR-0011; one is selected
+per simulation configuration:
+
+- **PA (Physical Address)** — direct PA, retained as PageFault fallback.
+- **VA (Virtual Address with MMU)** — currently implemented default.
+- **LA (Logical Address with BAAW)** — proposed, supports per-channel
+  HBM modelling (1:1 / n:1 mapping modes).
+
+VA model details (current default):
 
 - Tensors are assigned a contiguous virtual address (VA) range at deployment.
 - PE_MMU translates VA→PA per access; TLB overhead is configurable.
 - Mapping installation (MmuMapMsg) traverses the fabric with measured latency.
 - Replicate tensors use per-cube local PA mapping; sharded tensors broadcast.
-- PA-only fallback is retained for backward compatibility.
+- PA fallback is retained for backward compatibility.
 - Tensor placement is represented as a list of PA shards, each explicitly tagged
   with `(sip, cube, pe)`, plus a tensor-wide `va_base`.
 

benches/ipcq_allreduce.py

@@ -1,2 +0,0 @@
-def run(torch):
-    print("IPCQ all reduce kernel bench")

benches/loader.py

@@ -1,40 +0,0 @@
-from __future__ import annotations
-
-import importlib
-from collections.abc import Callable
-from typing import Any
-
-from kernbench.runtime_api.context import RuntimeContext
-
-BenchFn = Callable[[RuntimeContext], Any]
-
-
-def _load_module(bench_id: str):
-    bench_id = bench_id.strip()
-    if not bench_id:
-        raise ValueError("Bench id is empty.")
-    module_path = f"benches.{bench_id}"
-    try:
-        return importlib.import_module(module_path)
-    except ModuleNotFoundError as e:
-        raise ValueError(
-            f"Unknown bench '{bench_id}'. Expected module {module_path}.py"
-        ) from e
-
-
-def resolve_bench(bench_id: str) -> BenchFn:
-    """Resolve a bench id into its ``run(torch)`` callable.
-
-    Expected layout (repo root):
-        benches/<bench_id>.py
-            def run(torch: RuntimeContext) -> Any
-    """
-    mod = _load_module(bench_id)
-    run_fn = getattr(mod, "run", None)
-    if run_fn is None:
-        raise ValueError(
-            f"Bench module benches.{bench_id} must define 'run(torch)'."
-        )
-    if not callable(run_fn):
-        raise ValueError(f"'run' in benches.{bench_id} is not callable.")
-    return run_fn

ccl.yaml

@@ -6,7 +6,7 @@
 
 defaults:
   # Algorithm to run for this benchmark execution.
-  algorithm: intercube_allreduce
+  algorithm: lrab_hierarchical_allreduce
 
   # IPCQ ring buffer location.
   #   tcm  — PE-local TCM (fast, small, conflicts with compute TCM access)
@@ -37,9 +37,14 @@ algorithms:
   # exchange on root cube, then broadcast back. SIP topology is read
   # from topology.yaml → system.sips.topology. Kernel auto-selects
   # ring / torus / mesh inter-SIP exchange pattern.
-  intercube_allreduce:
-    module: kernbench.ccl.algorithms.intercube_allreduce
+  lrab_hierarchical_allreduce:
+    module: kernbench.ccl.algorithms.lrab_hierarchical_allreduce
     topology: none
     buffer_kind: tcm
     n_elem: 8
+    # root_cube: the kernel currently elects the root dynamically as the
+    # geometric center of the cube mesh (root = (h//2)*w + (w//2)) to
+    # minimize the intra-SIP critical path, so this value is NOT read today.
+    # Kept as a placeholder for a future explicit-root override / runtime
+    # election hook (see ADR-0032 D1 + Non-goals).
     root_cube: 15

docs/adr-history/ADR-0018-mem-logical-address.en.md

@@ -2,7 +2,7 @@
 
 ## Status
 
-Proposed
+Merged into ADR-0011 (Address Model: LA section).
 
 ## Context
 

docs/adr-history/ADR-0018-mem-logical-address.md

@@ -2,7 +2,7 @@
 
 ## Status
 
-Proposed
+Merged into ADR-0011 (Address Model: LA section).
 
 ## Context
 

docs/adr-history/ADR-0019-dev-noc-local-hbm.en.md

@@ -0,0 +1,5 @@
+# ADR-0019: Per-Channel and Aggregated HBM Connection Models within CUBE NOC
+
+## Status
+
+Merged into ADR-0017 (Cube NOC and HBM Connectivity).

docs/adr-history/ADR-0019-dev-noc-local-hbm.md

@@ -0,0 +1,5 @@
+# ADR-0019: CUBE NOC 내 Per-Channel 및 Aggregated HBM 연결 모델
+
+## Status
+
+Merged into ADR-0017 (Cube NOC and HBM Connectivity).

docs/adr-history/ADR-0021-dev-pe-pipeline-refactor.en.md

@@ -0,0 +1,5 @@
+# ADR-0021: PE Pipeline Refactoring — Component Separation + Scheduler-Based Routing
+
+## Status
+
+Merged into ADR-0014 (PE Pipeline Execution Model).

docs/adr-history/ADR-0021-dev-pe-pipeline-refactor.md

@@ -0,0 +1,5 @@
+# ADR-0021: PE 파이프라인 리팩토링 — 컴포넌트 분리 + Scheduler 기반 라우팅
+
+## Status
+
+Merged into ADR-0014 (PE Pipeline Execution Model).

docs/adr-history/ADR-0029-algo-hierarchical-allreduce.md

@@ -2,7 +2,9 @@
 
 ## Status
 
-Proposed
+Superseded by ADR-0032 (Intercube all-reduce). The 3-level kernel and
+`hierarchical_allreduce.py` module have been removed. The cube-mesh
+intercube + inter-SIP path is now the single all-reduce algorithm.
 
 ## Context
 

docs/adr-history/ADR-0031-mem-physaddr-pe-resource-extension.md

@@ -2,7 +2,11 @@
 
 ## Status
 
-Stub (Blocker for ADR-0030 — specific range allocations TBD)
+Superseded by ADR-0001 (Revision 2, 2026-04-27).
+PE_LOCAL / MCPU_LOCAL / CUBE_SRAM sub-unit tables are now defined in
+ADR-0001 D2.3.3-D2.3.5.
+
+Previous status: Stub (Blocker for ADR-0030 — specific range allocations TBD)
 
 ## Context
 
@@ -253,5 +257,5 @@ PhysAddr encoding. 호출자는 어느 range인지 몰라도 됨.
 |------|--------|
 | `src/kernbench/policy/address/phyaddr.py` | Range table (`PE_RESOURCE_MAP`), range-based decode, 신규 component-specific factory들 (`pe_ipcq_addr` 등), 기존 `pe_tcm_addr` 내부 인코딩 갱신 |
 | `src/kernbench/policy/address/allocator.py` | Range-aware pool 분리 (TCM pool / IPCQ pool / scratchpad pool 등 per-PE) |
-| `docs/adr/ADR-0001-physaddr-layout.md` | Amendment note: range-based PE resource partition |
+| `docs/adr/ADR-0001-mem-physaddr-layout.md` | Amendment note: range-based PE resource partition |
 | `tests/test_phyaddr.py` | Range table 검증, 각 factory의 encode/decode round-trip, 기존 `pe_tcm_addr` 회귀 |

docs/adr-ko/ADR-0001-mem-physaddr-layout.md

@@ -0,0 +1,358 @@
+# ADR-0001: 51비트 물리 주소 레이아웃 및 디코딩 계약
+
+## Status
+
+Accepted (Revision 2 — 2026-04-27: 구체적인 비트 레이아웃, rack_id 제거,
+Tray->SIP / SIP->DIE 명칭 변경, PE/MCPU/IOCPU 서브 유닛 표.
+ADR-0031을 대체함.)
+
+## Date
+
+2026-04-27 (original: 2026-02-27)
+
+## Context
+
+KernBench에는 다음과 같은 요건을 만족하는 안정적이고 파싱 가능한 물리 주소 체계가 필요하다.
+
+- 라우팅 도메인(SIP / die / HBM / PE-resource / IOCPU)으로 디코딩 가능
+- 토폴로지에 비의존적(개수를 하드코딩하지 않음)
+- 교체 가능한 정책과 DI-first 컴포넌트를 지원
+- 다수의 SIP, AHBM die, IO chiplet die를 통합된 공간에서 다룸
+
+### 연혁
+
+- 최초 ADR-0001은 `rack_id(4) + sip_id(4) + sip_seg(5) + local_offset(38)`
+  로 구성된 51비트 레이아웃을 정의했다. `rack_id`는 실제로 사용된 적이 없다.
+- ADR-0031(스텁)은 PE-resource 범위 분할을 요청했으나 구현되지 않았다.
+
+Revision 2에서는 `rack_id`를 제거하고 `sip_seg`를 `die_id`로 개명하며,
+PE, MCPU, CUBE_SRAM, IOCPU 리소스에 대한 구체적인 서브 유닛 표를 제공한다.
+ADR-0031은 본 ADR로 대체된다.
+
+## Decision
+
+**PhysAddr 값 객체**와, 정수 주소를 라우팅 도메인으로 변환하는
+**주소 디코딩 계약**을 정의한다.
+
+### D1. PhysAddr는 불변 값 객체이다
+
+- PhysAddr는 불변이며 순수한 값으로 비교 가능하다.
+- 모든 할당자는 **완전히 명세된 PhysAddr**(부분적인 메타데이터가 아님)를 반환한다.
+- PhysAddr를 해석하기 위해 전역 상태를 필요로 해서는 안 된다.
+
+### D2. 51비트 물리 주소 레이아웃
+
+51비트 물리 주소를 채택한다.
+
+#### 2.1 최상위 주소 맵
+
+```text
+[50:47] sip_id        (4)     -- 16 SIPs
+[46:42] die_id        (5)     -- 32 dies per SIP
+[41: 0] local_offset  (42)    -- 4 TB per die
+```
+
+```text
+50      47 46      42 41                      0
++---------+----------+-------------------------+
+| sip_id  | die_id   |      local_offset       |
++---------+----------+-------------------------+
+```
+
+#### 2.2 die_id 할당
+
+| die_id | 의미 |
+|--------|---------|
+| 0..15  | AHBM dies |
+| 16..20 | IOCHIPLET dies |
+| 21..31 | Reserved |
+
+#### 2.3 AHBM Die 레이아웃
+
+4 TB die-local 윈도우 중 하위 256 GB만 할당된다.
+
+```text
+[41:38] MBZ            (4)
+[37]    addr_space      (1)    -- 0 = local resource, 1 = HBM memory
+[36: 0] sub-address    (37)
+```
+
+| addr_space | 의미 |
+|------------|---------|
+| 0 | Local resource |
+| 1 | HBM memory |
+
+##### 2.3.1 HBM 윈도우 (addr_space = 1)
+
+```text
+[36:0] hbm_offset     (37)    -- 128 GB decode window
+```
+
+아키텍처상의 디코드 윈도우는 128 GB로 고정된다. 실제 구현 용량은
+SKU/토폴로지에 따라 더 작을 수 있다(D4 참조).
+
+##### 2.3.2 Resource 윈도우 (addr_space = 0)
+
+```text
+[36:34] resource_kind  (3)
+[33: 0] kind_local    (34)    -- 16 GB per kind
+```
+
+| resource_kind | 의미 |
+|---------------|---------|
+| 000 | PE_LOCAL |
+| 001 | MCPU_LOCAL |
+| 010 | CUBE_SRAM |
+| 011..111 | Reserved |
+
+각 kind는 16 GB 디코드 영역을 갖는다.
+
+##### 2.3.3 PE_LOCAL (resource_kind = 000)
+
+```text
+[33]    MBZ            (1)
+[32:29] pe_id          (4)     -- 0..15
+[28:25] pe_sub_unit    (4)
+[24: 0] sub_offset    (25)    -- 32 MB per slot
+```
+
+16 PE x 16 서브 유닛 슬롯 x 32 MB = 8 GB 활성 디코드.
+
+| pe_sub_unit | 이름 | 예산 |
+|-------------|------|--------|
+| 0 | PE_CPU_DTCM | 8 KB |
+| 1 | MATH_ENGINE_DTCM | 8 KB |
+| 2 | IPCQ | 256 KB |
+| 3 | PE_CPU_SFR | 16 KB |
+| 4 | MATH_ENGINE_SFR | 16 KB |
+| 5 | DMA_ENGINE_SFR | 192 KB |
+| 6 | PE_TCM | 2 MB |
+| 7..15 | Reserved | -- |
+
+##### 2.3.4 MCPU_LOCAL (resource_kind = 001)
+
+```text
+[33:30] MBZ            (4)
+[29:25] mcpu_sub_unit  (5)
+[24: 0] sub_offset    (25)    -- 32 MB per slot
+```
+
+1 GB 활성 디코드.
+
+| mcpu_sub_unit | 이름 | 예산 |
+|---------------|------|--------|
+| 0 | MCPU_ITCM | 512 KB |
+| 1 | MCPU_DTCM | 512 KB |
+| 2 | IPCQ | 256 KB |
+| 3 | MCPU_SFR | 8 KB |
+| 4 | MCPU_DMA_SFR | 16 KB |
+| 5 | MCPU_SRAM | 10 MB |
+| 6..31 | Reserved | -- |
+
+##### 2.3.5 CUBE_SRAM (resource_kind = 010)
+
+```text
+[33:25] MBZ            (9)
+[24: 0] sram_offset   (25)    -- flat 32 MB
+```
+
+#### 2.4 IOCHIPLET Die 레이아웃
+
+4 TB die-local 윈도우 중 하위 1 TB만 할당된다.
+
+```text
+[41:40] MBZ            (2)
+[39: 0] chiplet_offset (40)   -- 1 TB
+```
+
+주소 범위별 영역 구분:
+
+| 범위 | 의미 | 디코드 조건 |
+|-------|---------|------------------|
+| [0, 2 GB) | IOCPU resource | chiplet_offset < 0x8000_0000 |
+| [2 GB, 1 TB) | UAL | chiplet_offset >= 0x8000_0000 |
+
+##### 2.4.1 IOCPU 영역
+
+```text
+[30:27] iocpu_sub_unit (4)
+[26: 0] sub_offset    (27)    -- 128 MB per slot
+```
+
+16 x 128 MB 슬롯. 2 GB 활성 디코드.
+
+| iocpu_sub_unit | 이름 | 예산 |
+|----------------|------|--------|
+| 0 | IOCPU_ITCM | 512 KB |
+| 1 | IOCPU_DTCM | 512 KB |
+| 2 | IPCQ | 2 MB |
+| 3 | IOCPU_SFR | 8 KB |
+| 4 | IO_DMA_SFR | 16 KB |
+| 5 | IO_SRAM | 64 MB |
+| 6..15 | Reserved | -- |
+
+##### 2.4.2 UAL 영역
+
+서브 레이아웃은 별도 ADR에서 정의한다(TBD).
+
+#### 2.5 주소 지정 규칙
+
+1. MBZ 비트는 반드시 0이어야 한다. MBZ 비트가 0이 아닌 주소는
+   **아키텍처적으로 유효하지 않다**. 구현체는 디코드 폴트를 발생시키거나
+   오류를 반환할 수 있다 — 본 ADR은 동작을 규정하지 않는다.
+2. 단순한 하드웨어 디코드를 위해 고정된 슬롯 크기를 채택한다. 실제 구현
+   용량은 슬롯보다 작을 수 있다.
+3. 슬롯 내에서 서브 유닛의 구현 예산을 초과하는 접근은 **아키텍처적으로
+   유효하지 않다**(MBZ와 동일한 정책).
+
+### D3. 비트필드 디코딩은 결정론적이다
+
+정수 주소가 주어지면 필드 추출(`sip_id`, `die_id`, `kind`, `sub_unit`,
+`offset`)은 순수하게 위치 기반이다. 런타임 상태가 필요하지 않다.
+디코딩은 정수 주소를 결정론적으로 목적지 도메인(`sip_id`, `die_id`,
+타깃 종류 HBM / PE_LOCAL / MCPU_LOCAL / CUBE_SRAM / IOCPU / UAL)으로 매핑한다.
+
+### D4. 용량 검증은 토폴로지 설정에 의존할 수 있다
+
+디코딩된 주소가 **구현된 용량** 안에 들어가는지(예: 특정 SKU의 HBM 96 GB)는
+DI/설정을 통해 제공된 토폴로지 파라미터로 검증한다. 디코딩 자체(D3)는
+토폴로지를 참조하지 않으며 — 검증 단계에서만 참조한다. 이러한 파라미터는
+컴포넌트 구현이 아니라 토폴로지/설정 레이어에 존재해야 한다.
+
+### D5. 라우팅은 원시 비트가 아닌 디코딩된 도메인을 소비한다
+
+라우팅 정책은 디코딩된 도메인을 사용한다.
+
+- `src` 위치 (sip / die / pe 또는 node_id)
+- PhysAddr 디코딩에서 도출된 `dst` 도메인
+- 크기 인지 링크 레이턴시를 위한 `size_bytes`
+
+라우팅은 디코딩 모듈 내부를 제외하고는 원시 비트필드를 직접 들여다보아서는
+안 된다.
+
+## 고려된 대안
+
+1. **`rack_id`(4비트) 유지**: 기각 — 실제로 사용된 적이 없으며, 4비트를
+   소비함으로써 die-local 확장을 42비트(IOCHIPLET 1 TB)까지 가능하게 하는
+   기회를 막는다.
+
+2. **die당 256 GB로 균일화**: 기각 — IOCHIPLET UAL은 약 1 TB가 필요하다.
+   해제된 rack_id 비트를 활용하여 42비트 local_offset을 가능하게 한다.
+
+3. **가변 폭 die 윈도우(AHBM 256 GB, CHIPLET 1 TB를 다중 seg 스패닝으로 구현)**:
+   기각 — D3(결정론적 디코딩)를 복잡하게 만든다. MBZ 패딩을 갖는 균일한
+   4 TB 윈도우가 더 단순하다.
+
+4. **모든 곳에서 원시 정수를 사용하고, 라우팅에서 임시로 디코딩**: 기각 —
+   로직이 중복되고 라우팅이 일관성을 잃으며 가정이 숨겨진다.
+
+5. **토폴로지 크기(SIP/CUBE/PE 개수)를 디코딩에 하드코딩**: 기각 —
+   SPEC R3를 위반하고 교체 가능성을 깬다.
+
+6. **디코딩을 메모리 컨트롤러나 라우터 내부에 둠**: 기각 — 정책이 컴포넌트로
+   누출되며 SPEC R4 / D5를 위반한다.
+
+## 결과
+
+### 긍정적
+
+- 단순한 계층적 디코더: SIP -> die -> kind -> 서브 유닛.
+- 메모리(HBM)와 로컬 리소스(PE/MCPU/SRAM/IOCPU)의 깔끔한 분리.
+- 결정론적 라우팅 도메인은 명확한 테스트 불변식을 가능하게 한다(SPEC R1, R5).
+- 확장 가능: 11개의 예약된 die_id 슬롯, 예약된 resource_kind / 서브 유닛
+  슬롯, 예약된 MBZ 비트.
+- DI-first: 컴포넌트를 변경하지 않고도 디코더를 교체할 수 있다(SPEC R4).
+
+### 트레이드오프
+
+- power-of-2 슬롯 정렬로 인한 희소한 주소 공백.
+- 큰 예약/MBZ 영역(향후 확장을 위해 의도된 것).
+- 토폴로지에서 유도된 크기에 대해 명시적인 설정이 필요하다(D4).
+- 안정적이고 잘 테스트된 상태로 유지되어야 하는 단일 "정통" 디코딩 모듈이
+  도입된다.
+
+## 대체 대상
+
+- **ADR-0031 (PhysAddr PE-Resource Extension)**: 스텁 상태였음. D2.3.3-D2.3.5의
+  PE_LOCAL / MCPU_LOCAL / CUBE_SRAM 서브 유닛 표가 ADR-0031에서 제시한
+  목표를 충족한다.
+
+## 구현 메모 (비규범적)
+
+- 권장 모듈: `src/kernbench/policy/address/phyaddr.py`
+- 테스트는 다음을 커버해야 한다: kind별 인코딩/디코딩 라운드트립, MBZ 강제,
+  die_id 디스패치(AHBM / IOCHIPLET / 예약), 서브 유닛 경계값, 팩토리 API의
+  후방 호환성.
+- 팩토리 메서드: `hbm_addr`, `pe_hbm_addr`, `pe_tcm_addr`, `cube_sram_addr`는
+  시그니처를 유지한다(`rack_id` 제외). `cube_id` 파라미터는 `die_id`로
+  개명된다.
+- 신규 팩토리: `pe_resource_addr`, `mcpu_resource_addr`, `iocpu_resource_addr`,
+  `ual_addr`.
+
+## 부록 A. 주소 예시
+
+### A.1 AHBM HBM 접근
+
+sip=2, die=5, HBM offset=0x1000
+
+```text
+sip_id     = 2       -> [50:47] = 0b0010
+die_id     = 5       -> [46:42] = 0b00101
+addr_space = 1       -> [37]    = 1 (HBM)
+hbm_offset = 0x1000  -> [36:0]
+
+51-bit addr = (2 << 47) | (5 << 42) | (1 << 37) | 0x1000
+```
+
+### A.2 AHBM PE_LOCAL -- PE3 PE_TCM, offset=0x400
+
+```text
+sip_id        = 0  -> [50:47] = 0
+die_id        = 0  -> [46:42] = 0
+addr_space    = 0  -> [37]    = 0
+resource_kind = 0  -> [36:34] = 000 (PE_LOCAL)
+pe_id         = 3  -> [32:29] = 0011
+pe_sub_unit   = 6  -> [28:25] = 0110 (PE_TCM)
+sub_offset    = 0x400 -> [24:0]
+
+local_offset = (0 << 34) | (3 << 29) | (6 << 25) | 0x400
+```
+
+### A.3 AHBM MCPU_LOCAL -- MCPU_SRAM, offset=0x0
+
+```text
+sip_id        = 1  -> [50:47] = 0001
+die_id        = 3  -> [46:42] = 00011
+addr_space    = 0  -> [37]    = 0
+resource_kind = 1  -> [36:34] = 001 (MCPU_LOCAL)
+mcpu_sub_unit = 5  -> [29:25] = 00101 (MCPU_SRAM)
+sub_offset    = 0  -> [24:0]  = 0
+
+local_offset = (1 << 34) | (5 << 25)
+```
+
+### A.4 IOCHIPLET -- IOCPU IPCQ, offset=0x20000
+
+```text
+sip_id         = 1   -> [50:47] = 0001
+die_id         = 17  -> [46:42] = 10001 (IOCHIPLET[1])
+iocpu_sub_unit = 2   -> [30:27] = 0010 (IPCQ)
+sub_offset     = 0x20000 -> [26:0]
+
+chiplet_offset = (2 << 27) | 0x20000
+                 (< 0x8000_0000 -> IOCPU region)
+```
+
+### A.5 IOCHIPLET -- UAL 영역, offset=4 GB
+
+```text
+sip_id         = 0   -> [50:47] = 0
+die_id         = 16  -> [46:42] = 10000 (IOCHIPLET[0])
+chiplet_offset = 0x1_0000_0000 (4 GB >= 2 GB -> UAL region)
+```
+
+## 링크
+
+- SPEC.md: R1 (routing), R3 (configurable topology), R4 (DI-first),
+  R5 (multi-domain comm)
+- ADR-0031: Superseded

docs/adr-ko/ADR-0002-lat-routing-distance.md

@@ -0,0 +1,100 @@
+# ADR-0002: 라우팅 거리, 순서 및 우회 규칙
+
+## Status
+Accepted
+
+## Date
+2026-02-27
+
+## Context
+KernBench Graph Latency Simulator는 서로 다른 아키텍처·토폴로지에 대한
+커널 실행 시간을 비교해야 하며, 그래프 순회로부터 end-to-end 레이턴시를
+계산하여 이를 달성한다.
+
+의미 있는 비교를 지원하려면:
+- 라우팅이 결정론적이어야 한다
+- 레이턴시가 실제 인터커넥트 구조를 반영해야 한다
+- 로컬과 리모트 트래픽이 구분 가능해야 한다
+- "우회(bypass)" 최적화가 디버깅 가능성이나 정확성을 훼손해서는 안 된다
+
+또한 시뮬레이터는 소프트웨어가 관리하는 메타데이터 및 제어 경로를
+가리는 숨겨진 지름길을 피하는 것을 목표로 한다.
+
+## Decision
+
+### D1. 거리(distance)는 hop 수가 아니라 누적 레이턴시이다
+- 라우팅 "거리"는 **노드별·링크별 레이턴시의 합**으로 정의된다.
+- 순서 결정이나 경로 선택에 hop 수만을 사용해서는 안 된다.
+- 크기 인지(size-aware) 직렬화 레이턴시(bytes / BW)가 거리에 기여한다.
+
+### D2. 라우팅 순서는 그래프 순회에서 유도된다
+- 선택된 경로는 구성된 그래프와 라우팅 정책 하에서
+  누적 레이턴시가 최소인 경로이다.
+- 동일 입력(토폴로지 + 정책 + 요청)에 대해 결정론적 순서가 보장되어야 한다.
+
+### D3. 우회는 명시적이며 그래프로 표현된다
+- 모든 경로는 그래프에 명시적으로 표현되며 레이턴시 누적의 대상이 되어야 한다.
+- 예: PE_DMA는 NOC 라우터 메시(ADR-0017 D7)에 연결된다. 모든 목적지
+  (HBM, 공유 SRAM, 큐브 간 UCIe)는 명시적 메시 hop을 통해 도달한다.
+  로컬 HBM 접근은 hop 수가 최소(스위칭 오버헤드만)이며, 리모트 접근은
+  추가 라우터를 거친다.
+- 암묵적이거나 "마법 같은" 우회 경로는 금지된다.
+
+### D4. end-to-end 레이턴시가 0인 경로는 없다
+
+- 모든 라우팅 요청은 **end-to-end** 레이턴시가 > 0이어야 한다.
+- 개별 패브릭 세그먼트(예: NOC hop)는 패브릭이 분산되어 있고 해당 granularity에서
+  거리가 의미가 없을 때 distance_mm = 0을 가질 수 있다.
+  이는 같은 경로상의 다른 컴포넌트(예: PE_DMA, SRAM, UCIe 엔드포인트)가
+  0이 아닌 레이턴시에 기여하여 end-to-end 불변성을 유지하므로 허용된다.
+- end-to-end가 완전히 0 레이턴시인 경로는 금지된다. 단, 명시적으로
+  표시된 테스트 전용 stub만 예외이다.
+
+### D5. 정책과 토폴로지의 책임 분리
+- 토폴로지 빌더:
+  - 노드와 링크 및 그들의 레이턴시/BW 파라미터를 정의한다
+- 라우팅 정책:
+  - 디코딩된 도메인을 바탕으로 사용 가능한 그래프 경로 중에서 선택한다
+- 라우팅 정책은 누락된 링크를 가정해서는 안 된다. 누락된 연결성은
+  토폴로지 구성 오류이다.
+
+### D6. 소프트웨어 관리 라우팅 메타데이터 금지
+- 라우팅 결정은 그래프 모델 외부에서 거리·hop 수·순서를 추적하는
+  요청별 소프트웨어 관리 메타데이터에 의존해서는 안 된다.
+- 모든 거리·순서 계산은 순회 자체에서 유도된다.
+
+## Alternatives Considered
+
+1) **Hop 수 기반 라우팅**
+- 기각: 이질적인 레이턴시·BW를 무시하고 아키텍처 차이를 잘못 표현한다.
+
+2) **암묵적 로컬 지름길**
+- 기각: 디버깅 가능성을 해치고 순회 기반 레이턴시 원칙을 위반한다.
+
+3) **소프트웨어 관리 거리 메타데이터**
+- 기각: 제어 오버헤드를 증가시키고 라우팅 시맨틱을 모호하게 만든다.
+
+## Consequences
+
+### 긍정적
+- 명확하고 디버깅 가능한 hop-by-hop 트레이스 (SPEC R2, R4).
+- 아키텍처 비교가 실제 인터커넥트 구조를 반영한다.
+- 라우팅 동작이 재현 가능하고 결정론적이다.
+
+### 트레이드오프 / 비용
+- 그래프 구성이 정확하고 완전해야 한다.
+- 우회 모델링이 명시적 그래프 표현을 요구하므로 토폴로지 기술이
+  약간 더 복잡해진다.
+
+## Implementation Notes (Non-normative)
+- 권장 책임 분담:
+  - 그래프 빌더: 필요한 모든 경로가 존재함을 보장.
+  - 라우터: 디코딩된 도메인과 정책을 바탕으로 다음 hop 선택.
+- 테스트가 검증해야 할 항목:
+  - end-to-end 레이턴시 > 0
+  - 동일 입력에 대한 결정론적 라우팅
+  - 우회 경로가 출력 트레이스에 명시적으로 나타남
+
+## Links
+- SPEC.md: R1 (라우팅), R2 (레이턴시), R3 (토폴로지), R5 (다중 도메인 통신)
+- ADR-0001: PhysAddr 레이아웃 및 디코딩 계약

docs/adr-ko/ADR-0003-dev-target-system-hierarchy.md

@@ -0,0 +1,68 @@
+# ADR-0003: 타겟 시스템 계층 및 모델링 범위
+
+## Status
+
+Accepted
+
+## Context
+
+자사 AI Accelerator 플랫폼에서 LLM 커널 성능을 평가하기 위해 시스템 수준의 시뮬레이터가 필요하다.
+해당 플랫폼은 PCIe 또는 UAL을 통해 스위칭 패브릭으로 연결된 다수의 동일한 SIP를 포함하는 컴퓨트 트레이로 구성되며,
+호스트 CPU가 명령/커널을 발급한다.
+
+## Decision
+
+시스템 계층을 다음과 같이 명시적으로 모델링한다.
+
+### D1. Tray-level
+
+- 하나의 컴퓨트 트레이는 다음을 포함한다:
+  - 호스트 CPU (요청 발급 / 런타임 및 데이터 배치 조정)
+  - 다수의 동일한 SIP (가속기)
+  - SIP 간 인터커넥트 패브릭 (스위치를 통한 PCIe 및/또는 UAL)
+
+### D2. SIP-level
+
+- SIP는 다음으로 구성된 멀티 다이 패키지이다:
+  - 다수의 CUBE (HBM 다이 + 컴퓨트 PE + UCIe)
+  - 하나 이상의 IO 칩렛 (호스트/SIP 인터페이스)
+- IO 칩렛:
+  - 다음 인터페이스를 제공한다: PCIe-EP, IO_CPU, 선택적으로 UAL-EP
+  - SIP 당 다수가 존재할 수 있다
+  - 배치는 SIP shoreline(상/하/좌/우)으로 제약되며, 각 shoreline에는 1~2개의 IO 칩렛이 위치할 수 있다
+
+### D3. CUBE-level
+
+- 하나의 CUBE는 다음을 포함한다:
+  - HBM + 메모리 컨트롤러 (HBM_CTRL)
+  - NoC (on-die 패브릭): HBM 데이터, 큐브 간(UCIe) 트래픽, 명령(M_CPU↔PE_CPU),
+    공유 SRAM 액세스를 포함한 모든 큐브 내부 트래픽을 운반한다.
+    반드시 제공해야 하는 것: 풀-대역폭 PE↔로컬 HBM 경로, PE↔SRAM 연결성,
+    PE↔UCIe 연결성, M_CPU↔PE 명령 경로.
+    NoC 토폴로지는 구현 선택사항(예: 2D 메시, 링, 크로스바)이며,
+    현재 구현은 XY 라우팅 방식의 2D 메시를 사용한다(ADR-0017 참조).
+    HBM_CTRL은 각 PE의 로컬 NoC 포트에 부착된다(로컬 HBM = 최소 홉).
+  - 공유 SRAM: 모든 PE가 NoC를 통해 액세스 가능한 큐브 수준 공유 메모리
+  - PE 명령 분배 및 완료 집계를 조정하는 관리/제어 CPU (M_CPU)
+  - 다수의 PE
+  - CUBE↔CUBE 및 CUBE↔IO 연결성을 위한 최대 4개의 UCIe 엔드포인트 (N/E/W/S)
+
+### D4. PE-level
+
+- 하나의 PE는 하나의 커널 인스턴스를 실행할 수 있다
+- PE는 내부 제어 + 가속기를 포함한다 (PE 뷰 단위로 모델링):
+  - PE_CPU, 명령 핸들러, PE_TCM, DMA/GEMM/MATH 엔진, 내부 큐
+
+## Consequences
+
+- 시뮬레이터는 "뷰" 단위의 추상화를 지원한다:
+  - SIP 뷰는 PE 내부를 숨긴다
+  - CUBE 뷰는 각 PE를 단일 블록으로 다룬다
+  - PE 뷰는 PE 내부를 전개한다
+- 토폴로지는 매개변수화된 상태로 유지되며, 크기/개수/링크는 설정으로부터 주어진다.
+
+## Links
+
+- SPEC R3/R5
+- ADR-0005 (다이어그램 뷰)
+- ADR-0017 (큐브 NoC 2D 메시 아키텍처)

docs/adr-ko/ADR-0004-mem-memory-semantics-local-hbm.md

@@ -0,0 +1,78 @@
+# ADR-0004: 메모리 시맨틱 및 로컬 HBM 대역폭 보장
+
+## Status
+
+Accepted
+
+## Context
+
+PE↔HBM 동작을 정확하게 모델링하는 것은 커널 레이턴시 추정에 필수적이다.
+각 PE는 "로컬 HBM"이라는 개념을 가지며, 이는 중간 온칩 패브릭 대역폭과
+무관하게 HBM 전체 대역폭을 보장해야 한다.
+
+## Decision
+
+### D1. 로컬 HBM의 정의
+
+- 각 PE에는 논리적으로 정의된 "로컬 HBM" 영역이 할당된다.
+- 로컬 HBM은 NOC 메시(ADR-0017 D4) 내에서 해당 PE의 라우터에 직접 연결된
+  pseudo-channel 부분집합에 대응한다.
+- 경로는: PE_DMA → 로컬 라우터 → HBM_CTRL (스위칭 오버헤드만, 메시 hop 0개).
+- 매핑(HBM pseudo-channel → PE 로컬 영역)은 토폴로지 구성에서 유도된다.
+
+### D2. 로컬 HBM 대역폭 보장 계약
+
+- PE에서 자신의 로컬 HBM으로의 접근은 중간 패브릭 대역폭 제한과
+  무관하게 HBM의 유효 read/write 대역폭 전부를 보장해야 한다.
+- 유효 HBM 대역폭 = 스펙 대역폭 × 효율 계수.
+  효율 계수(`hbm_ctrl.attrs.efficiency`로 설정, 기본값 0.8)는 실세계 DRAM의
+  비효율(리프레시 사이클, 뱅크 충돌, 페이지 미스 등)을 모델링한다.
+  예: 256 GB/s 스펙 × 0.8 = 204.8 GB/s 유효 대역폭.
+- 토폴로지 빌더는 그래프 구성 시점에 router-to-hbm 에지의 대역폭에
+  효율 계수를 적용하므로, 이후의 모든 라우팅·레이턴시 계산은 유효 값을
+  사용한다.
+- 이 보장은 다음으로 모델링된다:
+  - PE-로컬-HBM 상호작용 지점에서 HBM 대역폭을 강제하는 전용 논리 경로
+    그리고/또는 서비스 모델,
+  - 명시적으로 모델링된 컴포넌트들을 따라 0이 아닌 레이턴시를 여전히 발생시킨다.
+- HBM CTRL 내부 모델링(PC 스트라이핑, cut-through, 스케줄링 충실도)은
+  ADR-0033 (레이턴시 모델: 가정 및 알려진 단순화)에 통합되어 있다.
+  여기서의 총 대역폭 보장은 계약으로 유지되며, ADR-0033은 PC 단위 모델이
+  이를 어떻게 실현하는지와 어떤 스케줄러 효과가 의도적으로 단순화되었는지를
+  기록한다.
+
+### D3. 리모트 PE HBM 시맨틱 (큐브 내)
+
+- 한 PE가 다른 PE의 로컬 HBM에 접근할 때는 NOC를 거친다:
+  - PE_DMA → NOC → (패브릭 hop) → 대상 PE의 NOC 포트 → HBM_CTRL
+- NOC의 대역폭과 hop 수에 의해 리모트 HBM 접근이 로컬 접근 대비 제한될 수 있다.
+
+### D4. 비로컬 HBM 시맨틱 (큐브 간 / SIP 간)
+
+- PE에서 다른 큐브나 SIP에 있는 HBM으로의 접근은 다음에 의해 제한될 수 있다:
+  - 큐브 내 NOC 대역폭,
+  - 큐브 간 UCIe 링크,
+  - SIP 간 패브릭 (PCIe/UAL).
+- 이 경로들은 명시적이고 추적 가능해야 한다.
+
+### D5. 공유 SRAM 시맨틱
+
+- 각 CUBE는 해당 CUBE의 모든 PE가 접근 가능한 공유 SRAM을 포함한다.
+- 접근 경로: PE_DMA → NOC → 공유 SRAM.
+- 공유 SRAM의 대역폭은 NOC↔SRAM 링크 대역폭으로 제한된다.
+- 공유 SRAM은 HBM 주소 공간의 일부가 아니라 별도의 메모리 도메인이다.
+
+## Verification Notes
+
+테스트가 다뤄야 할 케이스:
+
+- 로컬 HBM 케이스: 패브릭 BW 파라미터와 무관하게 대역폭이 HBM 대역폭과 일치
+- 리모트 PE HBM 케이스: 레이턴시가 메시 hop 순회를 포함
+- 비로컬 케이스(큐브 간/SIP 간): 패브릭/링크 파라미터에 대역폭·레이턴시가 반응
+- 공유 SRAM 케이스: NOC 경유 접근이 올바른 대역폭으로 수행됨
+
+## Links
+
+- SPEC R2/R5
+- ADR-0002 (거리/순서 및 명시적 우회)
+- ADR-0017 D7 (NOC를 통한 PE DMA → HBM 데이터 경로)

docs/adr-ko/ADR-0005-dev-diagram-views-distance-layout.md

@@ -0,0 +1,186 @@
+# ADR-0005: 다이어그램 뷰 및 거리 기반 레이아웃 규칙
+
+## Status
+
+Accepted
+
+## Context
+
+대규모, 매개변수화된 AI Accelerator 시스템에 대해 검증 가능하고 점검 가능한
+시스템 모델링이 필요하다.
+
+사람이 다음을 할 수 있어야 한다:
+
+- 모델링된 토폴로지를 시각적으로 점검하고,
+- 통신 구조와 상대적 거리에 대해 추론하고,
+- 세부 사항에 압도되지 않으면서 여러 추상화 수준에서 이를 수행한다.
+
+시뮬레이터는 거리(누적 레이턴시)를 1급 개념(first-class concept)으로 모델링한다.
+다이어그램은 기본적으로 이 거리를 반영해야 한다.
+
+---
+
+## Decision
+
+### D1. Global Defaults
+
+- 모든 다이어그램은 기본적으로 **거리 인식(distance-aware)** 이어야 한다.
+- 모든 다이어그램은 아키텍처의 **대표 뷰(representative view)** 를 렌더링해야 한다.
+- 인스턴스 인덱스(예: sip0, cube2, pe3)는 다이어그램 생성에 필수가 아니어야 한다.
+- 인스턴스 인덱스는 다음의 경우에만 사용될 수 있다:
+  - 비대칭 또는 디버깅 시나리오에서 거리 앵커를 정의하기 위한 경우, 또는
+  - 명시적으로 요청된 경우.
+
+---
+
+### D2. Representative Rendering Rule
+
+- 모든 CUBE는 동일한 내부 구조를 공유한다.
+- 모든 PE는 동일한 내부 구조를 공유한다.
+
+따라서:
+
+- SIP 수준 다이어그램은 대표 CUBE와 IO 칩렛을 렌더링한다.
+- CUBE 수준 다이어그램은 대표 PE를 불투명 블록으로 렌더링한다.
+- PE 수준 다이어그램은 내부가 완전히 전개된 대표 PE를 렌더링한다.
+
+다이어그램은 명시적으로 요청되지 않는 한
+특정 SIP, CUBE, 또는 PE 인덱스에 의존해서는 안 된다.
+
+---
+
+### D3. Diagram Views
+
+#### View A — SIP 수준 다이어그램
+
+**목적**
+시스템 규모의 구조와 연결성을 설명한다.
+
+**가시 요소**
+
+- SIP 경계 (선택사항)
+- CUBE (불투명 블록)
+- IO 칩렛 (불투명 블록)
+- 연결성 명확화에 필요한 경우에만 선택적 UCIe 스텁
+
+**비가시 요소**
+
+- PE 내부
+- CUBE 내부 패브릭
+- IO 칩렛 내부
+
+**가시 링크**
+
+- 호스트 ↔ IO 칩렛 (PCIe)
+- SIP ↔ SIP (스위치를 통한 PCIe / UAL)
+- IO ↔ CUBE (온패키지 링크)
+
+---
+
+#### View B — CUBE 수준 다이어그램
+
+**목적**
+큐브 내부 구조와 데이터/제어 흐름을 설명한다.
+
+**가시 요소**
+
+- 라우터 메시: NoC 라우터의 2D 격자 (cube_mesh.yaml로부터), 모든 트래픽은 메시를 통해 라우팅됨
+- PE 라우터에 부착된 HBM_CTRL (로컬 HBM = 0 홉)
+- HBM 서브시스템 (HBM_CTRL)
+- 공유 SRAM: 큐브 수준 공유 메모리
+- 관리 CPU (M_CPU)
+- 불투명 블록으로 표현된 PE (PE[0..N−1])
+- 포트로 표현된 UCIe 엔드포인트 (N/E/W/S)
+
+**비가시 요소**
+
+- PE 내부
+
+**가시 링크**
+
+- PE → 라우터 (메시를 통한 HBM + 비-HBM 데이터 경로)
+- 라우터 ↔ HBM_CTRL (로컬 HBM 액세스)
+- 라우터 ↔ 라우터 (원격 액세스를 위한 메시 홉)
+- 라우터 ↔ UCIe 엔드포인트
+- 라우터 ↔ 공유 SRAM
+- M_CPU ↔ 라우터 (명령 경로)
+- 라우터 → PE_CPU (명령 전달, PE 블록 내부로 축약됨)
+
+---
+
+#### View C — PE 수준 다이어그램
+
+**목적**
+PE 내부 동작과 실행 구조를 설명한다.
+
+**가시 요소**
+
+- PE_CPU
+- 명령 핸들러 / 스케줄러
+- PE_TCM (로컬 SRAM)
+- HW 가속기 (DMA, GEMM, MATH 등)
+- 로컬 HBM 인터페이스
+- 선택적 IPCQ / 메시징 엔드포인트
+
+**가시 링크**
+
+- 제어 경로 (CPU → 스케줄러 → 엔진)
+- 데이터 경로 (엔진 ↔ TCM, DMA ↔ 로컬 HBM)
+- 외부 패브릭 포트는 추상 포트로만 표현
+
+---
+
+### D4. 거리 기반 레이아웃 (기본)
+
+#### 거리 정의
+
+- 거리는 ADR-0002와 정합되도록 **누적 레이턴시(accumulated latency)** 로 정의된다.
+- 거리는 단일 앵커 노드로부터 계산된다.
+
+#### 기본 앵커 선택
+
+- SIP 뷰: IO 칩렛 (또는 존재한다면 호스트 CPU)
+- CUBE 뷰: 대표 PE
+- PE 뷰: PE_CPU 또는 명령 핸들러
+
+앵커는 **암묵적 기본값**이며, 지정이 강제되어서는 안 된다.
+
+#### 레이아웃 규칙
+
+- 다이어그램은 거리 버킷에 기반한 레이어로 배치되어야 한다.
+- 레이아웃 방향은 뷰 유형 내에서 일관되어야 한다
+  (선호: 좌→우).
+- 동일 거리의 노드는 결정론적으로 안정된 순서를 가져야 한다
+  (역할 또는 식별자 기준).
+
+가독성을 위해 사이클은 점선 또는 곡선 엣지로 렌더링될 수 있으며,
+이는 거리 의미에 영향을 주지 않는다.
+
+---
+
+### D5. 생성 컨트랙트 (도구 / Claude Code용)
+
+다이어그램 생성 시:
+
+- 기본적으로 거리 기반 레이아웃을 가정한다.
+- 기본적으로 대표 렌더링을 가정한다.
+- 필요한 경우가 아니면 SIP/CUBE/PE 인덱스를 묻지 않는다.
+- 숨겨진 추상화 수준을 전개하지 않는다.
+- 마이크로 홉의 정밀도보다 아키텍처적 명확성을 우선한다.
+
+---
+
+## Consequences
+
+- 다이어그램은 토폴로지 스케일링에 걸쳐 안정적으로 유지된다.
+- 거리 또는 라우팅 정책의 변경이 시각적으로 반영된다.
+- 다이어그램은 수작업으로 유지되는 문서가 아닌, 시뮬레이터 모델로부터
+  파생된 검증 가능한 산출물의 역할을 한다.
+
+---
+
+## Links
+
+- SPEC Section 4 (Output, Debuggability, and Diagrams)
+- ADR-0002 (라우팅 거리 의미)
+- ADR-0006 (토폴로지 컴파일 및 자동 다이어그램 생성)

docs/adr-ko/ADR-0006-dev-topology-compilation-distance-diagram.md

@@ -0,0 +1,130 @@
+# ADR-0006: 토폴로지 컴파일, 거리 추출, 그리고 자동 다이어그램 생성
+
+## Status
+
+Accepted
+
+## Context
+
+시뮬레이터는 토폴로지 설정(예: topology.yaml)을 명시적인 모델 그래프로 컴파일하고,
+라우팅 및 누적 레이턴시(거리)를 계산한다.
+정합성을 보장하고 수작업으로 유지되는 토폴로지 도면을 피하기 위해,
+다이어그램은 이 권위 있는 산출물로부터 생성되어야 한다.
+
+또한 사용성을 위해, 다이어그램은 안정적인 위치로 자동 방출되어
+개발자가 저장소 내에서 즉시 미리볼 수 있어야 한다.
+
+---
+
+## Decision
+
+### D1. 토폴로지 컴파일은 유일한 진실 공급원이다
+
+- topology.yaml(또는 동등한 설정)은 다음으로 컴파일된다:
+  - 명시적인 시스템 그래프,
+  - 노드/링크 속성,
+  - 라우팅 정책.
+이 컴파일된 그래프가 시스템의 권위 있는 표현이다.
+
+### D2. 컴파일 중 거리 추출
+
+- 토폴로지 컴파일 중 또는 그 직후, 시뮬레이터는 ADR-0002와 정합되는
+  거리 메타데이터(누적 레이턴시)를 계산해야 한다.
+- 거리 메타데이터는 ADR-0005에서 정의한 거리 기반 다이어그램 레이아웃을 지원하기에 충분해야 한다.
+- 분산된 패브릭 세그먼트(예: NoC)는 ADR-0002 D4에 따라 distance_mm = 0을 가질 수 있다.
+  이러한 노드의 레이아웃 배치는 거리 버킷이 아닌 명시적 위치 메타데이터를 사용한다.
+
+### D3. 다이어그램 생성은 파생 산출물이다
+
+- 다이어그램은 다음으로부터 생성되어야 한다:
+  - 컴파일된 토폴로지 그래프,
+  - 추출된 거리 메타데이터,
+  - ADR-0005에 정의된 뷰/레이아웃 규칙.
+- 다이어그램 생성은 추가적인 수작업 토폴로지 기술을 요구해서는 안 된다.
+
+### D4. 저장소로의 자동 다이어그램 방출
+
+- 토폴로지 컴파일의 일부로서, 구현은 기본적으로 다음 다이어그램을 생성해야 한다:
+  - SIP 수준 다이어그램 (대표, 거리 인식)
+  - CUBE 수준 다이어그램 (대표, 거리 인식)
+  - PE 수준 다이어그램 (대표, 거리 인식)
+- 기본 출력 디렉터리는 다음과 같다:
+  - `docs/diagrams/`
+- 생성기는 컴파일된 토폴로지(또는 다이어그램 규칙)가 변경되었을 때에만 덮어쓰기/업데이트해야 한다.
+
+### D5. 뷰별 투영 및 레이아웃
+
+각 뷰(SIP / CUBE / PE)에 대해:
+
+- 생성기는 컴파일된 그래프를 축소된 뷰 그래프로 투영해야 한다:
+  - ADR-0005에 따라 노드를 숨기거나 축약하고,
+  - 해당 뷰와 관련된 연결성 의미를 보존하고,
+  - 거리 버킷을 계산하여 레이아웃 레이어를 결정론적으로 할당한다.
+- CUBE 수준 투영은 다음을 포함해야 한다:
+  - 라우터 메시 (cube_mesh.yaml로부터), HBM_CTRL, 공유 SRAM, M_CPU, UCIe 포트,
+    그리고 불투명 블록으로 표현된 PE.
+  - 모든 경로(HBM, 비-HBM, 명령)는 동일한 라우터 메시를 통해 라우팅된다 (ADR-0017).
+- 기본 앵커는 암묵적이며 (ADR-0005) 인스턴스 인덱스를 요구해서는 안 된다.
+
+### D6. 출력 포맷과 결정론
+
+- 생성기는 다음 중 최소 하나를 출력해야 한다:
+  - Mermaid (Markdown 네이티브)
+  - Graphviz DOT (rank 기반 제어)
+  - SVG (mm 단위 정확도 레이아웃, 외부 의존성 없음)
+- 컴파일된 토폴로지로부터 mm 단위 정확도의 위치 메타데이터가 가용한 경우 SVG가 선호된다.
+- 출력은 결정론적이어야 한다:
+  - 동일한 토폴로지 + 동일한 규칙 → 동일한 다이어그램 텍스트
+- 파일 이름은 결정론적이고 안정적이어야 한다 (아래의 "출력 컨벤션" 참조).
+
+### D7. 성능 및 캐싱
+
+- 다이어그램 생성은 지연(lazy) 및/또는 캐시될 수 있으며, `docs/diagrams/`의 출력이
+  컴파일된 토폴로지와 정합을 유지하는 한 그렇다.
+- 구현은 다음을 기반으로 한 캐시 키를 사용해야 한다(SHOULD):
+  - 토폴로지 콘텐츠 해시,
+  - 라우팅 정책 버전,
+  - 다이어그램 규칙 버전,
+  - 뷰 유형 (SIP/CUBE/PE).
+
+---
+
+## 출력 컨벤션
+
+### 디렉터리
+
+- `docs/diagrams/`는 생성된 다이어그램의 표준 출력 디렉터리이다.
+
+### 파일 이름 (권장, 결정론적)
+
+- `system_view.svg` / `system_view.mmd` / `system_view.dot`
+- `sip_view.svg` / `sip_view.mmd` / `sip_view.dot`
+- `cube_view.svg` / `cube_view.mmd` / `cube_view.dot`
+- `pe_view.svg` / `pe_view.mmd` / `pe_view.dot`
+
+선택적으로, 멀티 토폴로지 워크플로우용:
+
+- `sip_view__{topology_id}.svg`
+- `cube_view__{topology_id}.svg`
+- `pe_view__{topology_id}.svg`
+
+### 저장소 정책
+
+- 생성된 다이어그램 파일은 diff 기반 리뷰가 가능하도록 저장소에 커밋될 수 있다.
+- 커밋된 경우, 이는 토폴로지 컴파일로부터 재현 가능해야 한다.
+
+---
+
+## Consequences
+
+- 다이어그램은 항상 시뮬레이터 동작과 정합한다.
+- 아키텍처 변경이 시각화에 자동으로 전파된다.
+- 다이어그램 diff는 아키텍처 변경의 의미 있는 지표가 된다.
+
+---
+
+## Links
+
+- SPEC Section 4 (Output, Debuggability, and Diagrams)
+- ADR-0002 (거리 의미)
+- ADR-0005 (다이어그램 뷰 및 레이아웃 규칙)

docs/adr-ko/ADR-0007-api-runtime-api-boundaries.md

@@ -0,0 +1,95 @@
+# ADR-0007: 런타임 API 및 시뮬레이션 엔진 경계
+
+## Status
+
+Accepted
+
+## Context
+
+시뮬레이터는 책임이 명확히 다른 여러 계층으로 구성된다:
+
+- 벤치마크와 사용자 코드가 사용하는 호스트 대상 API 계층,
+- 요청을 실행하는 이산 이벤트 시뮬레이션 엔진,
+- 하드웨어 동작을 모델링하는 디바이스 컴포넌트.
+
+엄격한 경계가 없으면 오케스트레이션 로직이 컴포넌트로 누출되거나
+시뮬레이션 내부가 사용자 대상 API와 얽힐 수 있다.
+
+본 ADR은 다음 사이의 명확한 책임 경계를 정의한다:
+
+- 런타임 API,
+- 시뮬레이션 엔진 (sip_engine),
+- 하드웨어 컴포넌트.
+
+---
+
+## Decision
+
+### D1. 런타임 API는 호스트 대상 오케스트레이션만 담당
+
+런타임 API는 호스트/드라이버 수준의 동작을 표현하며 다음을 해야 한다:
+
+- 고수준 동작 노출 (텐서 배포, 커널 launch),
+- 엔드포인트 컴포넌트(예: IO_CPU)에만 요청 제출,
+- futures/handles로 완료 대기,
+- 호스트측 메타데이터(텐서 할당 맵, 커널 바인딩)의 소유와 영속화.
+
+런타임 API가 해서는 안 되는 것:
+
+- hop-by-hop 라우팅 또는 fan-out 하드코딩,
+- 내부 컴포넌트(M_CPU, PE_CPU, 엔진) 직접 호출,
+- 토폴로지나 라우팅 관련 가정 내장.
+
+---
+
+### D2. 시뮬레이션 엔진은 컴포넌트를 연결하고 완료를 추적
+
+시뮬레이션 엔진(sim_engine)은 다음을 해야 한다:
+
+- 초기화 시점에 컴포넌트 연결 (컴포넌트 포트/와이어 프레임워크에 따라
+  포트 store 생성 + 와이어 프로세스 시작 — ADR-0015),
+- 컴파일된 토폴로지 그래프의 진입 컴포넌트(예: 메모리 동작은 PCIE_EP,
+  커널 launch는 IO_CPU)에 요청 주입,
+- 이산 이벤트 모델로 이벤트 스케줄링과 실행,
+- correlation ID와 완료 추적 관리.
+
+시뮬레이션 엔진이 해서는 안 되는 것:
+
+- 텐서 시맨틱 정의,
+- 커널 실행 정책 정의,
+- 런타임 API에 내부 그래프 세부사항 노출,
+- 요청 실행 중에 토폴로지 경로를 따라 걷기,
+- 컴포넌트의 `run()` 메서드 직접 호출,
+- hop별 레이턴시 추적 또는 fan-out 분해 (컴포넌트의 책임).
+
+---
+
+### D3. 컴포넌트가 fan-out과 집계를 담당
+
+디바이스측 컴포넌트는 다음을 해야 한다:
+
+- 요청을 하위 도메인으로 fan-out
+  (IO_CPU → M_CPU → PE_CPU → 스케줄러/엔진),
+- 완료·실패 신호 집계,
+- 결정론적으로 상위로 결과 전파.
+
+런타임 API와 시뮬레이션 엔진 모두 컴포넌트 수준의 fan-out을 명시적으로
+오케스트레이션해서는 안 된다.
+
+---
+
+## Consequences
+
+- 토폴로지와 라우팅이 변해도 런타임 API는 안정적이다.
+- 시뮬레이션 내부는 사용자 대상 코드에 영향을 주지 않고 변경 가능하다.
+- 컴포넌트 구현은 DI로 교체 가능한 상태가 유지된다.
+
+---
+
+## Links
+
+- SPEC R4, R7, R8
+- ADR-0008 (텐서 배포)
+- ADR-0009 (커널 실행)
+- ADR-0015 (컴포넌트 포트/와이어 모델과 엔진 역할)
+- ADR-0010 (CLI 표면과 실행 시맨틱 — 런타임 API 소비자)

docs/adr-ko/ADR-0008-api-tensor-deploy-and-allocation.md

@@ -0,0 +1,100 @@
+# ADR-0008: 텐서 배포 및 할당 (호스트 할당기, PA 우선)
+
+## Status
+
+Accepted
+
+## Context
+
+벤치마크는 PyTorch와 유사한 텐서 시맨틱을 요구한다:
+
+- 텐서 생성 (empty, fill),
+- 가속기 디바이스로의 배포 (tensor.to()).
+
+현실적인 시스템에서는 호스트 소프트웨어가 할당·매핑을 관리하고 DMA/MMU
+매핑을 설치한다. Phase 0에서는 (ADR-0011) 다음으로 단순화한다:
+
+- 디바이스 메모리 동작은 PA만 사용,
+- VA/MMU/IOMMU는 모델링하지 않는다.
+
+호스트↔디바이스 인터페이스를 최소로 유지하기 위해 별도의
+AllocateTensorMeta 메시지는 피한다. 대신 호스트 할당은 PA 샤드 맵을
+생성하여 MemoryWrite/Read와 KernelLaunch가 직접 사용한다.
+
+---
+
+## Decision
+
+### D1. Tensor는 PA 샤드 매핑을 가진 호스트 소유 핸들
+
+Tensor 객체는 다음을 캡슐화하는 호스트 소유 핸들이다:
+
+- shape과 dtype,
+- 초기화 의도,
+- PA 샤드 맵 형태의 디바이스 배치 및 할당 메타데이터.
+
+배포 이후 Tensor 핸들은 다음을 포함해야 한다:
+
+- 각각 (sip, cube, pe, pa, nbytes, offset_bytes)를 가진 샤드 리스트.
+
+이 PA 샤드 매핑이 커널 인수 바인딩의 단일 진실 원천이다.
+
+---
+
+### D2. 배포는 호스트 할당기를 사용한다 (Phase 0)
+
+Phase 0에서 텐서 배포는 호스트 할당기를 통해 PA 샤드 매핑을 생성한다:
+
+- 배치(split/replicate/hybrid)는 DP 정책에 의해 결정,
+- 할당은 PE 수준에서 PA 범위를 부여하고 샤드 매핑을 반환,
+- Tensor 핸들은 결정론적으로 결과 샤드 리스트를 저장.
+
+Phase 0에서는 호스트가 보는 별도의 디바이스 할당 RPC는 필요하지 않다.
+
+---
+
+### D3. 데이터 초기화와 전송은 MemoryWrite/Read만 사용
+
+텐서가 함의하는 모든 데이터 초기화나 전송(예: fill, copy)은
+Host ↔ IO_CPU 메시지만으로 표현되어야 한다:
+
+- MemoryWrite
+- MemoryRead
+
+규칙:
+
+- MemoryWrite/Read는 PA + (sip, cube, pe) 태그를 참조해야 한다 (ADR-0012).
+- 할당 메타데이터는 별도의 할당 메시지로 임베드되어서는 안 된다.
+- 대량 텐서 데이터는 Phase 0 메시지에 임베드되어서는 안 된다.
+
+시뮬레이션 엔진은 MemoryWrite/Read를 그래프를 통해 스케줄하므로 레이턴시는
+명시적 순회로 계산된다.
+
+---
+
+### D4. 확장 경로 (호환성 유지)
+
+향후 ADR이 다음을 추가하여 선택적인 VA/MMU/IOMMU 모델링을 도입할 수 있다:
+
+- 텐서 핸들에 가상 주소,
+- 매핑 설치 단계,
+- 변환 레이턴시·페이지 granularity.
+
+Phase 0의 PA 샤드 맵은 유효한 fast-path 구성으로 유지된다.
+
+---
+
+## Consequences
+
+- Host↔IO_CPU 계약이 최소(MemoryRead/Write + KernelLaunch)로 유지된다.
+- KernelLaunch가 샤드 태그를 통해 PE별 데이터 배치를 명시적으로 전달할 수 있다.
+- 초기 구현이 단순하고 테스트 가능하게 유지된다.
+
+---
+
+## Links
+
+- ADR-0011 (메모리 주소 지정 — PA / VA / LA)
+- ADR-0012 (Host↔IO_CPU 스키마)
+- ADR-0007 (runtime_api vs sim_engine 경계)
+- ADR-0009 (커널 실행)

docs/adr-ko/ADR-0009-api-kernel-execution-messaging.md

@@ -0,0 +1,138 @@
+# ADR-0009: 커널 실행 메시징 및 완료 시맨틱
+
+## Status
+
+Accepted
+
+## Context
+
+커널 실행은 호스트에서 시작되어 디바이스 측 제어 컴포넌트를 통해 진행된다:
+
+Host → IO_CPU → M_CPU → PE_CPU → 스케줄러 → 엔진
+
+완료는 역방향으로 전파된다.
+
+벤치마크를 단순하고 토폴로지에 비의존적으로 유지하기 위해, 커널 실행은
+엔드포인트 기반(endpoint-driven)이어야 하며 완료 집계는 결정론적이어야 한다.
+
+---
+
+## Decision
+
+### D1. 커널 런치는 엔드포인트 요청이다
+
+커널 런치는 IO_CPU 엔드포인트에 단일 KernelLaunch 요청을 제출함으로써
+시작된다.
+
+runtime API는 반드시:
+
+- 커널 런치 요청을 구성하고,
+- 이를 IO_CPU로 제출하며,
+- 단일 완료 결과를 대기해야 한다.
+
+runtime API는 내부 팬아웃(fan-out)을 직접 조율해서는 안 된다.
+
+---
+
+### D2. 텐서 인자는 메타데이터로 전달된다
+
+KernelLaunch 요청은 텐서 인자를 다음을 통해 참조해야 한다:
+
+- 호스트가 소유한 텐서 핸들, 또는
+- 그러한 핸들로부터 해석된 디바이스 주소 맵.
+
+대용량 텐서 데이터는 커널 런치 메시지에 임베드되어서는 안 된다.
+
+---
+
+### D3. 팬아웃과 집계는 컴포넌트의 책임이다
+
+- IO_CPU는 작업을 M_CPU들에게 팬아웃한다.
+- M_CPU는 작업을 PE_CPU들에게 팬아웃한다.
+- PE_CPU는 커널 실행과 엔진 디스패치를 관리한다.
+
+완료 시맨틱:
+
+- M_CPU는 대상 PE들이 모두 완료되거나 실패 정책이 트리거되면 완료된다.
+- IO_CPU는 대상 큐브들이 모두 완료되거나 실패 정책이 트리거되면 완료된다.
+
+---
+
+### D4. 완료 및 실패 전파
+
+- 모든 메시지는 correlation ID를 포함해야 한다.
+- 완료와 실패는 호스트로 결정론적으로 전파되어야 한다.
+- 시뮬레이션 엔진은 완료를 관찰할 수 있는 future/handle을 제공한다.
+
+---
+
+### D5. 런치 타이밍은 엔드포인트 동기화된다
+
+단일 커널 런치가 지정한 모든 PE는 런치 진입점으로부터의 디스패치 경로 길이와
+무관하게, 동일한 시뮬레이션 시각에 커널 본문 실행을 시작해야 한다.
+
+근거. 디스패치 트리 Host → IO_CPU → M_CPU → PE_CPU는 모든 레벨에서 가변
+레이턴시를 가진다. M_CPU에 가까운 PE는 멀리 있는 PE보다 런치를 더 일찍
+수신하고, IO_CPU에 가까운 큐브는 먼 큐브보다 더 일찍 수신한다. 동기화가
+없으면 각 PE의 커널은 서로 다른 `env.now`에서 시작되어, `pe_exec_ns`와 같은
+PE별 메트릭이 커널 자체의 동작이 아니라 디스패치 경로 기하 구조의 함수가
+된다 — 그 결과 커널 내부 대기(예: 큐브 간 또는 SIP 간 홉에서의 `tl.recv`)를
+타이밍하는 벤치마크에서 측정 아티팩트가 발생한다.
+
+메커니즘.
+
+- `KernelLaunchMsg`는 선택적 `target_start_ns: float | None`을 포함한다.
+- **IO_CPU**가 정식 스탬프 주체이다. M_CPU들로 팬아웃할 때, 모든 대상
+  (sip, cube, pe) 튜플에 대한 **두 단계 디스패치 체인**의 최대값을
+  `max_latency`로 하여 `target_start_ns = env.now + max_latency`를
+  계산한다:
+
+  ```
+  max_latency(sip, cube, pe) =
+      compute_path_latency_ns(find_node_path(io_cpu, m_cpu(sip, cube)))
+    + compute_path_latency_ns(find_node_path(m_cpu(sip, cube), pe_cpu))
+    - io_cpu.overhead_ns
+    - m_cpu.overhead_ns
+  ```
+
+  이는 실제 디스패치를 **두 개의 순차적 Transaction**(IO_CPU → M_CPU,
+  그리고 M_CPU → PE_CPU)으로 모델링한다. 각 구간의
+  `compute_path_latency_ns`는 양 끝점의 `overhead_ns`를 더하는데,
+  `io_cpu.overhead_ns`는 이 메서드가 실행되기 전 IO_CPU가 이미 지불했으므로
+  차감하고, `m_cpu.overhead_ns`는 구간1의 끝점인 동시에 구간2의 시작점으로
+  나타나지만 런타임에는 한 번만 지불되므로 한 번 차감한다. 단일
+  `find_node_path(io_cpu, pe_cpu)` 순회는 **동등하지 않다** — M_CPU를
+  우회하는 그래프 경로를 선택할 수 있어 먼 큐브에 대해 예측값이 조용히
+  과소평가되며, D5 불변식을 위반하게 된다.
+
+  팬아웃된 하위 Transaction은 `KernelLaunchMsg`에 대해
+  **`nbytes = 0`**을 운반한다(제어 메시지에 한함). 이를 적용하지 않으면
+  큰 커널 런치 페이로드가 공유되는 첫 홉의 패브릭 대역폭을 점유하여
+  큐브별 디스패치를 직렬화하고, 먼 M_CPU들이 `target_start_ns`를
+  지나가게 되어 늦은 도착 위반이 다시 발생한다.
+- **M_CPU**는 이미 스탬프된 `target_start_ns`를 변경 없이 그대로 전달한다.
+  값이 없는 경우(예: M_CPU로 직접 런치하는 단위 테스트)에만 M_CPU가 큐브별
+  배리어 `env.now + max(로컬 명령 경로 레이턴시)`를 계산한다.
+- **PE_CPU**는 `_execute_kernel`의 최상단에서 `pe_exec_start`를 기록하고
+  커널 본문을 호출하기 전에 `env.timeout(target_start_ns - env.now)`를
+  yield한다.
+- `target_start_ns is None`인 경우 PE_CPU는 레거시 비동기 동작으로 빠진다
+  — 하위 호환성을 보존한다.
+
+IO_CPU 레벨 스탬핑은 모든 대상 큐브의 모든 PE가 동일한 배리어 시뮬레이션
+시각을 사용하도록 보장하여, 큐브 내 디스패치 오프셋 아티팩트와 다중 큐브
+런치에서의 큐브 간 오프셋 아티팩트를 모두 제거한다. 실제 하드웨어의
+타이밍 브로드캐스트 런치(레이턴시 등화 디스패치 트리)를 모델링한다.
+
+이 동기화는 엔진 / IO_CPU / M_CPU / PE_CPU 제어 평면 내부에서 수행된다 —
+runtime API와 애플리케이션 커널은 변경되지 않는다.
+
+---
+
+## Links
+
+- SPEC R1, R2, R7, R8
+- ADR-0007 (Runtime API 경계)
+- ADR-0008 (텐서 배치)
+- ADR-0013 (검증 전략 — V2 팬아웃 테스트)
+- ADR-0015 D4 (커널 런치의 구체적 패브릭 경로)

docs/adr-ko/ADR-0010-api-cli-surface-and-semantics.md

@@ -0,0 +1,145 @@
+# ADR-0010: 명령줄 인터페이스 및 실행 시맨틱
+
+## Status
+
+Accepted
+
+## Context
+
+`kernbench` CLI는 시뮬레이터의 사용자 대면 진입점이다. 네 개의 서브명령을
+노출한다:
+
+- `run` — 토폴로지에 대해 벤치마크를 실행한다.
+- `list` — 등록된 벤치마크 목록을 출력한다.
+- `probe` — 레이턴시 / 대역폭 측정을 위한 진단 유틸리티.
+- `web` — 인터랙티브 토폴로지 뷰어.
+
+디바이스 열거는 CLI에 중앙 집중화되어 있다. runtime API와 시뮬레이션 엔진
+모두 디바이스를 열거하지 않는다. 벤치마크는 설계상 단일 디바이스를
+유지하며 입력으로 디바이스 식별자를 받는다.
+
+## Decision
+
+### D1. 벤치마크 계약 — 설계상 단일 디바이스
+
+- 벤치마크는 반드시 단일 디바이스에 대한 동작만 정의해야 한다.
+- 벤치마크는 반드시 디바이스 식별자를 입력으로 받아야 한다.
+- 벤치마크는 다중 디바이스를 열거하거나 루프해서는 안 된다.
+
+다중 디바이스 실행은 벤치마크의 관심사가 아니라 CLI의 관심사이다(D3).
+
+### D2. `kernbench run` — 벤치마크 실행
+
+필수 인자:
+
+- `--topology <path>`: 토폴로지 YAML 파일 경로. `resolve_topology()`를
+  통해 로드된다.
+- `--bench <identifier>`: 벤치마크 식별자. `kernbench.benches.registry.resolve()`를
+  통해 해석되며, 등록된 kebab-case 이름(예: `gemm-single-pe`) 또는
+  `kernbench list` 의 숫자 인덱스를 모두 받는다.
+
+선택 인자:
+
+- `--device <selector>` (기본값: `all`):
+  - `all` — 발견된 SIP마다 한 번씩 실행한다(D3 참고).
+  - `sip:<N>` — SIP N에서만 실행한다.
+  - `resolve_device()`를 통해 파싱된다.
+- `--verify-data` (기본값: off) — Phase 2 데이터 검증을 활성화한다
+  (ADR-0020 참고). 설정되면 `engine_factory`가 엔진을
+  `enable_data=True`로 구성한다. 벤치마크 실행 후, 기록된 op들의 진단
+  요약이 출력된다.
+
+각 호출은 단일 시뮬레이션 인스턴스 내에서 벤치마크를 한 번 실행한다.
+
+### D3. 다중 디바이스 실행은 논리적으로 병렬이다
+
+`--device all`(또는 생략) 상태이며 토폴로지에 SIP가 여러 개일 때:
+
+- 벤치마크 실행은 단일 시뮬레이션 엔진 인스턴스에 제출된다.
+- 시뮬레이션 시간 상에서 실행은 논리적으로 병렬이다.
+- 디바이스 간 경합(공유 패브릭 대역폭, SIP 간 트래픽 등)이 자연스럽게
+  모델링된다.
+
+CLI는 여러 OS 프로세스나 독립된 시뮬레이션 실행을 생성하지 **않는다** —
+병렬성은 단일 시뮬레이션 인스턴스 내부에서 일어난다.
+
+### D4. `kernbench list` — 등록된 벤치마크 목록 출력
+
+인자 없음. 각 등록된 벤치의 자동 부여된 인덱스, 등록된 이름,
+한 줄 설명을 출력한다.
+
+벤치는 `@bench(name=..., description=...)` 데코레이터
+(`kernbench.benches.registry`)를 통해 자기 자신을 등록한다.
+`kernbench.benches/` 아래의 언더스코어로 시작하지 않는 모든 모듈은
+반드시 최소 하나의 벤치를 등록해야 한다; 데코레이터가 누락되면
+패키지 import 시점에 `RuntimeError`가 발생한다.
+
+인덱스는 import 시점에 이름의 알파벳 순으로 부여된다. 인덱스는
+`--bench` 의 축약 표기를 위한 CLI 편의 기능이며 안정적인 API가
+아니다 — 알파벳 순으로 새 벤치가 끼면 이후 인덱스가 밀린다.
+
+### D5. `kernbench probe` — 레이턴시 / 대역폭 진단 유틸리티
+
+필수 인자:
+
+- `--topology <path>`: 토폴로지 YAML 파일 경로.
+
+선택 인자:
+
+- `--case <name>` (기본값: `all`) — 미리 정의된 트래픽 패턴을 실행하거나,
+  `all`로 정의된 모든 케이스를 실행한다.
+
+Probe는 시뮬레이션 엔진을 통해 각 패턴을 실행하고 케이스별로 다음을
+보고한다:
+
+- 종단 간 레이턴시(ns).
+- 유효 대역폭(nbytes / total_ns).
+- 병목 대역폭(선택된 경로상의 최소 엣지 BW).
+- 활용률(유효 / 병목).
+
+Probe는 추가로 단조성 불변식을 검증한다 — 예를 들어 local-HBM 접근 ≤
+큐브 내 PE 간 ≤ 큐브 간 ≤ SIP 간 — 그리고 위반을 보고한다. Probe는
+레이턴시 / 대역폭 모델을 검증하기 위한 개발자 도구이다; 벤치마크가
+아니다.
+
+### D6. `kernbench web` — 토폴로지 뷰어
+
+선택 인자:
+
+- `--port <N>` (기본값: `8765`) — HTTP 포트.
+- `--no-open` — 브라우저를 자동으로 열지 않는다.
+
+컴파일된 토폴로지를 브라우저에서 렌더링하는 로컬 HTTP 서버를 띄운다.
+정적인 `docs/diagrams/` 산출물과는 구별된다:
+
+- `docs/diagrams/` 파일은 토폴로지 컴파일 시점에 파생된다(ADR-0006).
+- `kernbench web`은 인터랙티브이다 — 팬/줌, 컴포넌트 속성 호버,
+  SIP / CUBE / PE 뷰 간 전환.
+
+### D7. runtime API와 시뮬레이션 엔진은 디바이스 스코프를 유지한다
+
+- runtime API 호출은 호출당 하나의 디바이스에서 동작한다.
+- 시뮬레이션 엔진은 모든 요청을 결정론적으로 스케줄링한다.
+- 어느 레이어도 디바이스를 열거하지 않는다.
+
+이 불변식은 각 레이어를 독립적으로 테스트 가능하게 유지한다; 디바이스
+열거와 다중 디바이스 팬아웃은 오직 CLI의 `run` 명령에만 존재한다(D3).
+
+`probe` 구현은 `kernbench.probes` 아래에 있다 (`kernbench.benches`와
+분리됨). 이는 probe가 등록된 벤치가 아니라 진단 유틸리티임을 반영한다.
+
+## Consequences
+
+- 벤치마크 작성자는 단일 디바이스 로직을 작성한다; 다중 디바이스 동작은
+  CLI가 SIP들에 걸쳐 디스패치함으로써 자연스럽게 도출된다.
+- 새로운 서브명령(예: 트레이스 내보내기, 리플레이) 추가는 벤치마크나
+  runtime API 변경을 요구하지 않는다 — CLI가 확장 포인트이다.
+- `probe`와 `web`은 진단/시각화 도구이며 벤치마크가 아니다; 벤치마크 로더
+  경로를 우회한다.
+
+## Links
+
+- SPEC R7, R8, R9
+- ADR-0007 (Runtime API와 시뮬레이션 엔진 경계)
+- ADR-0020 (Two-pass 데이터 실행 — `--verify-data`)
+- ADR-0006 (토폴로지 컴파일과 다이어그램 생성 — `kernbench web`의 배경)

docs/adr-ko/ADR-0011-mem-memory-addressing-simplification.md

@@ -0,0 +1,503 @@
+# ADR-0011: 메모리 주소 지정 — PA / VA / LA 주소 모델
+
+## Status
+
+Accepted.
+
+- **VA 모델: 현재 구현됨 (기본값).**
+- PA 모델: PE_DMA의 PageFault fallback으로 구현됨.
+- LA 모델: 제안됨, 미구현.
+
+## Context
+
+KernBench의 주소 모델은 각 단계마다 이전 단계의 한계를 해결하면서
+세 단계의 설계 지점을 거쳐 발전해 왔다. 본 ADR은 미래의 구현 작업이
+이 셋 중 하나를 선택해야 하므로 셋 모두를 한 곳에 기록한다.
+
+### PA 단독 베이스라인
+
+KernBench Phase 0는 모든 디바이스 메모리 동작(MemoryRead/MemoryWrite)을
+순수 물리 주소 전송으로 다뤘다. 호스트측 가상 주소 지정 없음, MMU/IOMMU
+변환 없음. 할당기는 PA 매핑을 반환하고, DMA 요청은 PA를 직접 운반했다.
+
+이는 초기 정확성·레이턴시 작업에는 충분했지만, 샤딩된 텐서에 대해
+`base_addr + offset` 패턴을 사용하는 표준 Triton 커널을 실행하기에는
+부족했다. 각 PE의 샤드는 서로 다른 PA를 갖지만, 커널은 offset을 계산하기
+위해 연속된 단일 주소 공간이 필요하기 때문이다.
+
+### VA/MMU를 채택한 이유 (현재 기본값)
+
+현실적인 시스템은 호스트측 가상 주소 지정과 DMA를 위한 MMU/IOMMU 스타일
+변환 경로를 사용한다. 호스트는 PE 수준에서 물리 메모리를 할당하고,
+그것을 가상 주소 공간에 매핑하여 매핑을 설치한 뒤, DMA 요청은 가상
+주소를 사용하며 그것이 물리 주소로 변환된다.
+
+이 모델을 채택하면 커널이 연속된 VA 범위에 대해 `base_addr + offset`을
+사용할 수 있고, 디바이스측 MMU가 각 접근을 적절한 PA로 변환한다.
+
+### LA/BAAW를 제안한 이유
+
+VA/MMU는 HBM을 단일 backing 공간으로 다룬다. KernBench는 HBM이 병렬로
+여러 pseudo channel로 구성된 아키텍처를 탐색해야 한다:
+
+- CUBE의 HBM은 32 또는 64개의 pseudo channel을 갖는다.
+- PE-Local-HBM 모델에서 각 PE에는 N개의 pseudo channel이 할당된다
+  (N = `hbm_pseudo_channels / pes_per_cube`).
+- 채널당 대역폭(예: 32 GB/s)이 PE의 총 대역폭을 결정한다
+  (N × 채널당).
+
+두 가지 채널 매핑 모드를 모델링할 수 있어야 한다:
+
+- **1:1 모드** — 하나의 논리 접근 → N개의 채널별 요청.
+  채널별 대역폭 경쟁을 정밀하게 모델링.
+- **n:1 모드 (기본값)** — 하나의 논리 접근 → 하나의 집계 요청.
+  채널들이 interleave된다고 가정; 집계된 대역폭 모델.
+
+VA의 `tl.load(va_ptr)`은 하나의 목표에 대한 하나의 DMA 요청을 생성한다.
+이를 PE_DMA 내부에서 채널별 요청으로 분해하려면 주소 계층이 채널을
+인지해야 한다. 이것이 BAAW(Logical-to-Physical Mapping Unit)를 가진
+LA(Logical Address) 추상화의 역할이다.
+
+LA 설계를 이끄는 핵심 요구사항:
+
+- PE_DMA → HBM_CTRL 유효 대역폭 시맨틱이 두 모드에서 동일해야 한다
+  (요청 형태와 자원 모델만 다름).
+- 커널 프로그래밍 모델은 변경되지 않는다 — 물리 채널 정보는 커널 코드에
+  절대 노출되지 않는다.
+- 모드 전환은 토폴로지 수준의 설정이다.
+
+### 설계 공간 요약
+
+| 모델 | 상태 | 핵심 아이디어 |
+|------|------|--------------|
+| PA | fallback (구현됨) | 직접 물리 주소 지정, 변환 없음 |
+| VA | 현재 기본값 (구현됨) | 텐서별 연속 VA 범위; MMU가 접근별로 변환 |
+| LA | 제안됨 | LA + BAAW가 (PA, 채널)로 해석; 1:1 및 n:1 채널 매핑 모드 지원 |
+
+---
+
+## Decision
+
+본 ADR은 세 개의 주소 모델을 정의한다. 어느 시점에도 시스템은 정확히
+한 모델로 동작한다. 선택은 토폴로지·설정 주도이며, 단일 시뮬레이션 실행
+내에서의 공존은 요구되지 않는다.
+
+---
+
+### 주소 모델: PA (물리 주소) — fallback
+
+#### D-PA1. PA 단독 시맨틱
+
+- 모든 디바이스 메모리 접근(MemoryRead/MemoryWrite)은 디바이스 물리 주소(PA)와
+  크기에 대해 동작한다.
+- PA 단독 모드는 PE_DMA의 PageFault fallback 경로를 통해 여전히 동작한다.
+  DMA src/dst 주소에 MMU 매핑이 없으면 PE_DMA는 그 값을 PA로 직접 다룬다.
+
+#### D-PA2. 할당은 PA 매핑을 생성한다
+
+디바이스 할당은 PE 로컬 메모리 영역을 선택하고 커널 실행 및 DMA 요청
+발행에 충분한 PA 매핑을 반환한다.
+
+PA 모델은 주로 PA 단독 테스트와의 하위 호환성을 위해, 그리고 VA / LA
+모델이 해석되어 들어가는 기저 물리 계층으로 유지된다.
+
+---
+
+### 주소 모델: VA (MMU를 동반한 가상 주소) — 현재 기본값
+
+#### D-VA1. 가상 주소 모델
+
+- 각 텐서는 하나의 연속된 VA 범위(`TensorHandle.va_base`)를 가진다.
+- `TensorShard`는 `va` 필드를 가지지 **않는다** — 샤드 VA는
+  `va_base + offset_bytes`로 유도된다.
+- 커널은 포인터 인수로 `va_base`를 받는다(`TensorArg.va_base` 경유).
+- `DmaReadCmd.src_addr`와 `DmaWriteCmd.dst_addr`는 VA(PA가 아님)를 운반한다.
+
+#### D-VA2. PE_MMU 컴포넌트
+
+- 하이브리드 설계: SimPy 컴포넌트(`MmuMapMsg`용 inbox) + 유틸리티
+  (PE_DMA가 호출하는 동기식 `translate()`).
+- 페이지 정렬 dict 조회로 O(1) VA → PA 변환.
+- `tlb_overhead_ns`로 접근당 레이턴시 설정 가능.
+- PageFault fallback: VA에 매핑이 없으면 PE_DMA가 그것을 PA로 직접
+  다룬다 (PA 모델과의 하위 호환성 유지).
+
+#### D-VA3. 매핑 설치
+
+- `MmuMapMsg`는 패브릭을 순회한다: Host → PCIE_EP → IO_CPU (큐브 fan-out)
+  → M_CPU (PE fan-out) → NOC → PE_MMU. 레이턴시는 end-to-end로 측정된다.
+- `MmuMapMsg.target_sips`는 SIP 수준 라우팅을 제어하여 복제 텐서의
+  cross-SIP 매핑 오염을 방지한다.
+- `DPPolicy.cube`에 기반한 매핑 전략:
+  - **Replicate** (`cube="replicate"`): (sip, cube)별 로컬 매핑만.
+    각 큐브의 PE들은 자신의 로컬 PA만 본다. cross-cube 매핑은 설치되지
+    않는다.
+  - **Sharded** (`cube="column_wise"` 등): 모든 샤드 매핑을 모든 대상
+    큐브로 브로드캐스트. cross-PE 및 cross-cube DMA를 가능하게 한다.
+
+#### D-VA4. 텐서 라이프사이클
+
+- `del tensor`는 `Tensor.__del__` + `RuntimeContext`에 대한 `weakref`를
+  통해 자동 정리를 트리거한다. 패브릭을 통해 `MmuUnmapMsg`를 보내고
+  VA와 PA 공간을 반환한다.
+- `with RuntimeContext(...) as ctx:`는 스코프 기반 일괄 정리를 제공한다.
+- `RuntimeContext._tensors`는 GC 방지를 피하기 위해 `weakref.ref`를 사용.
+- `PEMemAllocator`는 coalescing이 있는 free-list를 사용한다(bump allocator 아님).
+- `VirtualAllocator`는 VA 공간에 대해 coalescing이 있는 free-list를 사용한다.
+
+#### D-VA5. 할당기
+
+- `VirtualAllocator`: 디바이스 전체의 VA 공간, coalescing을 동반한
+  페이지 정렬 alloc/free.
+- `PEMemAllocator`: PE별 HBM/TCM, coalescing을 동반한 free-list 기반
+  alloc/free.
+- 페이지 크기는 `topology.yaml`의 `pe_mmu` attrs로 설정 가능
+  (기본 4096).
+
+#### Consequences (VA 모델)
+
+- Triton 커널은 샤딩된 텐서에 대해 `base_addr + offset` 패턴을 자연스럽게
+  사용한다.
+- 모든 레이턴시는 MMU 매핑 설치와 접근당 TLB 오버헤드를 포함하여
+  그래프 순회를 통해 명시적이다.
+- PA 단독 모드는 fallback으로 유지된다 (PageFault → PA로 처리).
+- IPCQ와 그 외 고정 주소 자원은 MMU를 우회한다 (PA 직접 사용).
+
+---
+
+### 주소 모델: LA (BAAW를 동반한 논리 주소) — 제안됨
+
+LA는 채널 수준 HBM 모델링이 필요할 때 VA를 대체한다.
+이 모델을 채택하면 VA/MMU 인프라가 제거된다 (D-LA1이 제거되는 산출물을
+나열한다). 동일 실행 내에서 VA와의 공존은 목표가 아니다.
+
+#### D-LA1. LA 도입 — VA 인프라 대체
+
+LA는 커널 코드(`tl.load`, `tl.store`, `tl.composite`)가 사용하는
+유일한 주소 공간이다. 속성:
+
+- Tensor를 연속된 논리 공간에 매핑할 수 있다 (VA처럼).
+- `(논리 버퍼 + offset)`을 표현한다.
+- 물리 채널 정보를 직접 포함하지 **않는다**.
+- 물리적 해석이 일어나기 전까지는 중간 추상화로 유지된다.
+
+LA 주소 공간:
+
+| 항목 | 값 |
+|------|-------|
+| LA 시작 | `0x1_0000_0000` (4 GB, 이전 VA 시작과 동일) |
+| LA 공간 크기 | PE당 64 GB |
+| 정렬 단위 | segment (D-LA3 참조) |
+
+LA는 PE 로컬이다: 서로 다른 PE가 동일한 LA 값을 사용할 수 있지만,
+BAAW segment 테이블이 다르므로 서로 다른 PA로 해석된다.
+
+LA가 채택되면 제거되는 VA 인프라:
+
+| 제거 | 대체 |
+|---------|-------------|
+| `policy/address/va_allocator.py` (VirtualAllocator) | LA allocator (동일한 free-list 접근, 이름 변경) |
+| `policy/address/pe_mmu.py` (PeMMU) | BAAW segment 테이블 (PE_DMA 내부) |
+| `components/builtin/pe_mmu.py` (PeMmuComponent) | 제거 — BAAW는 별도 컴포넌트가 아니라 PE_DMA 내부 로직 |
+| `runtime_api/kernel.py`: `MmuMapMsg`, `MmuUnmapMsg` | `BaawSegmentInstallMsg` |
+| `runtime_api/context.py`: VA alloc + MMU install | LA alloc + BAAW segment install |
+| `runtime_api/tensor.py`: `va_base` | `la_base` |
+| `topology.yaml`: `pe_mmu` 컴포넌트 entry | 제거 |
+
+#### D-LA2. 매핑 모드 설정
+
+토폴로지 수준(큐브) 설정:
+
+```yaml
+cube:
+  memory_map:
+    hbm_mapping_mode: n_to_one    # one_to_one | n_to_one
+    hbm_pseudo_channels: 64       # 전체 pseudo channel 수
+    hbm_channels_per_pe: 8        # PE당 로컬 채널 수
+    hbm_channel_bw_gbs: 32.0      # 채널당 대역폭
+```
+
+그래프 컴파일러(토폴로지 빌더)와 BAAW 초기화가 이 값을 소비한다.
+
+#### D-LA3. Segment와 BAAW
+
+Segment는 LA 공간을 분할한다. 각 segment는 특정 HBM 채널 또는 채널
+그룹에 매핑된다. 텐서 deploy 시점에 런타임 할당기가 생성한다. BAAW는
+segment 테이블을 사용하여 LA → 물리 요청(들)로 해석한다.
+
+```python
+@dataclass
+class BaawSegment:
+    la_base: int          # segment 시작 LA
+    la_size: int          # segment 크기 (bytes)
+    mode: str             # "one_to_one" | "n_to_one"
+    # 1:1 모드 필드
+    channel_count: int    # 이 segment에 할당된 채널 수 (예: 8)
+    pa_bases: list[int]   # 채널별 PA base (len = channel_count)
+    channel_ids: list[int]   # 채널별 논리 ID (예: [0..7])
+    channel_size: int     # 채널당 크기 (la_size // channel_count)
+    # n:1 모드 필드
+    agg_pa_base: int      # 집계 PA base
+    agg_node_id: str      # 집계 라우터 node_id
+```
+
+Segment 라이프사이클:
+
+1. **할당** (텐서 deploy): RuntimeContext가 LA allocator에서 LA를
+   할당한다. PEMemAllocator가 채널별 PA(1:1) 또는 집계 PA(n:1)를
+   할당한다. `BaawSegmentInstallMsg`가 segment를 PE_DMA에 등록한다.
+2. **사용** (커널 실행): 커널 `tl.load(la_ptr)` → `DmaReadCmd
+   (src_addr=LA)`. PE_DMA의 BAAW 프론트엔드가 segment를 조회하여
+   PA(들)로 변환한다.
+3. **해제** (텐서 free): segment가 테이블에서 제거되고 LA와 PA가
+   반환된다.
+
+#### D-LA4. BAAW 해석 로직
+
+BAAW는 PE_DMA 내부의 프론트엔드 단계이며, 별도의 SimPy 컴포넌트가 아니다.
+PE_DMA의 `handle_command()` 시작 시점에 실행되는 동기식 주소 해석 로직.
+
+입력: `(LA, nbytes)`. 출력:
+
+- **1:1 모드**: `list[PhysicalRequest]` — 채널당 하나.
+- **n:1 모드**: 단일 `PhysicalRequest`.
+
+```python
+@dataclass
+class PhysicalRequest:
+    pa: int           # 51-bit 물리 주소
+    nbytes: int       # 이 요청의 전송 크기
+    dst_node: str     # 대상 node_id (채널 라우터 또는 집계 라우터)
+
+
+def resolve(self, la: int, nbytes: int) -> list[PhysicalRequest]:
+    seg = self._find_segment(la)  # la_base <= la < la_base + la_size
+    offset = la - seg.la_base
+
+    if seg.mode == "n_to_one":
+        pa = seg.agg_pa_base + offset
+        return [PhysicalRequest(pa=pa, nbytes=nbytes, dst_node=seg.agg_node_id)]
+
+    # one_to_one
+    requests = []
+    per_ch_size = seg.channel_size
+    for i, (pa_base, ch_id) in enumerate(zip(seg.pa_bases, seg.channel_ids)):
+        ch_offset = offset % per_ch_size
+        ch_nbytes = nbytes // seg.channel_count
+        pa = pa_base + ch_offset
+        dst_node = f"{self._pe_prefix}.ch_r{ch_id}"
+        requests.append(PhysicalRequest(pa=pa, nbytes=ch_nbytes, dst_node=dst_node))
+    return requests
+```
+
+BAAW의 책임:
+
+- 논리 접근 → 물리 요청 단위로 변환.
+- 모드에 따라 fan-out(1:1) 또는 pass-through(n:1) 적용.
+- PA와 대상 노드 계산.
+
+BAAW가 하지 않는 것:
+
+- 실제 데이터 이동 수행.
+- NOC 라우팅 실행.
+- 대역폭 점유 시뮬레이션 (하위 컴포넌트의 역할).
+
+BAAW의 출력은 추가적인 주소 디코딩 없이 시뮬레이터의 라우팅·자원
+모델에서 바로 사용 가능하다.
+
+#### D-LA5. PE_DMA `handle_command()` 변경
+
+현재(VA 기반) 흐름:
+
+```
+DmaReadCmd.src_addr (VA)
+  → MMU.translate(VA) → PA
+  → PhysAddr.decode(PA) → PhysAddr 객체
+  → resolver.resolve(PhysAddr) → dst_node_id
+  → router.find_path(pe_prefix, dst_node_id) → path
+  → 1 sub-Transaction → 패브릭 주입
+```
+
+LA 기반 흐름:
+
+```
+DmaReadCmd.src_addr (LA)
+  → BAAW.resolve(LA, nbytes) → list[PhysicalRequest]
+  → 각 PhysicalRequest에 대해:
+      → router.find_path(pe_prefix, req.dst_node) → path
+      → compute_drain_ns(path, req.nbytes) → drain
+      → sub-Transaction → 패브릭 주입
+  → 모든 sub-Transaction 대기
+  → pe_txn.done.succeed()
+```
+
+주요 변경:
+
+- MMU 참조 제거 → BAAW resolve.
+- `PhysAddr.decode()` + `resolver.resolve()` → BAAW가 `dst_node`를
+  직접 반환.
+- 1 요청 → 1:1 모드에서 N개의 병렬 요청.
+
+#### D-LA6. 1:1 모드 상세
+
+- 하나의 논리 접근 → N개의 물리 요청 (N = `channels_per_pe`).
+- N = `hbm_pseudo_channels / pes_per_cube`.
+- 각 요청: 완전히 해석된 51-bit PA, 특정 채널 라우터를 대상으로 함
+  (`{pe_prefix}.ch_r{channel_id}`).
+- 채널별 링크가 대역폭 경쟁을 모델링.
+- PE_DMA가 N개의 sub-transaction을 동시에 주입.
+
+예: `hbm_pseudo_channels=64`, `pes_per_cube=8` → `channels_per_pe=8`.
+PE0은 ch0-7을 소유.
+
+```text
+Tensor A (4 KB) → LA 0x1_0000_0000, size=4096 bytes
+BAAW segment: {
+    la_base: 0x1_0000_0000, la_size: 4096,
+    mode: "one_to_one", channel_count: 8,
+    pa_bases: [PA_ch0, PA_ch1, ..., PA_ch7],
+    channel_ids: [0, 1, 2, 3, 4, 5, 6, 7],
+    channel_size: 512,
+}
+
+BAAW resolve 결과 (8 요청):
+  → PhysicalRequest(pa=PA_ch0, nbytes=512, dst_node="sip0.cube0.pe0.ch_r0")
+  → PhysicalRequest(pa=PA_ch1, nbytes=512, dst_node="sip0.cube0.pe0.ch_r1")
+  → ...
+  → PhysicalRequest(pa=PA_ch7, nbytes=512, dst_node="sip0.cube0.pe0.ch_r7")
+
+PE_DMA: 8개 sub-transaction 병렬 주입
+  채널별 라우터 → hbm_ctrl 링크 (channel_bw_gbs) per channel
+  전체 유효 BW = 8 × channel_bw_gbs
+```
+
+다른 N 값:
+
+- `hbm_pseudo_channels=32`, `pes_per_cube=8` → `channels_per_pe=4`,
+  4 요청
+- `hbm_pseudo_channels=64`, `pes_per_cube=4` → `channels_per_pe=16`,
+  16 요청
+
+#### D-LA7. n:1 모드 상세
+
+- 하나의 논리 접근 → 하나의 집계 요청.
+- 대상: 집계 라우터 → hbm_ctrl (ADR-0017 D8 참조).
+- 집계 링크 BW = `channels_per_pe × channel_bw_gbs`
+  (예: 8 × 32 = 256 GB/s).
+- 모델링을 위한 단일 큐 / 자원.
+- 채널별 PA 분해 없음.
+
+```text
+Tensor A (4 KB) → LA 0x1_0000_0000, size=4096 bytes
+BAAW segment: {
+    la_base: 0x1_0000_0000, la_size: 4096,
+    mode: "n_to_one",
+    agg_pa_base: PA_agg,
+    agg_node_id: "sip0.cube0.pe0.agg_router",
+}
+
+BAAW resolve 결과:
+  → PhysicalRequest(pa=PA_agg, nbytes=4096, dst_node="sip0.cube0.pe0.agg_router")
+
+PE_DMA: 1 sub-transaction
+  집계 라우터 → hbm_ctrl 링크 (256 GB/s)
+```
+
+#### D-LA8. 커널 모델 보존
+
+- 커널은 여전히 단일 메모리 op(`tl.load`, `tl.store`,
+  `tl.composite`)을 발행한다.
+- LA가 커널 코드에 노출되는 주소 체계이다.
+- 채널 분해·집계는 PE_DMA의 BAAW 내부에서 일어난다.
+- 커널 코드는 물리 채널 정보를 절대 보지 않는다.
+
+#### Consequences (LA 모델, 제안됨)
+
+긍정적:
+
+- 1:1 vs n:1 시맨틱이 한 곳(BAAW)에 모인다.
+- 커널 추상화 보존 — 커널 코드 변경 없음.
+- 토폴로지 기반 정책 제어 (yaml로 모드 전환).
+- 시뮬레이션 모델의 정합성·디버깅 가능성 향상.
+- Segment 기반 매핑이 페이지 테이블보다 단순하며 오버헤드도 적다.
+
+부정적:
+
+- 전체 VA/MMU 코드 리팩터가 필요하다.
+- 요청 생성 경로가 더 복잡 (1:1 모드에서 N 요청).
+- n:1 모드에서 채널별 가시성 감소.
+- VA 관련 테스트 재작성 필요.
+
+---
+
+## Migration Path
+
+- **PA → VA**는 확장이었다. PA 모드는 PE_DMA 내부의 PageFault fallback으로
+  유지된다. 전환은 PA 코드 제거를 요구하지 않는다.
+- **VA → LA**는, 채택될 경우, 공존이 아닌 대체이다. VA 인프라 제거
+  목록은 D-LA1 참조. PA fallback은 테스트를 위해 PE_DMA 내부에 직교적으로
+  유지될 수 있다.
+
+## Alternatives Considered (LA 모델)
+
+1. **VA 유지 + MMU에서 fan-out**: MMU가 채널별 PA를 반환한다.
+   기각: MMU의 역할이 변환을 넘어 요청 분해까지 확장되며, 집계(n:1)를
+   표현하기 어색해진다.
+2. **채널 인지 커널 API**: 커널이 채널별 load/store를 직접 호출한다.
+   기각: 추상화 누출, 이식성 손실, 모든 벤치마크 재작성 필요.
+3. **항상 PA (LA 없음)**: 런타임이 커널에 채널별 PA를 직접 전달한다.
+   기각: 집계와 양립 불가; 변환 시점이 불명확; 채널 정보가 커널로 누출.
+
+## Test Requirements
+
+### VA 모델 (현재, regression)
+
+- 설치된 매핑을 따라 cross-PE / cross-cube DMA 경로.
+- 측정된 레이턴시를 동반한 `MmuMapMsg` / `MmuUnmapMsg`의 패브릭 순회.
+- 접근당 TLB 오버헤드 타이밍.
+- PageFault fallback 경로가 PA 단독 동작을 보존하는지.
+
+### LA 모델 (구현 시)
+
+- 1:1 모드: 동일 논리 접근 → N개의 채널별 요청.
+- n:1 모드: 동일 논리 접근 → 1개의 집계 요청.
+- 동일 워크로드에 대해 두 모드 사이의 대역폭 동치.
+- 1:1 모드: 채널별 경쟁이 올바르게 모델링됨.
+- n:1 모드: 집계된 대역폭이 올바르게 반영됨.
+- 모드 전환에 걸쳐 커널 코드가 변경되지 않음.
+- BAAW segment install / uninstall 정확성.
+- 별개 segment 안의 여러 텐서가 충돌하지 않음.
+
+## Implementation Order (LA, 일정 잡힐 때)
+
+1. LA 타입 (`policy/address/la_allocator.py`).
+2. BAAW segment 테이블 (`policy/address/baaw.py`).
+3. `BaawSegmentInstallMsg` (`runtime_api/kernel.py`).
+4. PE_DMA BAAW 통합 (`components/builtin/pe_dma.py`
+   `handle_command()`).
+5. RuntimeContext: LA alloc + segment install
+   (`runtime_api/context.py`).
+6. `Tensor.va_base` → `Tensor.la_base` (`runtime_api/tensor.py`).
+7. VA/MMU 코드 제거.
+8. `topology.yaml`에서 `pe_mmu` 제거; 매핑 모드 설정 추가.
+9. 테스트 이전:
+
+| 테스트 파일 | 조치 |
+|-----------|--------|
+| `tests/test_mmu_component.py` | 제거 → BAAW segment install 테스트 |
+| `tests/test_mmu_fabric.py` | 제거 → BAAW + 패브릭 통합 테스트 |
+| `tests/test_pe_mmu.py` | 제거 |
+| `tests/test_va_allocator.py` | LA allocator 테스트로 교체 |
+| `tests/test_va_integration.py` | LA + BAAW 통합 테스트로 교체 |
+| `tests/test_va_offset.py` | LA offset 테스트로 교체 |
+
+## Links
+
+- ADR-0007 (runtime_api vs sim_engine 경계)
+- ADR-0008 (텐서 배포)
+- ADR-0009 (커널 실행)
+- ADR-0014 (PE 내부 실행 모델)
+- ADR-0015 (컴포넌트 포트/와이어 모델)
+- ADR-0017 (큐브 NOC와 HBM 연결성 — LA 모델 토폴로지 소비자)
+- ADR-0013 (검증 전략 — V1 PA 태깅)
+- SPEC R2 (순회 기반 레이턴시), R10 (메모리 주소 지정)

docs/adr-ko/ADR-0012-api-host-io-message-schema.md

@@ -0,0 +1,239 @@
+# ADR-0012: Host ↔ IO_CPU 메시지 스키마 (PA-우선, PE-태깅)
+
+## Status
+
+Accepted
+
+## Context
+
+Phase 0은 PA-우선 메모리 모델을 사용한다(ADR-0011):
+
+- 메모리 연산은 디바이스 물리 주소(PA)만 사용한다,
+- VA/MMU/IOMMU는 모델링하지 않는다.
+
+호스트 대면 runtime API는 IO_CPU 엔드포인트를 통해 디바이스와
+상호작용한다. 다음을 보장하기 위해 Host ↔ IO_CPU에 대한 안정적이고
+최소한의 메시지 스키마를 정의한다:
+
+- 벤치마크는 안정적으로 유지된다,
+- IO_CPU 내부의 팬아웃/집계는 독립적으로 진화할 수 있다,
+- 완료와 실패 전파는 결정론적이다.
+
+또한 PE-태깅(A 방식)을 요구한다: 각 샤드는 (sip,cube,pe)를 명시적으로
+운반하여, IO_CPU가 PA 디코딩에 의존하지 않고 결정론적으로
+라우팅/팬아웃할 수 있도록 한다.
+
+---
+
+## Decision
+
+### D1. 계약 범위
+
+본 스키마는 오직 Host ↔ IO_CPU에 대해서만 안정적인 계약이다.
+
+IO_CPU를 넘어선 메시지(M_CPU, PE_CPU, 스케줄러, 엔진으로 가는 것)는
+컴포넌트 내부 사항이며 Phase 0에서 이 호스트 계약의 일부가 아니다.
+
+---
+
+### D2. 필수 메시지 집합
+
+runtime API는 Host ↔ IO_CPU에 대해 오직 다음 메시지 타입만 사용해야 한다:
+
+- MemoryWrite
+- MemoryRead
+- KernelLaunch
+
+벤치마크가 필요로 하는 모든 연산(텐서 초기화/복사, 커널 실행)은 이
+메시지들로 표현 가능해야 한다.
+
+---
+
+### D3. 공통 envelope (모든 요청에 필수)
+
+모든 Host ↔ IO_CPU 요청은 반드시 다음을 포함해야 한다:
+
+- `msg_type: str`
+- `correlation_id: str`
+  - 호스트에서 생성
+  - 응답을 결정론적으로 매칭하는 데 사용
+- `request_id: str`
+  - correlation_id 내에서 고유함
+- `target_device: str`
+  - 디바이스 식별자(예: "sip:0")
+- `timestamp_tag: str | None` (선택)
+  - 디버그 태그 전용; 결정성에 영향을 주어서는 안 됨
+
+모든 Host ↔ IO_CPU 응답은 반드시 다음을 포함해야 한다:
+
+- `correlation_id: str`
+- `request_id: str`
+- `completion: Completion`
+
+---
+
+### D4. Completion 스키마 (필수)
+
+`Completion`은 반드시 다음을 가져야 한다:
+
+- `ok: bool`
+- `error_code: str | None`
+- `error_message: str | None`
+
+규칙:
+
+- `ok == true`이면 `error_code`와 `error_message`는 반드시 null이어야 한다.
+- `ok == false`이면 `error_code`는 반드시 null이 아니어야 한다.
+- 완료 시맨틱은 결정론적이어야 한다.
+
+---
+
+### D5. MemoryWrite 스키마 (PA-우선, PE-태깅)
+
+`MemoryWrite`는 호스트에서 시작된 디바이스 메모리 쓰기/초기화 연산을
+나타낸다.
+
+필수 필드:
+
+- 공통 envelope 필드 (D3)
+- 목적지 배치 태그 (A 방식):
+  - `dst_sip: int`
+  - `dst_cube: int`
+  - `dst_pe: int`
+- `dst_pa: int`
+  - 목적지 PE의 주소 공간 내 목적지 물리 주소
+- `nbytes: int`
+- `src_kind: "pattern" | "host_buffer_ref"`
+  - Phase 0은 반드시 "pattern"을 지원해야 한다
+- `pattern: Pattern | None`
+  - `src_kind == "pattern"`인 경우 필수
+
+`Pattern` (Phase 0 필수 지원):
+
+- `pattern_kind: "zero" | "fill_u8" | "fill_u16" | "fill_u32" | "fill_fp16" | "fill_fp32"`
+- `value: number | None`
+  - fill_*에 필요; zero에서는 무시됨
+
+선택 필드:
+
+- `dst_mem_kind: "HBM" | "TCM" | "AUTO"` (기본값 "AUTO")
+- `debug_label: str | None`
+
+비고:
+
+- 이 메시지는 Phase 0에서 대용량 텐서 데이터를 임베드해서는 안 된다.
+- 모든 레이턴시는 명시적인 그래프 순회 및 모델링된 컴포넌트로부터
+  발생해야 한다.
+
+---
+
+### D6. MemoryRead 스키마 (PA-우선, PE-태깅)
+
+`MemoryRead`는 호스트에서 시작된 디바이스 메모리 읽기를 나타낸다.
+
+필수 필드:
+
+- 공통 envelope 필드 (D3)
+- 소스 배치 태그 (A 방식):
+  - `src_sip: int`
+  - `src_cube: int`
+  - `src_pe: int`
+- `src_pa: int`
+- `nbytes: int`
+
+선택 필드:
+
+- `dst_kind: "host_sink" | "discard"` (기본값 "host_sink")
+- `debug_label: str | None`
+
+응답 페이로드:
+
+- Phase 0에서는 실제 바이트는 필요하지 않다(레이턴시/트레이스 중심)
+- 구현은 추후 새로운 ADR을 통해 가벼운 통계나 해시를 반환할 수 있다
+
+---
+
+### D7. KernelLaunch 스키마 (PA-우선, PE-태깅된 샤드)
+
+`KernelLaunch`는 IO_CPU를 통해 대상 디바이스에서 커널을 런치하는 것을
+나타낸다.
+
+필수 필드:
+
+- 공통 envelope 필드 (D3)
+- `kernel_ref: KernelRef`
+- `args: list[KernelArg]`
+
+`KernelRef`는 반드시 다음을 가져야 한다:
+
+- `name: str`
+- `kind: "deployed" | "builtin"`
+- `deploy_pa: int | None` — 커널 바이너리가 배치된 PA("deployed"에 필수)
+- `deploy_sip: int` — 바이너리가 위치한 SIP
+- `deploy_cube: int` — 바이너리가 위치한 큐브
+- `deploy_pe: int` — 바이너리가 위치한 PE
+- `nbytes_code: int` — 커널 바이너리 크기(BW 모델링용)
+
+커널 바이너리는 MemoryWrite를 통해 디바이스 메모리에 사전 배치되어야 한다.
+KernelLaunch는 커널 소스 코드나 IR을 런치 메시지에 임베드해서는 안 된다.
+
+`KernelArg`는 PA 매핑을 통한 텐서 인자와 값을 통한 스칼라 인자를 지원한다.
+
+텐서 인자 (필수):
+
+- `arg_kind: "tensor"`
+- `tensor_pa_map: TensorPAMap`
+
+`TensorPAMap`은 반드시 다음을 가져야 한다:
+
+- `shards: list[TensorShard]`
+
+`TensorShard`는 반드시 다음을 가져야 한다 (A 방식 강제):
+
+- `sip: int`
+- `cube: int`
+- `pe: int`
+- `pa: int`
+- `nbytes: int`
+- `offset_bytes: int`
+
+스칼라 인자 (필수):
+
+- `arg_kind: "scalar"`
+- `dtype: "i32" | "i64" | "fp16" | "fp32" | "bool"`
+- `value: number | bool`
+
+KernelLaunch 선택 필드:
+
+- `grid: dict | None`
+- `meta: dict | None`
+- `failure_policy: "fail_fast" | "collect_all"` (기본값 "fail_fast")
+- `debug_label: str | None`
+
+비고:
+
+- KernelLaunch는 대용량 텐서 데이터를 임베드해서는 안 된다.
+- KernelLaunch는 오직 IO_CPU 엔드포인트에만 제출되어야 한다.
+- IO_CPU는 샤드의 (sip,cube,pe) 태그를 사용하여 내부적으로 작업을
+  팬아웃해야 한다.
+
+---
+
+## Verification Notes
+
+테스트는 다음을 검증해야 한다:
+
+- 스키마 검증이 필수 필드 누락을 거부함,
+- 결정론적 correlation/응답 매칭,
+- MemoryWrite/Read/KernelLaunch가 명시적인 홉 트레이스를 생성함,
+- 라우팅된 모든 요청은 레이턴시 > 0을 가짐.
+
+---
+
+## Links
+
+- ADR-0011 (메모리 주소 지정 — PA / VA / LA)
+- ADR-0007 (runtime_api와 sim_engine 경계)
+- ADR-0009 (커널 실행 팬아웃/집계)
+- ADR-0013 (검증 전략 — V1 메시지 스키마 검증)
+- SPEC R2, R7, R8

docs/adr-ko/ADR-0013-ver-verification-strategy.md

@@ -0,0 +1,145 @@
+# ADR-0013: 검증 전략 및 Phase 1 테스트 계획
+
+## Status
+
+Accepted
+
+## Context
+
+KernBench는 시스템 레벨 시뮬레이터이며, 그 정확성은 다음으로 정의된다:
+
+- SPEC에 정의된 불변식 준수,
+- 결정성과 디버깅 가능성,
+- 라우팅과 레이턴시의 명시적 모델링.
+
+진화하는 구현을 고려할 때, 점진적 개발을 허용하면서도 아키텍처적
+편향(drift)을 방지하는 안정적인 검증 전략이 필요하다.
+
+본 ADR은 Phase 1 검증 계획과 초기 구현에 대해 "올바른 동작"이 무엇인지를
+정의한다.
+
+---
+
+## Decision
+
+### D1. 검증은 계약 기반이다
+
+검증은 반드시 다음으로부터 도출되어야 한다:
+
+- SPEC 요구사항,
+- 채택된 ADR들.
+
+테스트는 부수적인 구현 세부사항이 아니라 아키텍처 계약을 검증해야 한다.
+
+---
+
+### D2. Phase 1 검증 범위
+
+Phase 1 검증은 다음에 초점을 둔다:
+
+- 메시지 계약 유효성 (ADR-0012),
+- IO_CPU 경계에서의 라우팅과 팬아웃 시맨틱 (ADR-0009),
+- PA-우선 메모리 주소 지정 및 샤드 태깅 (ADR-0011),
+- 핵심 레이턴시 및 트레이스 불변식 (SPEC 0.1, R2).
+
+마이크로아키텍처 정확도, 대역폭 경합, 사이클 레벨 동작은 Phase 1의
+범위에서 명시적으로 제외된다.
+
+---
+
+### D3. 필수 Phase 1 검증 케이스
+
+다음 검증 케이스는 구현에서 반드시 지원되어야 한다:
+
+#### V1. 메시지 스키마 검증
+
+- 텐서 샤드 중 어느 하나라도 `(sip, cube, pe)`가 누락된 KernelLaunch
+  요청은 반드시 거부되어야 한다.
+- 목적지/소스 배치 태그가 누락된 MemoryWrite/MemoryRead 요청은 반드시
+  거부되어야 한다.
+- Completion 결과는 반드시 `ok / error_code / error_message` 계약을
+  따라야 한다.
+
+#### V2. IO_CPU 팬아웃과 집계
+
+다음 조건이 주어졌을 때:
+
+- SIP 1개, CUBE 1개, PE 2개로 구성된 토폴로지,
+- 서로 다른 PE를 대상으로 하는 두 개의 텐서 샤드를 포함하는
+  KernelLaunch 요청,
+
+시스템은 반드시:
+
+- 단일 KernelLaunch를 IO_CPU에 제출하고,
+- 내부적으로 두 PE에 작업을 팬아웃하며,
+- 완료를 집계하여 호스트에 단일의 결정론적 완료를 반환해야 한다.
+
+#### V3. 레이턴시 및 트레이스 불변식
+
+모든 유효한 요청에 대하여:
+
+- 홉별 트레이스는 반드시 비어 있지 않아야 한다,
+- 총 레이턴시는 반드시 0보다 커야 한다,
+- 동일한 입력으로 반복 실행 시 반드시 동일한 트레이스를 생성해야 한다.
+
+#### V4. 토폴로지 독립성과 교차 도메인 커버리지
+
+검증 케이스는 다음을 포함한 다양한 토폴로지 형태에서 통과해야 한다:
+
+- 최소: (SIP 1, CUBE 1, PE 1)
+- 다중 PE: (SIP 1, CUBE 1, PE N개)
+- SIP 내 다중 CUBE: (SIP 1, CUBE M개, CUBE당 PE ≥1)
+- 다중 SIP 트레이: (SIP K개, SIP당 CUBE ≥1, CUBE당 PE ≥1)
+
+다중 CUBE 및 다중 SIP 토폴로지에 대해 Phase 1 검증은 다음에 초점을
+둔다:
+
+- 명시적 연결성(필요한 링크가 존재함),
+- 결정론적 라우팅과 제어 경로 순회,
+- 대표적인 교차 도메인 요청(CUBE 간 및 SIP 간 경로)에 대해 비어 있지
+  않은 트레이스와 레이턴시 > 0.
+
+테스트는 토폴로지 크기, 노드 ID, 링크 수를 하드코딩해서는 안 된다.
+대신 컴파일된 토폴로지 메타데이터로부터 기대값을 도출해야 한다.
+
+---
+
+### D4. Phase 1 산출물
+
+Phase 1은 다음을 포함할 수 있다:
+
+- 검증 전용 테스트 코드,
+- 토폴로지 픽스처,
+- 트레이스 검사 유틸리티.
+
+Phase 1은 다음을 요구해서는 안 된다:
+
+- 단지 테스트를 만족시키기 위한 프로덕션 코드 변경,
+- 진행을 위한 테스트의 약화 또는 제거.
+
+---
+
+### D5. Phase 2 강제
+
+Phase 2(Apply)는 반드시:
+
+- Phase 1 검증 케이스를 실행하고,
+- 검증이 실패하면 모든 변경을 롤백하며,
+- 테스트를 권위 있는 계약으로 보존해야 한다.
+
+---
+
+## Consequences
+
+- 아키텍처 정확성은 초기에 강제된다.
+- 테스트는 시스템 동작의 실행 가능한 문서로 기능한다.
+- 구현은 엄정성을 잃지 않으면서도 유연성을 유지한다.
+
+---
+
+## Links
+
+- SPEC 0.1, R2, R6
+- ADR-0011 (메모리 주소 지정 — PA / VA / LA)
+- ADR-0012 (Host ↔ IO_CPU 메시지 스키마)
+- ADR-0009 (커널 실행 시맨틱)

docs/adr-ko/ADR-0014-dev-pe-pipeline-execution-model.md

@@ -0,0 +1,441 @@
+# ADR-0014: PE 파이프라인 실행 모델
+
+## Status
+
+Accepted
+
+## Context
+
+본 ADR은 PE 내부 커널 실행 모델을 정의한다:
+
+- PE 내부 컴포넌트의 역할 분담
+- 명령 디스패치 경로 (simple / composite / epilogue를 포함한 multi-op composite)
+- TileToken 기반 자가-라우팅 파이프라인 (스케줄러는 디스패치와 완료 처리만 담당)
+- 레지스터 파일을 매개로 한 TCM 중심 데이터플로우
+- 엔진 자원 모델
+- 관측 가능성 및 트레이스 계약
+- 토폴로지 표현
+
+PE 내부 구조 (본 ADR 범위 7개 컴포넌트 + 외부 참조 2개):
+
+- `pe_cpu`, `pe_scheduler`, `pe_dma`, `pe_fetch_store`, `pe_gemm`, `pe_math`,
+  `pe_tcm` — 본 ADR에서 정의
+- `pe_mmu` — VA 모델, ADR-0011 D-VA에서 정의
+- `pe_ipcq` — 집합 통신, ADR-0023에서 정의
+
+목표는 결정론적이고 트레이스 친화적인 실행 계약을 통해 각 블록이 독립적으로
+교체 가능하도록 유지하는 것이다.
+
+## Decision
+
+### D1. PE 내부 컴포넌트의 역할
+
+**PE_CPU**
+
+- 커널 명령어 스트림 / 제어 로직을 실행한다.
+- PE 명령을 생성하여 `PE_SCHEDULER`에 제출한다 (`PeInternalTxn`을 통해).
+- 엔진 큐에 직접 작업을 넣지 않는다.
+
+**PE_SCHEDULER**
+
+- PE 내부의 유일한 디스패처.
+- `PE_CPU`로부터 명령을 수신한다. 명령 타입별 디스패치:
+  - Simple 명령 (`DmaReadCmd`, `DmaWriteCmd`, `GemmCmd`, `MathCmd`)
+    → 대상 엔진으로 직접 전달.
+  - `CompositeCmd` → `TilePlan`을 생성하고, 단일 `_feed_loop`를 통해
+    파이프라인에 타일을 공급한다 (D6).
+- composite 내부의 stage-to-stage 체이닝에는 관여하지 않는다;
+  이는 토큰 자가-라우팅(D6)으로 처리된다.
+
+**PE_DMA**
+
+- 큐브 NoC를 통해 TCM과 외부 메모리 도메인(HBM, 공유 SRAM, 큐브 간 UCIe)
+  사이의 메모리 전송을 처리한다.
+- 두 개의 실행 채널:
+  - `DMA_READ` (capacity = 1) 및 `DMA_WRITE` (capacity = 1) — D4 참조.
+- 추가 가상 채널:
+  - `vc_compute` — GEMM/MATH 타일의 load/store/writeback 트래픽.
+  - `vc_comm` — IPCQ 집합 통신 송신 데이터 (ADR-0023 D8에서 정의).
+
+**PE_FETCH_STORE**
+
+- TCM ↔ 레지스터 파일 전송 유닛.
+- 레지스터 파일 접근 시맨틱을 컴퓨트 엔진으로부터 격리하여
+  GEMM/MATH가 순수한 컴퓨트 컴포넌트로 유지되도록 한다.
+- BW 기반 레이턴시 모델; TCM 접근 경합은 `PE_TCM`의 BW 자원을 통해
+  자연스럽게 직렬화된다.
+
+**PE_GEMM**
+
+- MAC 어레이. 레지스터 파일에서 피연산자를 읽고, 결과를 레지스터 파일에
+  쓴다. `PE_TCM`에 직접 접근하지 않는다.
+
+**PE_MATH**
+
+- 원소별 / 리덕션 / SIMD 유닛. 레지스터 파일을 읽고 쓴다.
+
+**PE_TCM**
+
+- BW로 직렬화된 접근을 갖는 tightly-coupled 스크래치패드. 소유권에 따라
+  두 개의 논리 영역으로 분할된다 (D5 참조).
+
+**외부 참조 컴포넌트** (다른 곳에서 정의됨):
+
+- `pe_mmu` — 접근마다 VA→PA 변환 (ADR-0011 D-VA).
+- `pe_ipcq` — 집합 통신 링 버퍼와 피어 엔드포인트 메타데이터
+  (ADR-0023).
+
+### D2. 명령 생명주기와 큐
+
+`PE_SCHEDULER`는 세 개의 논리적 구조를 유지한다:
+
+**SubmissionQueue** — `PE_CPU`가 쓰고, 스케줄러가 소비한다.
+
+**InflightTable** — `PE_SCHEDULER`만 소유하고 변경한다; 전개된 sub-command,
+의존성 상태, 엔진 할당, 완료 상태를 추적한다.
+
+**CompletionQueue** — `PE_SCHEDULER`가 쓴다; 최종 완료 레코드를 보관한다.
+
+**Single-writer 규칙**: `PE_SCHEDULER`만이 명령 완료 상태를 변경한다.
+엔진은 명시적 이벤트 / 메시지로 완료를 보고하며, 이는 스케줄러가
+소비한다.
+
+**명령 완료**: 모든 sub-command가 완료되면 `PE_SCHEDULER`가 완료 레코드를
+발행한다.
+
+### D3. 디스패치 모드
+
+#### D3.1 Simple 명령
+
+simple 명령은 정확히 하나의 엔진 sub-command로 전개된다:
+
+- `DmaReadCmd` / `DmaWriteCmd` → `PE_DMA`
+- `GemmCmd` → `PE_GEMM`
+- `MathCmd` → `PE_MATH`
+
+흐름:
+
+```text
+PE_CPU → SubmissionQueue → PE_SCHEDULER → engine queue → engine execution
+       → completion → PE_SCHEDULER → CompletionQueue
+```
+
+#### D3.2 Composite 명령 (단일-op 타일 파이프라인)
+
+기본 `CompositeCmd`는 단일 컴퓨트 op를 타일 파이프라인 시퀀스로 실행한다:
+
+```text
+DMA_READ → FETCH (TCM → RF) → COMPUTE (GEMM | MATH) → STORE (RF → TCM) → DMA_WRITE
+```
+
+`PE_SCHEDULER`는 DMA 페이로드를 하드웨어 타일로 분할하고, 단조 증가하는
+`tile_id`를 갖는 `TileToken`을 타일마다 하나씩 발행한다.
+
+타일 의존성 (단일 타일 `t` 내부):
+
+```text
+DMA_READ(t) → FETCH(t) → COMPUTE(t) → STORE(t) → DMA_WRITE(t)
+```
+
+엔진 자원이 허용하는 한 타일 간 오버랩이 허용된다
+(D4가 제약을 규정):
+
+```text
+DMA_READ(t+1) ∥ COMPUTE(t)
+DMA_WRITE(t-1) ∥ COMPUTE(t)
+```
+
+#### D3.3 Multi-op composite (스코프를 갖는 head + epilogue)
+
+`CompositeCmd`는 `ops: tuple[OpSpec, ...]`를 운반하여 multi-op
+파이프라인을 표현할 수 있다:
+
+```python
+@dataclass(frozen=True)
+class OpSpec:
+    kind: str         # "gemm" | "math.exp" | "math.bias_add" | ...
+    scope: Scope      # "per_k_tile" | "per_output_tile" | "once"
+    ...
+```
+
+- `ops[0]` (head)이 타일 기하 구조를 정의한다 (예: head GEMM이 M/K/N
+  분할을 결정).
+- `ops[1:]` (epilogue)는 후속 stage이며 `scope`에 따라 실행 빈도가
+  결정된다:
+  - `per_k_tile` — 모든 K-리덕션 스텝마다.
+  - `per_output_tile` — 출력 타일당 한 번.
+  - `once` — 커널당 한 번.
+
+크로스-엔진 체인(예: GEMM head → MATH epilogue)은 자연스럽다 —
+각 stage는 토큰 자가-라우팅(D6)을 통해 디스패치되므로, GEMM과 MATH는
+동일한 컴퓨트 슬롯(D4)을 공유하더라도 동일 composite 내에서 직렬적으로
+참여한다.
+
+비어 있는 `ops` 형식은 레거시 단일-op 경로이다.
+
+### D4. 엔진 자원 모델
+
+**DMA 엔진**:
+
+- `DMA_READ`: `simpy.Resource(capacity=1)`.
+- `DMA_WRITE`: `simpy.Resource(capacity=1)`.
+- 두 채널은 동시에 실행된다 (READ ∥ WRITE 허용).
+- 채널 내부에서는 요청이 직렬화된다 (READ ∥ READ 불가; WRITE도 동일).
+- `vc_comm`은 IPCQ 트래픽을 위한 직교 채널로 ADR-0023 D8에서 정의됨 —
+  본 ADR 범위 밖.
+
+**컴퓨트 엔진**:
+
+- `accel_slot`: `PE_GEMM`과 `PE_MATH`가 공유하는 `simpy.Resource(capacity=1)`.
+- PE 내에서 동시에 최대 한 개의 컴퓨트 op만 실행된다.
+- Multi-op composite 체인(D3.3)은 이 슬롯을 통해 컴퓨트 stage를 직렬로
+  실행한다; 토큰 자가-라우팅(D6)이 이전 컴퓨트가 슬롯을 해제한 후에만
+  다음 stage가 시작되도록 보장한다.
+
+**엔진 완료**: 각 엔진은 완료 이벤트를 발행하며, 이는 스케줄러 /
+`PipelineContext`(D6)가 소비한다.
+
+### D5. 데이터플로우
+
+**입력 경로 (HBM 소스)**:
+
+```text
+HBM → cube NOC → PE_DMA (DMA_READ) → PE_TCM
+PE_TCM → PE_FETCH_STORE → Register File
+Register File → PE_GEMM | PE_MATH
+```
+
+**입력 경로 (공유 SRAM 소스)**:
+
+```text
+Shared SRAM → cube NOC → PE_DMA (DMA_READ) → PE_TCM
+PE_TCM → PE_FETCH_STORE → Register File
+```
+
+**출력 경로 (HBM 목적지)**:
+
+```text
+Register File → PE_FETCH_STORE → PE_TCM
+PE_TCM → PE_DMA (DMA_WRITE) → cube NOC → HBM
+```
+
+GEMM/MATH는 `PE_TCM`에 직접 접근하지 않는다 — `PE_FETCH_STORE`가
+TCM↔레지스터 파일의 유일한 게이트웨이이다. 이를 통해 TCM BW 경합이
+명시적으로 드러나며, fetch 유닛 정책(예: 프리패치)을 컴퓨트 엔진과
+독립적으로 교체할 수 있다.
+
+#### D5.1 PE_TCM 분할
+
+`PE_TCM`은 두 개의 논리 영역으로 분할된다:
+
+**SchedulerReservedTCM**
+
+- `PE_SCHEDULER`가 단독으로 소유한다.
+- composite 명령의 타일 버퍼를 보관한다.
+- `PE_SCHEDULER`가 이 영역을 분할하고, DMA_READ / COMPUTE / DMA_WRITE
+  stage마다 버퍼를 할당하며, 입출력 분리를 보장하고, 타일-버퍼 수명을
+  관리한다.
+
+**AllocatableTCM**
+
+- `PEMemAllocator`가 관리하는 범용 영역.
+- 호스트 / DP 가시 할당에 사용된다.
+
+**가시성 규칙 (강한 격리)**: `PEMemAllocator`는 `SchedulerReservedTCM`을
+보거나 그 내부에 할당해서는 안 된다. 예약 영역은 구성 시점에 할당자가
+관리하는 범위에서 제외된다.
+
+**타일 버퍼 규칙**:
+
+- 타일이 활성 수명 동안 `SchedulerReservedTCM` 내부의 입력 버퍼와 출력
+  버퍼는 겹쳐서는 안 된다.
+- 타일 버퍼는 해당 `DMA_WRITE`가 완료될 때까지 유효하다.
+- 버퍼 재사용은 소비하는 타일의 수명이 끝난 후에만 허용된다.
+
+### D6. TileToken 자가-라우팅 파이프라인
+
+composite의 stage-to-stage 진행은 스케줄러를 거치지 **않고** 일어난다.
+각 컴포넌트는 토큰의 `plan`을 사용해 토큰을 다음 stage의 컴포넌트로
+직접 전달한다:
+
+```text
+Scheduler → DMA → Fetch → GEMM → Math (epi) → Store → DMA_WB → (complete)
+              ↑ chaining: no scheduler hop                          ↑
+                                                  PipelineContext.complete_tile()
+```
+
+이는 실제 HW의 done-wire 체인을 반영한다. 스케줄러는 **초기 디스패치 +
+완료 집계**만 담당한다.
+
+#### TilePlan / Stage
+
+```python
+class StageType(Enum):
+    DMA_READ = 0
+    FETCH = 1
+    GEMM = 2
+    MATH = 3
+    STORE = 4
+    DMA_WRITE = 5
+
+@dataclass(frozen=True)
+class Stage:
+    stage_type: StageType
+    component: str         # topology node id (e.g., "sip0.cube0.pe0.pe_dma")
+    params: dict           # stage-specific parameters
+
+@dataclass(frozen=True)
+class TilePlan:
+    tile_id: int
+    stages: tuple[Stage, ...]
+```
+
+#### TileToken
+
+```python
+@dataclass
+class TileToken:
+    tile_id: int
+    pipeline_ctx: PipelineContext
+    plan: TilePlan
+    stage_idx: int
+    params: dict             # cached current stage params
+    data_op: bool = True     # op_log opt-in (ADR-0020 D4)
+```
+
+단일 소유자 불변식: 토큰은 한 시점에 정확히 한 컴포넌트가 소유한다.
+생명주기: 스케줄러가 `stage_idx=0`으로 생성 → 컴포넌트 `_process()` →
+`stage_idx` 증가 → 다음 stage의 `in_port`에 put → 마지막 stage가
+`pipeline_ctx.complete_tile()` 호출.
+
+#### PipelineContext (정확히 한 번 완료)
+
+```python
+@dataclass
+class PipelineContext:
+    id: str
+    total_tiles: int
+    completed_tiles: int = 0
+    done_event: simpy.Event = None
+
+    def complete_tile(self) -> None:
+        self.completed_tiles += 1
+        if self.completed_tiles == self.total_tiles:
+            self.done_event.succeed()
+```
+
+각 타일의 마지막 stage는 `complete_tile()`을 정확히 한 번 호출해야
+한다. 중복 호출은 버그이다 (SimPy `Event`는 최대 한 번만 succeed
+가능).
+
+#### Feed 순서
+
+`PE_SCHEDULER`는 `_pending_feeds` FIFO를 소비하는 `_feed_loop` 프로세스를
+정확히 하나 갖는다. composite 명령은 제출 순서대로 인큐되며, 한 명령의
+타일 feed는 다음 명령의 feed가 시작되기 전에 완료까지 실행된다.
+**명령 간 타일-feed 인터리빙은 허용되지 않는다.**
+
+단일 명령의 타일들 내부에서는 다운스트림 파이프라인 오버랩이 자연스럽게
+발생한다 — 이전 타일이 후행 stage를 진행하는 동안 feeder는 남은 타일을
+첫 stage 큐로 계속 푸시한다 (SimPy Store 백프레셔가 흐름 제어를
+관장한다). 첫 stage 큐가 가득 차면 feeder만 블록되며, 스케줄러 워커의
+inbox 처리는 계속된다.
+
+#### 토큰 라우팅 패턴 (기본 클래스)
+
+```python
+def _pipeline_worker(self, env):
+    while True:
+        token = yield self._inbox.get()
+        yield from self._process(env, token)       # stage-specific logic
+        next_idx = token.stage_idx + 1
+        if next_idx < len(token.plan.stages):
+            next_stage = token.plan.stages[next_idx]
+            token.stage_idx = next_idx
+            token.params = next_stage.params
+            yield self.out_ports[next_stage.component].put(token)
+        else:
+            token.pipeline_ctx.complete_tile()
+```
+
+각 컴포넌트는 `_process()`만 구현한다; 체이닝은 기본 클래스에 존재한다.
+
+### D7. 관측 가능성 및 트레이스 계약
+
+시뮬레이터는 결정론적 트레이스 이벤트를 발행한다:
+
+- `command_submitted`
+- `sub_command_dispatched`
+- `engine_start`
+- `engine_complete`
+- `tile_ready`
+- `command_complete`
+
+동일한 입력에 대해 트레이스 순서는 결정론적이어야 한다.
+
+### D8. 토폴로지 표현
+
+PE 내부 컴포넌트는 `cube.pe_template`에 선언된다:
+
+```yaml
+pe_template:
+  components:
+    pe_cpu:         { kind: pe_cpu,         impl: builtin.pe_cpu,         attrs: { overhead_ns: ... } }
+    pe_scheduler:   { kind: pe_scheduler,   impl: builtin.pe_scheduler,   attrs: { overhead_ns: ... } }
+    pe_dma:         { kind: pe_dma,         impl: builtin.pe_dma,         attrs: { rd_engines: 1, wr_engines: 1 } }
+    pe_fetch_store: { kind: pe_fetch_store, impl: builtin.pe_fetch_store, attrs: { ... } }
+    pe_gemm:        { kind: pe_gemm,        impl: builtin.pe_gemm,        attrs: { shared_resource: accel_slot, ... } }
+    pe_math:        { kind: pe_math,        impl: builtin.pe_math,        attrs: { shared_resource: accel_slot, ... } }
+    pe_tcm:         { kind: pe_tcm,         impl: builtin.pe_tcm,         attrs: { size_mb: ..., read_bw_gbs: ..., write_bw_gbs: ... } }
+    pe_mmu:         { kind: pe_mmu,         impl: builtin.pe_mmu,         attrs: { ... } }   # ADR-0011 D-VA
+    pe_ipcq:        { kind: pe_ipcq,        impl: builtin.pe_ipcq,        attrs: { ... } }   # ADR-0023
+  links:
+    # Scheduler dispatch edges (initial)
+    scheduler_to_dma_mm:         0.0
+    scheduler_to_fetch_store_mm: 0.0
+    scheduler_to_gemm_mm:        0.0
+    scheduler_to_math_mm:        0.0
+    # Pipeline chaining edges (token self-routing per D6)
+    dma_to_fetch_store_mm:       0.0
+    fetch_store_to_gemm_mm:      0.0
+    fetch_store_to_math_mm:      0.0
+    gemm_to_fetch_store_mm:      0.0
+    gemm_to_math_mm:             0.0
+    math_to_fetch_store_mm:      0.0
+    fetch_store_to_dma_mm:       0.0
+    fetch_store_to_tcm_bw_gbs:   ...
+```
+
+템플릿은 PE마다 한 번 인스턴스화된다. PE 인스턴스는 `cube.pe_layout`
+(코너 배치)으로부터 파생된다. 외부 연결성(PE_DMA ↔ cube NoC ↔ HBM 등)은
+큐브 수준에서 모델링된다 (ADR-0017 D4).
+
+## Consequences
+
+### Positive
+
+- 각 블록이 독립적인 토폴로지 노드이다 — DI(ADR-0015)를 통해 개별
+  교체 가능하다.
+- PE 내부 구조가 토폴로지 그래프에 가시화된다.
+- 컴포넌트는 자신의 다운스트림을 알지 못한다 — plan 기반 라우팅이
+  유연성을 제공한다 (예: epilogue 체인에 스케줄러 변경이 불필요).
+- DMA와 컴퓨트가 SimPy Store 백프레셔를 통해 자연스럽게 오버랩된다.
+- Multi-op composite가 융합 연산(예: GEMM + bias_add)을 엔진 수준
+  결합 없이 표현한다.
+- TCM 접근 경합이 현실적이다 — `PE_FETCH_STORE`가 TCM↔RF의 유일한
+  게이트웨이이다.
+
+### Negative
+
+- PE 내부 컴포넌트 수가 더 거친 모델보다 많다 (기본 7개 + 외부 참조
+  2개) — 더 많은 토폴로지 노드/엣지.
+- PE 내부 토큰 전달이 트레이스에 명시적으로 드러난다 (HW 충실도와의
+  허용 가능한 trade-off).
+
+## Links
+
+- ADR-0011 D-VA (PE_MMU 컴포넌트, VA 변환)
+- ADR-0015 D4 (컴포넌트 포트/와이어 모델)
+- ADR-0020 (greenlet 커널 실행 / two-pass)
+- ADR-0023 (PE_IPCQ + PE_DMA 가상 채널)
+- SPEC R3, R4

docs/adr-ko/ADR-0015-dev-component-port-wire-model.md

@@ -0,0 +1,202 @@
+# ADR-0015: 컴포넌트 포트/와이어 모델과 패브릭 라우팅
+
+## Status
+
+Accepted
+
+## Context
+
+현실적인 하드웨어 모델링 — 큐, 경합, fan-out — 을 위해서는
+컴포넌트가 패브릭 순회를 소유하고, 시뮬레이션 엔진은 초기화와 완료
+관측만 처리해야 한다. 컴포넌트 간의 직접 메서드 호출이나 엔진 내부의
+경로 탐색은 큐잉과 경합 시맨틱을 무력화한다.
+
+본 ADR은 다음을 정의한다:
+
+- 컴포넌트가 타입드 포트 큐를 통해 통신하는 방식,
+- 전파 지연을 모델링하는 방식 (BW 점유를 포함한 와이어 프로세스),
+- Memory R/W (M_CPU 우회)와 Kernel Launch (M_CPU 경유)의 패브릭 경로,
+- 엔진의 축소된 역할 (와이어 초기화 + 완료 관측만),
+- M_CPU의 내부 서브컴포넌트로서의 M_CPU.DMA.
+
+---
+
+## Decision
+
+### D1. 컴포넌트 포트 모델
+
+각 컴포넌트는 SimPy Store로 모델링된 타입드 입출력 포트를 갖는다:
+
+```text
+in_ports:  dict[str, simpy.Store]   # keyed by source node_id
+out_ports: dict[str, simpy.Store]   # keyed by destination node_id
+```
+
+포트는 그래프 엣지를 기반으로 엔진 초기화 시 생성된다.
+각 유향 엣지(src → dst)는 다음을 생성한다:
+
+- `src.out_ports[dst]`  — 송신측
+- `dst.in_ports[src]`   — 수신측
+
+---
+
+### D2. 와이어 프로세스 (전파 지연 + BW 점유)
+
+토폴로지 그래프의 각 유향 엣지 (src, dst)에 대해 SimPy 와이어 프로세스가
+전파 지연과 BW 점유를 모델링한다:
+
+```python
+def wire_process(env, out_port, in_port, delay_ns, bw_gbs):
+    available_at = 0.0
+    while True:
+        cmd = yield out_port.get()
+        if bw_gbs > 0:
+            nbytes = getattr(cmd, "nbytes", 0)
+            if nbytes > 0:
+                wait = available_at - env.now
+                if wait > 0:
+                    yield env.timeout(wait)
+                available_at = env.now + (nbytes / bw_gbs)
+        yield env.timeout(delay_ns)
+        yield in_port.put(cmd)
+```
+
+와이어 프로세스는 엔진 초기화 시점에 시작된다.
+각 유향 엣지는 링크가 다음 트랜잭션을 위해 비워지는 시점을 추적하는
+`available_at` 타임스탬프를 유지한다. 한 트랜잭션이 링크를 점유하는 동안,
+동일 유향 링크의 다음 트랜잭션은 점유가 해제될 때까지 대기해야 한다
+(연속 직렬화). TX와 RX 방향은 독립적이다 (각각의 `available_at` 상태를
+갖는 별개의 와이어 프로세스).
+
+---
+
+### D3. 엔진 역할 (축소)
+
+시뮬레이션 엔진은 다음을 수행해야 한다:
+
+- 초기화 시점에 컴포넌트 와이어링 (포트 Store 생성, 와이어 프로세스 시작),
+- 각 요청 타입별 진입 컴포넌트 식별 (PCIE_EP),
+- 진입 컴포넌트의 in_port에 요청을 put,
+- 완료 이벤트 대기.
+
+시뮬레이션 엔진은 다음을 해서는 안 된다:
+
+- 요청 실행 중 토폴로지 경로 탐색,
+- 컴포넌트 `run()` 메서드 직접 호출,
+- hop별 레이턴시 추적이나 fan-out 분해.
+
+---
+
+### D4. Memory R/W와 Kernel Launch의 패브릭 경로
+
+Memory R/W와 Kernel Launch는 **서로 다른** 패브릭 경로를 사용한다.
+메모리 연산은 M_CPU를 우회하여 크로스바를 통해 직접 HBM으로 라우팅된다.
+Kernel Launch는 PE fan-out을 위해 M_CPU를 경유한다.
+
+**Memory R/W forward 경로 (pcie_ep → hbm_ctrl, M_CPU 우회):**
+
+```text
+pcie_ep → io_noc → io_ucie
+  → [transit cubes: ucie_in → noc → ucie_out]  (zero or more)
+  → target cube: ucie_in → router mesh → hbm_ctrl
+```
+
+**Memory R/W 완료 경로:**
+
+```text
+hbm_ctrl → router mesh → [transit cubes: ucie → router mesh → ucie]
+  → io_ucie → io_noc → pcie_ep
+```
+
+**Kernel Launch forward 경로 (pcie_ep → io_cpu → M_CPU → PE):**
+
+```text
+pcie_ep → io_noc → io_cpu → io_noc → io_ucie
+  → [transit cubes: ucie_in → noc → ucie_out]  (zero or more)
+  → target cube: ucie_in → noc → M_CPU → PE[0..n] (parallel fan-out)
+```
+
+**Kernel Launch 완료 경로:**
+
+```text
+PE[0..n] all complete → M_CPU (aggregation)
+  → noc → [transit cubes: ucie → noc → ucie]
+  → io_ucie → io_noc → io_cpu → io_noc → pcie_ep
+```
+
+**Memory R/W가 M_CPU를 우회하는 근거:**
+
+메모리 write/read 연산은 명령 해석이나 PE 디스패치가 필요하지 않다 —
+HBM으로의/로부터의 직접 데이터 전송이다. M_CPU를 경유하면 기능적 이득
+없이 불필요한 오버헤드(5ns)를 추가한다. IO 칩렛 내부의 io_noc가 라우팅
+결정을 처리한다: 메모리 연산은 큐브 패브릭으로 직접 가고, kernel
+launch는 io_cpu로 먼저 전달된다.
+
+---
+
+### D5. M_CPU.DMA는 M_CPU의 내부 서브컴포넌트이다
+
+M_CPU.DMA는 별개의 토폴로지 노드가 아니다.
+M_CPU 컴포넌트 구현이 소유하는 내부 서브컴포넌트이다.
+
+M_CPU.DMA는:
+
+- DMA READ 및 DMA WRITE 큐를 소유한다 (각 capacity=1, ADR-0014 D4),
+- NoC를 통해 hbm_ctrl에 메모리 요청을 발행한다,
+- NoC를 통해 hbm_ctrl로부터 완료를 수신한다,
+- M_CPU에 완료를 보고한다,
+- M_CPU의 `__init__`과 `run()` 내부에서 생성·관리된다.
+
+M_CPU.DMA는 컴파일된 토폴로지 그래프에서 노드로 나타나지 않는다.
+
+---
+
+### D6. Transit 큐브 포워딩
+
+메모리나 커널 요청의 대상이 아닌 큐브는 transit 노드로 동작한다.
+Transit 큐브는 요청을 소비하지 않고 포워딩한다:
+
+```text
+ucie_in (from upstream) → noc → ucie_out (to downstream)
+```
+
+Transit 포워딩은 ucie_in 컴포넌트 내부에서 전적으로 구현된다.
+transit 큐브의 noc와 ucie_out 컴포넌트는 패킷을 수정 없이 포워딩한다.
+
+---
+
+### D7. _formula_latency는 하한 교차 검증 용도로 유지된다
+
+경로 기반 공식 레이턴시 함수(`_formula_latency`)는 정확성 검증을 위한
+하한값으로 엔진 내에 유지된다.
+
+불변식:
+
+- Phase 0: `_formula_latency == component model total_ns`
+- Phase 1+: `_formula_latency <= component model total_ns` (경합이
+  큐잉을 추가)
+
+이 함수는 포트/와이어 모델과 독립적이며 토폴로지 그래프만 요구한다.
+`_route_kernel`의 샤드 비교와 회귀 가드로 사용된다.
+
+---
+
+## Consequences
+
+- 컴포넌트가 현실적인 하드웨어 동작(큐, 경합, fan-out)을 모델링한다.
+- 전파 지연이 엣지마다 정확하게 모델링된다.
+- 엔진이 라우팅 정책으로부터 분리된다.
+- 컴포넌트 구현이 DI(ADR-0007 D3)를 통해 교체 가능하게 유지된다.
+
+---
+
+## Links
+
+- ADR-0007 D2 (엔진 역할 경계)
+- ADR-0009 D3 (커널 실행 fan-out 계층)
+- ADR-0014 D4 (DMA 엔진 capacity=1)
+- ADR-0012 D1 (호스트 ↔ IO_CPU 메시지 스키마; M_CPU.DMA는 컴포넌트
+  내부)
+- ADR-0016 (IOChiplet NoC와 메모리 데이터 경로)
+- ADR-0017 (큐브 NoC 2D 메시 아키텍처)
+- ADR-0033 (이러한 메커니즘 위에 구축된 레이턴시 모델 가정)

docs/adr-ko/ADR-0016-dev-iochiplet-noc-and-memory-path.md

@@ -0,0 +1,99 @@
+# ADR-0016: IOChiplet NoC와 메모리 데이터 경로
+
+## Status
+
+Accepted
+
+## Context
+
+ADR-0003 D2는 IO chiplet을 PCIe-EP 및 IO_CPU 인터페이스를 제공하는 SIP
+레벨 컴포넌트로 정의하지만, IO chiplet 내부의 라우팅은 명세하지 않는다.
+ADR-0015 D4는 Memory R/W에 대한 M_CPU 우회를 문서화하도록 갱신되었지만,
+이 라우팅을 가능하게 하는 IO chiplet의 내부 NoC 아키텍처는 형식적으로
+문서화되지 않았다.
+
+IO chiplet은 다음을 위해 내부 라우팅 패브릭(io_noc)을 필요로 한다:
+
+- pcie_ep, io_cpu, 그리고 큐브당 UCIe PHY 포트들을 연결한다
+- 메모리 연산(MemoryWrite/Read)을 io_cpu를 거치지 않고 큐브 패브릭으로
+  직접 라우팅한다
+- 커널 런치 명령을 명령 해석을 위해 io_cpu를 통해 라우팅한다
+
+## Decision
+
+### D1. IOChiplet 내부 NoC (io_noc)
+
+각 IO chiplet 인스턴스는 다음을 연결하는 내부 NoC 노드(`io_noc`)를
+포함한다:
+
+- `pcie_ep` — 호스트 대면 PCIe 엔드포인트
+- `io_cpu` — 커널 런치 해석용 명령 프로세서
+- `io_ucie-{PHY}.conn{N}` — 큐브 UCIe 포트들로 가는 PHY별 연결 노드
+
+io_noc은 오버헤드가 0인 포워딩 전용 패브릭(`forwarding_v1` 구현)이다.
+모든 라우팅 결정은 io_noc 자체가 아니라 메시지 타입에 기반하여 시뮬레이션
+엔진이 내린다.
+
+### D2. IOChiplet UCIe 분해
+
+각 IO chiplet PHY 포트는 다음으로 분해된다:
+
+- `io_ucie-{PHY}` — UCIe 프로토콜 엔드포인트(overhead = 8ns)
+- `io_ucie-{PHY}.conn{N}` — io_noc과 io_ucie 사이의 N개 연결 노드
+
+이는 큐브 측 UCIe 분해(ADR-0015 D1)를 미러링하며, PHY당 여러 독립적인
+NoC-UCIe 연결을 허용한다.
+
+### D3. Memory R/W 경로 (M_CPU 우회)
+
+메모리 연산(MemoryWrite, MemoryRead)은 pcie_ep에서 io_noc을 거쳐 대상
+큐브로 직접 라우팅되며, io_cpu를 완전히 우회한다:
+
+```text
+pcie_ep → io_noc → conn → io_ucie → [cube UCIe] → router mesh → hbm_ctrl
+```
+
+이는 순수 데이터 전송에 대해 10ns의 io_cpu 오버헤드를 회피한다.
+시뮬레이션 엔진의 `_process_memory_direct()` 메서드는 pcie_ep에서 대상
+HBM 노드까지의 최단 경로를 해석하는 `find_memory_path()`를 사용한다.
+
+### D4. 커널 런치 경로 (io_cpu 경유)
+
+커널 런치 명령은 명령 해석 및 PE 팬아웃 설정을 위해 io_cpu를 필요로
+한다:
+
+```text
+pcie_ep → io_noc → io_cpu → io_noc → conn → io_ucie → [cube UCIe]
+  → noc → m_cpu → PE
+```
+
+엔진의 `_entry_points()` 메서드는 KernelLaunchMsg를 pcie_ep(진입)와
+io_cpu(명령 처리) 양쪽 모두를 통해 라우팅한다.
+
+### D5. IOChiplet-to-큐브 포트 매핑
+
+각 IO chiplet 인스턴스는 자신이 연결되는 큐브 포트를 선언한다:
+
+```yaml
+cube_ports:
+  - { cube: {xy: [0,0]}, cube_side: N, phy: P0, distance_mm: 2.0 }
+  - { cube: {xy: [1,0]}, cube_side: N, phy: P1, distance_mm: 2.0 }
+```
+
+토폴로지 빌더는 io_ucie PHY 노드에서 해당 큐브 UCIe 포트 노드로의 엣지를
+지정된 거리 및 IO chiplet의 `per_connection_bw_gbs`를 링크 대역폭으로
+하여 생성한다.
+
+## Consequences
+
+- IO chiplet은 잘 정의된 내부 라우팅 패브릭을 가진다
+- 메모리 연산은 불필요한 io_cpu 오버헤드를 회피한다
+- 커널 런치 명령은 여전히 적절한 명령 해석을 받는다
+- io_noc 패턴은 큐브 레벨 NoC 설계와 일관된다
+- ADR-0003 D2는 본 ADR에 의해 확장된다(모순되지 않는다)
+
+## Links
+
+- ADR-0003 D2 (IO chiplet 정의)
+- ADR-0015 D4 (Memory R/W 및 커널 런치의 패브릭 경로)
+- ADR-0012 D1 (호스트-IO_CPU 메시지 스키마)

docs/adr-ko/ADR-0017-dev-cube-noc-and-hbm-connectivity.md

@@ -0,0 +1,282 @@
+# ADR-0017: 큐브 NoC와 HBM 연결성
+
+## Status
+
+Accepted
+
+## Context
+
+CUBE 레벨의 NoC는 모든 큐브 내부 요청을 운반하는 2D 라우터 메시이다:
+PE-HBM 데이터, PE-PE 트래픽, 명령 경로(M_CPU↔PE_CPU), 공유 SRAM 접근,
+큐브 간 UCIe 트래픽.
+
+CUBE의 HBM은 PE 라우터에 부착된 PE별 컨트롤러 엔드포인트를 통해 노출된다.
+이러한 PE별 분할 덕분에 로컬-vs-원격 HBM이 메시 거리로 구분 가능하다:
+PE 자신의 HBM 파티션은 자신의 라우터에 위치하고(스위칭 오버헤드만 발생),
+다른 PE의 HBM 파티션은 해당 PE의 라우터로 메시 hop을 거쳐 도달 가능하다.
+
+설계 공간에서는 두 가지 채널 매핑 모드를 지원한다:
+
+- **n:1 (default, 구현됨)** — 각 PE의 HBM 파티션이 `channels_per_pe`
+  pseudo-channel을 하나의 엔드포인트로 집계한다. 유효 PE당 BW =
+  N × per-channel BW.
+- **1:1 (future)** — 각 PE 라우터가 채널별 미니 라우터로 분해된다;
+  채널별 BW 경합을 직접 모델링한다.
+
+두 모드 모두 PE당 유효 BW는 동일하다; 연결 입도만 다르다.
+
+## Decision
+
+### D1. 2D 라우터 메시
+
+각 큐브는 `mesh_gen.py`가 생성하는 2D 라우터 메시를 포함한다.
+
+- 노드 명명: `sip{S}.cube{C}.r{row}c{col}` (예: `sip0.cube0.r0c0`).
+- 구현: `forwarding_v1`. NoC `overhead_ns = 0`.
+- 기본 6×6 그리드 (PE 코너 배치 + UCIe 부착 개수로 산정); 더 큰 PE
+  개수는 그리드를 확장한다.
+- HBM 제외 영역: HBM 다이가 물리적으로 점유하는 중앙 행/열을 제외한다
+  (예: 6×6의 경우 r2c2, r2c3, r3c2, r3c3).
+- 레이턴시 = Manhattan 거리 × `ns_per_mm`.
+
+### D2. XY 라우팅 알고리즘
+
+결정론적 XY 라우팅:
+
+1. 수평 구간: 소스 X에서 목적지 X까지 소스 Y에서 라우팅.
+2. 수직 구간: 소스 Y의 목적지 X에서 목적지 Y까지 라우팅.
+
+각 유향 구간은 고유 키를 운반한다:
+
+- 수평: `("H", y_band, x_min, x_max, direction)`
+- 수직:   `("V", x_band, y_min, y_max, direction)`
+
+그리드 위치는 HBM 영역을 제외하고 라우터 그리드에 스냅된다.
+
+### D3. 구간별 경합 모델
+
+각 유향 XY 구간은 `simpy.Resource(capacity=1)`이다. 동일 구간을 공유하는
+트랜잭션(동일한 행 또는 열 밴드, 동일한 방향)은 자원을 두고 경합한다 —
+wormhole 라우팅 메시에서의 링크 수준 직렬화를 모델링한다.
+
+경합이 없을 때 NoC 순회 레이턴시는 Manhattan 거리 × `ns_per_mm`이다.
+경합이 있을 때는 SimPy의 자원 스케줄링이 큐잉 지연을 추가한다.
+
+### D4. NoC 부착 지점 (PE별 HBM 파티션)
+
+모든 PE 라우터는 세 개의 부착을 갖는다: `pe{idx}.dma`, `pe{idx}.cpu`,
+그리고 `pe{idx}.hbm`. 마지막은 PE별 HBM 컨트롤러 엔드포인트로
+`sip{S}.cube{C}.hbm_ctrl.pe{idx}`이며, 큐브 HBM의 한 슬라이스를
+소유한다 (하나의 pseudo-channel 그룹; D8 참조).
+
+기타 부착:
+
+- M_CPU와 공유 SRAM은 각각 전용 edge 라우터를 점유한다.
+- UCIe 엔드포인트(N/S/E/W)는 각각 해당 변에 분산된 4개의 연결 라우터를
+  노출한다 (D6 참조).
+
+```text
+                    UCIe-N (conn x4)
+                         |
+           +---------+---+---+---------+
+           |         |       |         |
+PE0.dma ---+  r0c0   |  ...  |  r0c5  +--- PE2.dma
+PE0.cpu <--+ +hbm.pe0|       | +hbm.pe2+--< PE2.cpu
+           |         |       |         |
+UCIe-W ----+  ...    | [HBM] |  ...   +---- UCIe-E
+(conn x4)  |         | zone  |         |  (conn x4)
+           |  r2c0   |       |         |
+M_CPU <--->+         |       |         |
+           |  r3c0   |       |         |
+SRAM <---->+         |       |         |
+           |         |       |         |
+PE4.dma ---+  r4c0   |  ...  |  r4c5  +--- PE6.dma
+PE4.cpu <--+ +hbm.pe4|       | +hbm.pe6+--< PE6.cpu
+           |         |       |         |
+           +---------+---+---+---------+
+                         |
+                    UCIe-S (conn x4)
+```
+
+PE별 HBM 분할은 로컬 vs 크로스-PE HBM을 메시 거리로 구분 가능하게 만드는
+핵심 불변식이다 (D7 참조).
+
+### D5. NoC 엣지 대역폭과 거리
+
+| Connection                    | BW (GB/s)  | Distance      | Notes                                       |
+| ----------------------------- | ---------- | ------------- | ------------------------------------------- |
+| PE_DMA → NOC                  | 256.0      | Physical (PE) | 로컬-HBM 집계 BW와 일치                     |
+| NOC → PE_CPU                  | —          | 0.0 mm        | 명령 경로 전용                              |
+| Router ↔ hbm_ctrl.pe{idx}     | 256.0      | 0.0 mm        | PE 라우터당; N × per-channel BW (D8 참조)   |
+| NOC ↔ M_CPU                   | —          | 0.0 mm        | 명령 경로                                   |
+| NOC ↔ SRAM                    | 128.0 × 4  | 0.0 mm        | 512 GB/s 집계                               |
+| NOC ↔ UCIe conn               | 128.0      | 0.0 mm        | 연결당; 포트당 4개 conn                     |
+
+`0.0 mm` 거리는 NoC의 분산 특성을 반영한다; 실제 순회 거리는 라우터
+그리드 내에서 Manhattan 거리로 계산된다.
+
+### D6. UCIe 분해와 큐브 간 트래픽
+
+4개의 UCIe 포트(N, S, E, W) 각각은 다음으로 분해된다:
+
+- `ucie-{PORT}` 노드 1개: UCIe 프로토콜 엔드포인트 (`overhead = 8.0 ns`).
+- `ucie-{PORT}.conn{0-3}` 노드 4개: NoC와 UCIe 간 연결 브리지.
+
+이 분해로 포트당 4개의 독립 NoC↔UCIe 연결이 생성되며, 각각 128 GB/s
+대역폭을 갖는다 (포트당 집계 512 GB/s).
+
+큐브 간 트래픽 경로:
+
+```text
+Source: PE_DMA → NOC → conn{i} → ucie-{PORT}
+                  [UCIe link: 512 GB/s, 1.0mm seam distance]
+Target: ucie-{PORT} → conn{i} → r{x}c{y} → (mesh hops) → hbm_ctrl.pe{idx}
+```
+
+UCIe 오버헤드(8.0 ns)는 각 `ucie-{PORT}` 노드에서 적용되므로 전체 횡단은
+16 ns(TX 포트 + RX 포트)가 소요된다.
+
+### D7. NoC를 통한 데이터 경로
+
+모든 큐브 내부 트래픽은 동일한 라우터 메시를 사용한다 — 별도의 fast path는
+없다.
+
+**로컬 HBM** (동일 PE의 자신 파티션; 0 메시 hop):
+
+```text
+PE_DMA → r{x}c{y} → hbm_ctrl.pe{idx}   (switching overhead only)
+```
+
+**큐브 내 크로스-PE HBM** (대상 PE의 파티션, 메시로 도달):
+
+```text
+PE_DMA → r{x}c{y} → (mesh hops) → r{x'}c{y'} → hbm_ctrl.pe{idx'}
+```
+
+예시: PE0(`r0c0` 위)이 PE2의 HBM(PE2는 `r1c4` 위)에 접근:
+
+```text
+PE0.pe_dma → r0c0 → r0c1 → r0c2 → r0c3 → r0c4 → r1c4 → hbm_ctrl.pe2
+```
+
+Dijkstra가 메시 내 최단 경로를 계산한다.
+
+**큐브 간 HBM** (UCIe 횡단):
+
+```text
+PE_DMA → r{x}c{y} → conn → ucie-{PORT} → [seam] → ucie-{PORT'} → conn
+       → r{x'}c{y'} → hbm_ctrl.pe{idx'}
+```
+
+**PE로의 커널 launch 명령**:
+
+```text
+[from io_noc] → ucie → conn → r{x}c{y} → (mesh) → M_CPU → (mesh) → PE_CPU
+```
+
+**공유 SRAM 접근**:
+
+```text
+PE_DMA → r{x}c{y} → (mesh) → SRAM
+```
+
+### D8. HBM 채널 매핑 모드
+
+채널 매핑은 큐브 범위에서 구성된다:
+
+```yaml
+cube:
+  memory_map:
+    hbm_mapping_mode: n_to_one       # one_to_one | n_to_one
+    hbm_pseudo_channels: 64          # total pseudo-channel count
+    hbm_channels_per_pe: 8           # per-PE local channel count
+    hbm_channel_bw_gbs: 32.0         # per-channel bandwidth (GB/s)
+    hbm_slices_per_cube: 8           # number of per-PE partitions
+    hbm_total_gb_per_cube: 48
+```
+
+**n:1 모드 (default, 구현됨).** 각 PE의 HBM 파티션은 `channels_per_pe`
+pseudo-channel을 집계하는 단일 엔드포인트 `hbm_ctrl.pe{idx}`이다.
+`Router ↔ hbm_ctrl.pe{idx}` 링크 대역폭은 `channels_per_pe ×
+hbm_channel_bw_gbs`와 같다. Pseudo-channel은 인터리브된다고 가정하며,
+PE당 집계 BW만 모델링한다. 별도의 집계 라우터 노드는 존재하지 않는다 —
+PE별 라우터 자체가 그 역할을 한다.
+
+**1:1 모드 (future).** 각 PE 라우터가 N개의 채널 미니 라우터로
+분해된다; 채널별 라우팅이 완전히 해석된 PA + channel ID를 운반한다.
+`ChannelSplitter`가 논리적 접근을 N개의 채널별 물리 요청으로 해결한다.
+채널별 링크가 BW 경합을 모델링한다. 크로스-PE 채널 접근 시맨틱은
+구현 ADR로 연기된다.
+
+**BW 계산 (default 값).**
+
+| Parameter                          | Value                      |
+| ---------------------------------- | -------------------------- |
+| 큐브당 pseudo channel              | 64 (parameter)             |
+| 큐브당 PE                          | 8 (parameter)              |
+| PE당 channel (N)                   | 64 / 8 = 8                 |
+| 채널당 BW                          | 32 GB/s (parameter)        |
+| PE당 로컬 BW                       | N × 32 = 256 GB/s          |
+| 큐브 전체 HBM BW                   | 64 × 32 = 2048 GB/s        |
+
+두 모드 모두 PE당 유효 BW는 동일하다; 요청 형태와 경합 모델만 다르다.
+
+### D9. AddressResolver — PE별 HBM 엔드포인트
+
+주소 리졸버는 PA의 HBM 오프셋을 소유 PE의 파티션으로 디코딩한다:
+
+```python
+# policy/routing/router.py
+hbm_slice_bytes = hbm_total_gb_per_cube * (1 << 30) // hbm_slices_per_cube
+
+if addr.kind == "hbm":
+    pe_id = int(addr.hbm_offset) // hbm_slice_bytes
+    return f"sip{s}.cube{d}.hbm_ctrl.pe{pe_id}"
+```
+
+pe_id 계산은 라우팅 레이어의 본질적 일부이다 (토폴로지 시점 관심사가
+아니다). 모든 HBM PA는 정확히 하나의 파티션에 속하므로 결정론적 라우팅이
+보장된다.
+
+외부 호출자(예: M_CPU DMA, PCIE_EP로부터의 Memory R/W)도 동일한 리졸버
+경로를 따른다 — 별도의 fast path는 존재하지 않는다.
+
+### D10. 메시 생성 파라미터
+
+`mesh_gen.py`는 다음으로부터 `cube_mesh.yaml`을 생성한다:
+
+- `cube.pe_layout`: 코너 배치(NW, NE, SW, SE)와 코너당 PE 개수.
+- `cube.geometry`: 큐브 물리 치수와 HBM 영역.
+- `cube.ucie.n_connections`: UCIe 부착용 라우터 개수를 결정.
+
+출력 `mesh_data` 딕셔너리는 다음을 포함한다:
+
+- 위치 및 HBM 제외 영역을 갖는 라우터 그리드.
+- PE-라우터 부착 (PE별 `pe{idx}.dma`, `pe{idx}.cpu`, `pe{idx}.hbm`).
+- UCIe-라우터 부착 (N/S/E/W가 edge 라우터에 분산).
+- M_CPU와 SRAM 라우터 부착.
+
+## Consequences
+
+- 로컬 HBM(0 메시 hop, 스위칭 오버헤드만)과 크로스-PE HBM(메시 hop)이
+  자연스럽게 구분되어 SPEC R5(다중 도메인 통신)와 ADR-0002(end-to-end
+  제로 레이턴시 경로 금지)를 만족한다.
+- 모든 큐브 내부 트래픽이 하나의 메시를 통해 라우팅된다 — 단일 경합
+  모델, 단일 레이아웃, 단일 엣지 BW 집합.
+- PE별 HBM 분할이 LA 모델(ADR-0011)에 깔끔하게 매핑된다: 각 PE의
+  파티션은 할당된 pseudo-channel의 n:1 집계이다.
+- 1:1 모드 확장이 구조적으로 자연스럽다 — 각 PE 라우터를 N개의 채널
+  라우터로 분해한다.
+- 메시 생성이 `topology.yaml`로 완전히 파라미터화된다; PE/큐브 기하
+  변경이 코드 수정 없이 전파된다.
+
+## Links
+
+- ADR-0002 (라우팅 거리, 순서, 제로 레이턴시 경로 금지)
+- ADR-0003 D3 (큐브 레벨 NoC 정의 — 본 ADR에서 확장)
+- ADR-0004 (메모리 시맨틱, 로컬 HBM)
+- ADR-0011 (메모리 주소 지정 — LA 모델이 PE별 파티션을 소비)
+- ADR-0014 D1 (라우터 메시를 통한 PE_DMA egress)
+- ADR-0015 D4 (Memory R/W와 Kernel Launch의 패브릭 경로)
+- ADR-0016 (IOChiplet io_noc — IO 칩렛 레벨에서의 유사 패턴)
+- ADR-0033 (레이턴시 모델: PC당 병렬성, 스위치 패널티)

docs/adr-ko/ADR-0020-prog-data-execution-two-pass.md

@@ -2,7 +2,7 @@
 
 ## Status
 
-Proposed
+Accepted
 
 ## Context
 
@@ -16,21 +16,6 @@ Proposed
 2. PE_GEMM, PE_MATH가 실제 행렬 연산을 수행하고 결과를 검증할 수 있어야 한다
 3. 시뮬레이션 성능 저하를 최소화해야 한다
 
-### 기존 커널 실행 구조의 한계
-
-현재 커널 실행은 3단계로 분리되어 있다:
-
-```
-Phase 0: TLContext에서 커널 함수 실행 → PeCommand 리스트 생성 (SimPy 밖, 데이터 없음)
-Phase 1: PE_CPU가 PeCommand 리스트를 SimPy로 replay (타이밍만)
-```
-
-Phase 0에서 커널이 **전부 실행 완료**된 후에야 SimPy가 시작된다.
-`tl.load()`는 TensorHandle(placeholder)을 반환하므로 실제 데이터에 접근할 수 없다.
-따라서 데이터 값에 따른 분기(dynamic control flow)가 불가능하다.
-
-본 ADR은 이 한계를 **메모리 연산에 한해** 해소한다 (D1, D3 참조).
-
 ### 제약 조건
 
 - SimPy는 single-thread 이벤트 루프 — numpy matmul을 안에서 하면 전체가 block
@@ -529,22 +514,3 @@ dtype별 tolerance 정책:
   (연산은 Phase 2에서 실행, Phase 1에서 결과 값 미확정).
   메모리 데이터 기반 분기는 greenlet으로 지원된다.
 - greenlet C 확장 의존성 추가 (pip install greenlet)
-
----
-
-## 영향받는 파일
-
-| 파일 | 변경 |
-|------|------|
-| `src/kernbench/components/base.py` | `_on_process_start/end` hook 추가 |
-| `src/kernbench/common/pe_commands.py` | `data_op = True` 추가, metadata 필드 확장 |
-| `src/kernbench/sim_engine/op_log.py` | 신규: OpRecord, OpLogger |
-| `src/kernbench/sim_engine/data_executor.py` | 신규: DataExecutor, MemoryStore |
-| `src/kernbench/sim_engine/engine.py` | op_logger 주입 (optional) |
-| `src/kernbench/triton_emu/tl_context.py` | `tl.load()` 등 내부에서 greenlet switch 호출 |
-| `src/kernbench/triton_emu/kernel_runner.py` | 신규: KernelRunner (greenlet ↔ SimPy 연결) |
-| `src/kernbench/components/builtin/pe_cpu.py` | Phase 0 제거, KernelRunner 호출로 변경 |
-| `pyproject.toml` | greenlet 의존성 추가 |
-
-컴포넌트 구현 파일 (pe_gemm.py, pe_dma.py, hbm_ctrl.py 등): **변경 없음**
-벤치마크 커널 (benches/*.py): **사용자 API 변경 없음**

docs/adr-ko/ADR-0022-prog-program-id-2d-grid.md

@@ -0,0 +1,100 @@
+# ADR-0022: 2D 그리드 program_id 시맨틱
+
+## Status
+
+Accepted
+
+## Context
+
+Triton 커널은 `tl.program_id(axis)`를 사용해 launch 그리드 내 자신의
+위치를 식별한다. 본 하드웨어는 2단계 계층을 갖는다: **큐브**가 **PE**를
+포함한다. 이전 구현은 `axis` 파라미터를 무시하고 항상 평탄화된 PE
+인덱스를 반환했기 때문에, 커널이 큐브 내부 위치와 큐브 식별자를 구분할
+수 없었다.
+
+## Decision
+
+`tl.program_id`와 `tl.num_programs`를 2D 하드웨어 그리드에 매핑한다:
+
+| Call | Returns | Description |
+|------|---------|-------------|
+| `tl.program_id(axis=0)` | `local_pe_id` | 큐브 내 PE 인덱스 |
+| `tl.program_id(axis=1)` | `cube_id` | 큐브 인덱스 |
+| `tl.num_programs(axis=0)` | `num_pes_per_cube` | 큐브당 PE 개수 |
+| `tl.num_programs(axis=1)` | `num_cubes` | 전체 큐브 개수 |
+
+전역 PID는 다음과 같이 도출된다:
+
+```python
+global_pid = tl.program_id(axis=1) * tl.num_programs(axis=0) + tl.program_id(axis=0)
+```
+
+### 축 매핑 근거
+
+- **axis=0 = PE (최내부)**: 큐브 내부 PE들은 HBM을 공유하고 로컬 NoC
+  메시를 통해 통신한다. 빠르고 강하게 결합된 차원이다 — 블록 내부의
+  스레드와 유사하다.
+- **axis=1 = 큐브 (외부)**: 큐브 간 통신은 더 높은 레이턴시의 UCIe를
+  통한다. 더 거친 스케줄링 차원이다 — 그리드 내의 블록과 유사하다.
+
+## Implementation
+
+### TLContext (`triton_emu/tl_context.py`)
+
+`cube_id`와 `num_cubes` 생성자 파라미터를 추가했다. `program_id()`와
+`num_programs()`가 `axis`로 디스패치한다:
+
+```python
+def program_id(self, axis: int = 0) -> int:
+    if axis == 1:
+        return self._cube_id
+    return self._pe_id
+
+def num_programs(self, axis: int = 0) -> int:
+    if axis == 1:
+        return self._num_cubes
+    return self._num_programs
+```
+
+### PE_CPU (`components/builtin/pe_cpu.py`)
+
+- `ctx.spec["system"]["sips"]["cubes_per_sip"]`에서 `num_cubes`를
+  추출한다.
+- `cube_id`(이미 `self._cube_idx`로 사용 가능)와 `num_cubes`를
+  TLContext에 전달한다.
+
+### KernelRunner (`triton_emu/kernel_runner.py`)
+
+- PE_CPU로부터 `num_cubes`를 수신한다.
+- greenlet 모드에서 `cube_id`와 `num_cubes`를 TLContext에 전달한다.
+
+## Backward Compatibility
+
+- `tl.program_id(0)` 또는 `tl.program_id()`를 사용하는 기존 코드는
+  변경되지 않는다 — 이전과 동일한 PE 인덱스를 반환한다.
+- `cube_id`와 `num_cubes`는 기본값이 `0`과 `1`이므로, 이를 제공하지
+  않는 호출자(예: 유닛 테스트)도 계속 동작한다.
+
+## Usage Example
+
+```python
+def sharded_gemm_kernel(a_ptr, b_ptr, out_ptr, M, K, N, tl):
+    local_pid = tl.program_id(axis=0)      # PE within cube
+    cube_id   = tl.program_id(axis=1)      # which cube
+    global_pid = cube_id * tl.num_programs(axis=0) + local_pid
+
+    # 전역 PID에 걸친 column-wise 샤딩
+    n_per_pid = N // (tl.num_programs(axis=1) * tl.num_programs(axis=0))
+    col_start = global_pid * n_per_pid
+
+    a = tl.load(a_ptr, shape=(M, K), dtype="f16")
+    b = tl.ref(b_ptr + col_start * K * 2, shape=(K, n_per_pid), dtype="f16")
+    h = tl.composite(op="gemm", a=a, b=b, out_ptr=out_ptr + col_start * M * 2)
+    tl.wait(h)
+```
+
+## Consequences
+
+- 벤치마크가 토폴로지 차원을 하드코딩하지 않고 큐브 인식 샤딩과 주소
+  지정을 표현할 수 있다.
+- 필요 시 axis=2(SIP 레벨)를 동일한 패턴을 따라 향후 추가할 수 있다.

docs/adr-ko/ADR-0023-dev-ipcq-pe-collective.md

@@ -2,7 +2,7 @@
 
 ## Status
 
-Proposed
+Accepted
 
 ## Context
 
@@ -17,14 +17,6 @@ Queue)를 통해** 일어난다.
 core-local 통신 큐와 유사하다. 호스트 레벨 collective(`dist.all_reduce`)는
 **미래 작업**으로 미루고, 본 ADR은 커널 collective 인프라에만 집중한다.
 
-### 현재 상태
-
-- ADR-0021 PE 파이프라인 리팩토링: PE 내부가 컴포넌트 단위로 분리됨
-  (PE_CPU, PE_SCHEDULER, PE_DMA, PE_FETCH_STORE, PE_GEMM, PE_MATH, PE_TCM, PE_MMU)
-- PE 간 직접 통신 채널 없음. 모든 데이터 이동은 PE_DMA → cube_noc/UCIe/PCIE → HBM 경로
-- 호스트 CCL skeleton (ADR 없음, ad-hoc 구현): `dist.init_process_group(backend="ahbm")`,
-  `_run_ccl_bench`가 rank별 greenlet로 동시 실행. collective는 stub 상태.
-
 ### 풀어야 할 문제
 
 1. PE 간 직접 데이터 이동 (peer's memory에 write)
@@ -365,23 +357,39 @@ data 경로의 piggyback 모델과 달리, credit return은 일반 vc_comm fabri
 거치지 않고 **별도 fast path**로 처리한다. 이는 실제 HW의 NVLink/UCIe
 credit return fast path를 추상화한 것이다.
 
-**Latency 계산**: magic constant가 아니라 **라우팅 경로의 bottleneck BW**
-기준으로 산출한다.
+**Latency 계산**: magic constant가 아니라 **라우팅 경로의 full path
+latency** (per-node overhead + edge propagation + drain) 기준으로
+산출한다.
 
 ```
 credit_size_bytes = 16  (ccl.yaml: ipcq_credit_size_bytes)
-path = router.find_path(self_pe, peer_pe)
-latency = compute_drain_ns(path, credit_size_bytes)
-        = credit_size_bytes / bottleneck_bw_on_path
+path = router.find_path(self_pe, peer_pe.pe_dma)
+latency = compute_path_latency_ns(path, credit_size_bytes)
+        = sum(edge.distance_mm * ns_per_mm)
+        + sum(node_overhead_ns[n] for n in path)
+        + credit_size_bytes / bottleneck_bw_on_path
 ```
 
+router는 source에만 `.pe_dma`를 자동 부여하므로 destination에는 반드시
+`.pe_dma` suffix를 명시해야 한다. 그렇지 않으면 `find_path`가 raise하고
+credit이 0 cost로 silently teleport되는 latent bug가 발생한다 (이번
+업데이트에서 수정됨).
+
+`tl.recv`는 credit-emit 완료를 yield-from으로 기다린다 (이전에는
+`env.process`로 fork). 이로써 credit-return cost가 receiver의
+`pe_exec_ns`에 반영되어, IPCQ control-plane이 consume-acknowledgement를
+완료한 뒤에야 recv가 kernel에 반환된다 — RAW DMA의 non-posted `tl.store`가
+HBM ack-trip을 기다리는 것의 protocol-level 등가물이다.
+
 이로써:
 - **토폴로지 비례 approximation**: cube 내 credit return과 cross-SIP credit이
-  자동으로 다른 latency를 가짐 (정확한 값은 아니지만 magic constant보다 의미 있음)
-- **Magic constant 없음**: 별도 `ipcq_ctrl_latency_ns` 같은 임의 값 불필요
-- **Deadlock 위험 없음**: piggyback과 달리 B가 A에게 보낼 데이터가 없어도
-  credit이 자동 발행됨
-- **기존 utility 재사용**: `ComponentContext.compute_drain_ns` 그대로 사용
+  자동으로 다른 latency를 가짐
+- **Magic constant 없음**: 모든 ns 값이 데이터 트래픽과 동일한 edge_map
+  및 `node_overhead_ns`에서 산출되는 `compute_path_latency_ns`로부터 옴
+- **Deadlock 위험 없음**: `peer_credit_store.put`은 unbounded, B가 A에게
+  보낼 데이터가 없어도 credit이 자동 발행됨
+- **`IPCQ ≥ raw DMA`** 보장: matched physical move에 대해 credit-emit이
+  RAW의 ack-trip cost와 균형을 이룸
 
 ```
 PE B: tl.recv(W) → 데이터 가져감 → my_tail++
@@ -426,11 +434,22 @@ backend init에서 IpcqInitMsg fan-out 시 양방향 fast path channel을 함께
 
 #### PE_DMA의 책임 추가
 
-PE_DMA(vc_comm)는 token 수신 시 다음 atomic 시퀀스로 처리한다.
-**두 동작 사이에 SimPy yield를 두어서는 안 된다** (I6 MUST 규칙 참조):
+PE_DMA(vc_comm)는 token 수신 시 다음 시퀀스로 처리한다: Transaction
+terminal의 BW drain을 먼저 지불하고, 이어서 atomic하게 data write +
+metadata forward 수행. **data write와 metadata forward 사이에는 SimPy
+yield를 두어서는 안 된다** (I6 MUST 규칙 참조). drain yield는 atomic
+구간 안이 아니라 그 앞에 위치해야 한다:
 
 ```python
-def _on_vc_comm_recv(self, env, token):
+def _on_vc_comm_recv(self, env, txn):
+    # Sender PE_DMA가 찍어 둔 drain_ns (= nbytes / bottleneck_bw) 를
+    # 여기서 지불. atomic 구간보다 앞이어야 한다 — recv는 bytes가
+    # "도착"한 이후에만 깨어나야 하므로.
+    drain = getattr(txn, "drain_ns", 0.0)
+    if drain > 0:
+        yield env.timeout(drain)
+
+    token = txn.request
     # ── ATOMIC: 두 동작 사이에 yield 금지 ──
     # 1. data를 dst_addr에 write (dst의 메모리 공간은 token.dst_endpoint.buffer_kind)
     data = self._memory_store.read(token.src_space, token.src_addr,
@@ -446,6 +465,32 @@ wire로 capacity가 unbounded인 store를 사용하므로 즉시 완료된다 (
 single-step). 이 최종 put이 atomic 구간의 끝이며, 그 이전에 다른 yield가
 삽입되면 안 된다.
 
+#### Drain-at-inbound semantics (D9 timing model)
+
+Transaction은 sender PE_DMA가 `drain_ns = nbytes / bottleneck_bw_on_path`
+를 찍어 둔 상태로 fabric에 들어간다. 이 simulator에서 per-hop `overhead_ns`
+는 각 forwarding component의 `run()` 에서 지불되고, 남은 BW drain은
+Transaction의 terminal node에서 한 번 지불된다. IPCQ가 아닌 모든
+Transaction (raw DMA, kernel-launch fanout 등) 은
+`ComponentBase._forward_txn` 이 terminal에서 이 drain을 지불한다. IPCQ의
+경우 목적지 PE_DMA가 `_handle_ipcq_inbound` 핸들러로 Transaction을
+가로채서 (IPCQ 전용 data write + metadata forward를 해야 하므로)
+**이 핸들러 최상단에서 drain을 명시적으로 지불해야 한다** — 그래야 IPCQ의
+timing model이 다른 모든 fabric Transaction과 동일선상에 놓인다.
+
+여기서 drain을 지불할 때의 side-effect:
+
+- **SRC `tl.send`**: 동작 불변. sender PE_DMA가 `sub_done` 을 `yield`
+  하지 않으므로 fire-and-forget 의미가 보존된다. metadata forward 이후
+  호출되는 `sub_done.succeed()` 는 sender 입장에서 listener가 없는 이벤트.
+- **DST `tl.recv`**: `drain_ns` 만큼 늦게 깨어난다. recv는 local PE_IPCQ
+  의 `IpcqMetaArrival` 수신 시에만 wake되며, metadata forward가 drain
+  이후로 이동했으므로 recv는 bandwidth까지 포함한 전체 fabric transfer
+  시간을 관측하게 된다.
+
+물리적 그림과 일치: send는 dispatch하고 바로 반환; recv는 bytes가 실제로
+자신의 inbox로 drain될 때까지 대기.
+
 #### Backpressure latency 정확도
 
 backpressure 해제까지 걸리는 시간:
@@ -924,7 +969,7 @@ tail 갱신은 D9 fast path SimPy Store 채널로 처리된다.
 
 ### D13. 테스트 전략
 
-ADR-0021의 D8 패턴을 따라 단위/통합/regression 테스트를 명시한다.
+단위/통합/regression 테스트를 명시한다.
 
 #### T1. 단위 테스트 (component-level)
 
@@ -1057,7 +1102,7 @@ F5. **Slot full + 무한 backpressure**:
 ### D15. 알고리즘 작성자 가이드 (요약)
 
 본 섹션은 알고리즘 작성자가 한 화면으로 시작점을 잡을 수 있도록 한다.
-자세한 step-by-step 가이드는 [docs/ccl-author-guide.md](../ccl-author-guide.md) 참조.
+자세한 step-by-step 가이드는 [docs/onboarding/ccl-author-guide.md](../onboarding/ccl-author-guide.md) 참조.
 
 #### 만지는 것 / 만지지 않는 것
 
@@ -1130,7 +1175,416 @@ def neighbors(rank, world_size, neighbor_map) -> dict | None:
 2. **send/recv 짝 맞지 않음** — peer 측 recv 없으면 hang (slot full backpressure)
 3. **dtype/shape 불일치** — 첫 구현은 검증 안 함, 작성자 책임
 
-자세한 step-by-step과 hello-world 예제는 `docs/ccl-author-guide.md` 참조.
+자세한 step-by-step과 hello-world 예제는 `docs/onboarding/ccl-author-guide.md` 참조.
+
+---
+
+## HW Realization Notes (Informative)
+
+**Status of this section**: Forward-looking. Describes how the simulator
+contract (D1–D15) would map to silicon. Not currently implemented;
+subject to revision before tapeout. The simulator implements the
+contract via Python/SimPy equivalents in
+[pe_ipcq.py](../../src/kernbench/components/builtin/pe_ipcq.py) and
+[pe_dma.py](../../src/kernbench/components/builtin/pe_dma.py).
+
+### D16. Proposed HW Block Diagram and End-to-End Dataflow
+
+![PE Baseline Architecture](../diagrams/pe_baseline.png)
+
+> Source: [`../diagrams/pe_baseline.d2`](../diagrams/pe_baseline.d2) — `d2 --layout=elk --scale 1.5`.
+
+![PE Proposed Architecture](../diagrams/pe_proposed.png)
+
+> Source: [`../diagrams/pe_proposed.d2`](../diagrams/pe_proposed.d2) — `d2 --layout=elk`.
+
+**Baseline → Proposed 핵심 변경**:
+
+- 단일 FIFO inbox → **compute port / IPCQ port 분리 + WRR Arbiter** (NEW)
+- PE_IPCQ (SimPy component) → **IPCQ Controller** (HW register + combinational logic)
+- TCM 내 **IPCQ Slot Region 예약 영역** 명시
+- Credit Injector / Receiver가 Fabric Port를 통해 NoC에 직접 연결
+
+#### End-to-End Sequence (HW view)
+
+```mermaid
+sequenceDiagram
+    participant CPU_A as PE_A: PE_CPU
+    participant IPCQ_A as PE_A: IPCQ Ctrl
+    participant DMA_A as PE_A: DMA
+    participant NOC as NoC Fabric
+    participant DMA_B as PE_B: DMA
+    participant IPCQ_B as PE_B: IPCQ Ctrl
+    participant TCM_B as PE_B: TCM
+    participant CPU_B as PE_B: PE_CPU
+
+    Note over CPU_A: tl.send(dir="E", src=0x1000)
+
+    CPU_A->>IPCQ_A: MMIO: send request
+    Note over IPCQ_A: Backpressure check:<br/>(head - peer_tail_cache) < n_slots → PASS<br/>Slot addr gen:<br/>dst = peer_rx_base + (head%n) × slot_size
+    IPCQ_A->>DMA_A: IpcqDmaToken {src, dst, sender_seq=head}
+    Note over IPCQ_A: my_head++
+    IPCQ_A-->>CPU_A: send returns (fire-and-forget)
+
+    Note over DMA_A: TCM read → snapshot in read buffer<br/>Flit pack: data + {sender_seq, dst_addr}
+    DMA_A->>NOC: IPCQ data flit(s)
+
+    Note over NOC: hop latency + BW drain
+
+    NOC->>DMA_B: IPCQ data flit(s)
+    Note over DMA_B: Terminal BW drain<br/>Slot write latency
+
+    rect rgb(255, 240, 220)
+        Note over DMA_B,IPCQ_B: ATOMIC (I6): same cycle, no stall
+        DMA_B->>TCM_B: write data → slot address
+        DMA_B->>IPCQ_B: Meta Extractor: {sender_seq, dst_addr}
+    end
+
+    Note over IPCQ_B: Range match dst_addr → direction "W"<br/>peer_head_cache["W"] = sender_seq + 1
+    IPCQ_B-->>CPU_B: recv_wake signal
+
+    Note over CPU_B: tl.recv(dir="W") wakes up
+    CPU_B->>IPCQ_B: recv request
+    Note over IPCQ_B: peer_head_cache > my_tail → YES<br/>slot_addr = rx_base + (tail%n) × slot_size
+    IPCQ_B-->>CPU_B: return slot_addr
+    CPU_B->>TCM_B: read data from slot
+    Note over IPCQ_B: my_tail++
+
+    IPCQ_B->>NOC: Credit (16B): {consumer_seq, dst_rx_base_pa}
+    Note over NOC: credit traversal (NoC latency)
+    NOC->>IPCQ_A: Credit arrival
+
+    Note over IPCQ_A: Match dst_rx_base_pa → direction "E"<br/>peer_tail_cache["E"] = consumer_seq<br/>Backpressure deassert (if stalled)
+```
+
+### D17. IPCQ Controller HW Module (신규)
+
+PE_CPU와 DMA Engine 사이에 위치하는 하드웨어 제어 블록. 시뮬레이터의
+`PeIpcqComponent`에 대응한다.
+
+#### QPair Register File
+
+방향별 queue pair 상태를 flip-flop으로 유지. PE_CPU가 MMIO(CSR)로 읽기/쓰기
+가능하며, init 시점에 소프트웨어가 채워넣는다.
+
+```
+Per-direction registers (each 64-bit):
+  my_head          — sender write position (monotonic)
+  my_tail          — receiver read position (monotonic)
+  peer_head_cache  — last known peer head (updated by Meta Extractor)
+  peer_tail_cache  — last known peer tail (updated by Credit Receiver)
+  rx_base_pa       — this PE's rx buffer base physical address
+  peer_rx_base_pa  — peer's rx buffer base physical address
+  n_slots          — ring depth (power-of-2 제약, D21 참조)
+  slot_size        — bytes per slot
+  peer_credit_tgt  — peer PE의 credit receive 주소
+
+Directions: 최대 8 (N/S/E/W/parent/child_left/child_right + spare)
+Total: 8 dirs × 9 regs × 8B = 576B flip-flops
+```
+
+#### Slot Address Generator (combinational)
+
+```
+Input:  pointer (my_head or my_tail), n_slots, slot_size, base_pa
+Output: slot_addr = base_pa + (pointer % n_slots) * slot_size
+
+Implementation:
+  n_slots power-of-2 → pointer & (n_slots - 1)   (AND mask, 1 gate)
+  slot_size power-of-2 → barrel shift             (1 cycle)
+  64-bit add → ripple/kogge-stone adder           (1 cycle)
+
+Latency: 1-2 cycles combinational
+```
+
+#### Backpressure Comparator (combinational)
+
+```
+full = (my_head - peer_tail_cache) >= n_slots
+
+Implementation: 64-bit subtract + unsigned compare
+Output: stall signal → PE_CPU (IPCQ send blocked) or DMA issue hold
+Latency: 1 cycle
+```
+
+#### Meta Extractor (inbound datapath sideband)
+
+DMA Engine의 inbound vc_comm path에 wired. 도착하는 IPCQ flit의 header에서
+metadata를 추출하여 queue pair 상태를 갱신한다.
+
+```
+Trigger: DMA inbound write completion (same cycle)
+Extract: {sender_seq, dst_addr} from flit header
+
+Direction matching (ADR-0025 D2):
+  for each dir:
+    match = (base_pa[dir] <= dst_addr) && (dst_addr < base_pa[dir] + n_slots[dir] * slot_size[dir])
+  8× parallel range comparators + priority encoder
+
+Update: peer_head_cache[matched_dir] = max(peer_head_cache, sender_seq + 1)
+Output: recv_wake signal → PE_CPU interrupt/flag
+Latency: 1 cycle (pipelined with DMA write — I6 atomicity 자연 보장)
+```
+
+#### Credit Injector (outbound)
+
+```
+Trigger: recv completion (my_tail 증가 후)
+Action:  pack 16B credit packet → DMA vc_comm (또는 dedicated credit VC)
+
+Packet: {consumer_seq = my_tail, dst_rx_base_pa = my_rx_base_pa}
+Latency: 1 cycle to generate, then NoC traversal
+```
+
+#### Credit Receiver (inbound sideband)
+
+```
+Trigger: 16B credit packet arrival (from NoC)
+Extract: {consumer_seq, dst_rx_base_pa}
+
+Direction matching (ADR-0025 D3):
+  for each dir:
+    match = (peer_rx_base_pa[dir] == credit.dst_rx_base_pa)
+
+Update: peer_tail_cache[matched_dir] = max(peer_tail_cache, consumer_seq)
+Output: send_wake signal → deassert backpressure stall
+Latency: 1 cycle
+```
+
+### D18. DMA Engine vc_comm IPCQ-aware Mode
+
+기존 vc_comm 채널(D8)에 IPCQ flit 처리 모드를 추가한다.
+
+**Outbound**:
+
+1. IPCQ Controller로부터 command 수신: `{src_addr, dst_addr, nbytes, sender_seq}`
+2. TCM에서 src_addr read → DMA read buffer에 snapshot (standard DMA behavior)
+3. Flit pack: data + piggyback metadata (sender_seq, dst_addr)
+4. NoC fabric port에 inject
+5. Fire-and-forget (completion 미대기)
+
+**Inbound**:
+
+1. NoC로부터 IPCQ flit 수신
+2. Terminal BW drain charge (`drain_ns = nbytes / bottleneck_bw`)
+3. Slot write latency charge (backing memory tier)
+4. **ATOMIC** (same pipeline stage, no stall insertion):
+   - TCM write: data → slot address
+   - Meta Extractor trigger: sender_seq + dst_addr → IPCQ Controller
+5. Done
+
+**I6 atomicity 하드웨어 보장**: TCM write completion과 Meta Extractor trigger가
+동일 pipeline stage에서 발생하므로 별도 synchronization이 불필요. 시뮬레이터의
+"no SimPy yield between MemoryStore.write and IpcqMetaArrival put" (D9, I6)이
+자연스럽게 보장된다.
+
+#### Data Snapshot Semantics
+
+DMA read buffer에 latch된 데이터는 src memory의 이후 수정에 영향받지 않는다.
+이는 DMA standard read-then-write behavior이므로 추가 HW 불필요.
+
+#### Credit Virtual Channel (선택적)
+
+- **옵션 A**: vc_comm에 credit을 multiplexing (16B header-only flit으로 구분).
+- **옵션 B**: 3rd dedicated credit VC 추가 (strict priority > data).
+
+옵션 B가 deadlock prevention에 유리하나, 16B credit의 BW 영향이 무시 가능하므로
+옵션 A로도 충분.
+
+### D19. Fabric Flit Format Extension
+
+```
+일반 data flit (예: 512-bit):
+┌──────────────────────────────────────────┐
+│ [511:480] routing header (32b)           │
+│ [479:0]   payload (480b = 60B)           │
+└──────────────────────────────────────────┘
+
+IPCQ data flit (첫 flit에만 metadata 포함):
+┌──────────────────────────────────────────┐
+│ [511:480] routing header (32b)           │
+│   [511]    ipcq_flag (1b)                │  ← IPCQ vs normal DMA 식별
+│   [510:509] vc_id (2b)                   │
+│   [508:480] route + hop count            │
+│ [479:416] ipcq_metadata (64b)            │  ← piggyback
+│   [479:448] sender_seq (32b)             │
+│   [447:416] dst_addr[31:0] (32b)         │  ← direction matching용
+│ [415:0]   payload (416b = 52B)           │
+└──────────────────────────────────────────┘
+후속 flits: full 60B payload (metadata 없음)
+
+Credit-only flit (128-bit, header-only):
+┌──────────────────────────────────────────┐
+│ [127:96]  routing header (32b)           │
+│   [127]   credit_flag (1b)               │
+│ [95:64]   consumer_seq (32b)             │
+│ [63:0]    dst_rx_base_pa (64b)           │
+└──────────────────────────────────────────┘
+```
+
+첫 flit의 payload가 60B → 52B로 감소 (13% overhead). Multi-flit transfer에서는
+후속 flit이 full payload이므로 대형 전송에서 overhead < 1%.
+
+### D20. TCM IPCQ Slot Region Layout
+
+```
+TCM Memory Map (16MB):
+┌─────────────────────────────┐ 0x000000
+│  Kernel Working Memory      │
+│  (compute tensors)          │
+│  ~14MB                      │
+├─────────────────────────────┤ 0xE00000
+│  IPCQ RX Buffers            │
+│  Dir N: slots × slot_size   │
+│  Dir S: slots × slot_size   │
+│  Dir E: slots × slot_size   │
+│  Dir W: slots × slot_size   │
+│  ~1MB                       │
+├─────────────────────────────┤ 0xF00000
+│  IPCQ Metadata / Scratch    │
+│  ~1MB                       │
+└─────────────────────────────┘ 0xFFFFFF
+```
+
+IPCQ region을 TCM의 상위 bank에 배치하여 compute access와의 bank conflict를
+최소화한다 (Risk D22 참조).
+
+### D21. 2nm Implementation Analysis
+
+#### Area Estimate
+
+| Module | Gate Count | Area (2nm est.) | Notes |
+|---|---|---|---|
+| QPair Register File | ~4.6K FF | 0.002 mm² | 576B flip-flops |
+| Slot Addr Gen + Backpressure | ~5K gates | 0.001 mm² | Combinational |
+| Meta Extractor + Credit Logic | ~3K gates | 0.001 mm² | 8× parallel comparators |
+| **IPCQ Controller subtotal** | **~12.6K** | **~0.004 mm²** | **PE 전체 대비 < 0.1%** |
+| DMA vc_comm 확장 | ~2K gates | 0.002 mm² | Flit pack/unpack |
+| **Total 변경분** | **~14.6K** | **~0.006 mm²** | |
+
+#### Timing
+
+| Path | Delay (2nm est.) | Target Clock | Margin |
+|---|---|---|---|
+| Backpressure (sub + cmp) | ~0.3 ns | 1 GHz (1 ns) | 3× |
+| Slot Addr Gen (mask + shift + add) | ~0.5 ns | 1 GHz | 2× |
+| Meta Extractor (8× range match) | ~0.4 ns | 1 GHz | 2.5× |
+| Credit Receiver (8× equality) | ~0.3 ns | 1 GHz | 3× |
+
+모든 critical path가 1 cycle 이내. Timing closure 문제 없음.
+
+#### Power
+
+- Active: ~1 mW (register R/W + comparators, send/recv 동작 시)
+- Idle: leakage only
+- PE 전체 전력 대비 무시 가능
+
+#### Constraints
+
+| 항목 | 제약 | 근거 |
+|---|---|---|
+| `n_slots` | **반드시 power-of-2** | mod → AND mask (1 gate). 임의 값은 divider 필요 (~10 cycles) |
+| `slot_size` | **power-of-2 권장** | mul → barrel shift. 임의 값은 multiplier 필요 |
+| TCM IPCQ region | **전용 bank 배치** | Compute access와 bank conflict 방지 |
+
+### D22. Risk Assessment
+
+#### TCM Bank Conflict
+
+- **Risk**: IPCQ slot write와 compute read가 동일 bank 접근 시 stall
+- **Mitigation**: IPCQ region을 TCM 상위 address의 전용 bank에 배치 (D20)
+- **Cost**: TCM banking flexibility 소폭 감소
+- **Severity**: Medium (성능 영향), Low (correctness 문제 아님)
+
+#### Credit Return Latency under Congestion
+
+- **Risk**: NoC 혼잡 시 credit return 지연 → sender backpressure stall
+- **Mitigation**:
+  - Credit을 별도 VC로 분리 + strict priority (16B로 BW impact 미미)
+  - 또는 n_slots를 넉넉히(8+) 설정하여 credit 지연을 buffer로 흡수
+- **Severity**: Low (credit 16B는 congestion에 거의 기여하지 않음)
+
+#### Inter-Direction Ordering
+
+- **Risk**: 같은 PE에서 여러 방향으로 동시 send 시 순서
+- **Mitigation**: Per-direction monotonic seq으로 충분. Inter-direction ordering은
+  kernel(소프트웨어) 책임 — 현재 시뮬레이터 모델과 동일 (D2 + D4)
+- **Severity**: Low (아키텍처 설계에 의해 해소)
+
+### D23. HW Alternatives Considered
+
+#### Doorbell + Polling (전통적 방식)
+
+```
+Send: DMA write data → DMA write doorbell register at peer → peer polls doorbell
+Recv: Polling loop on doorbell, or interrupt-driven
+```
+
+| 장점 | 단점 |
+|---|---|
+| 단순한 HW (IPCQ controller 불필요) | 2번의 DMA transaction (data + doorbell) |
+| 기존 DMA 재사용 | Data/doorbell 사이 ordering 보장 필요 (fence) |
+| | Polling은 전력 낭비, interrupt는 latency overhead |
+
+**평가**: Piggyback 대비 latency 2-3× 증가. **불채택.**
+
+#### Hardware Message Queue (NVIDIA NVLink 스타일)
+
+```
+Send: CPU → HMQ에 descriptor push → HW가 peer HMQ로 자동 전달
+Recv: HMQ에서 descriptor pop → data pointer 확인
+```
+
+| 장점 | 단점 |
+|---|---|
+| CPU는 descriptor만 작성 | 별도 HMQ engine 필요 (~0.05 mm²) |
+| Descriptor/data 분리 → 유연 | DMA와 별개 datapath → area/power 중복 |
+| | Large tensor에는 결국 DMA 필요 |
+
+**평가**: CCL의 large tensor 패턴에서 DMA 필수이므로 HMQ + DMA 이중 구조는
+면적 낭비. **불채택.**
+
+#### RDMA-style Completion Queue (CQ)
+
+```
+Send: DMA write → peer에 CQE 자동 생성
+Recv: CQ poll/interrupt → data 위치 확인
+```
+
+| 장점 | 단점 |
+|---|---|
+| InfiniBand/RoCE 성숙 모델 | CQ 관리 logic + CQE memory overhead |
+| Multi-tenant/isolation 용이 | CQE/data ordering 보장 추가 필요 |
+| | PE-to-PE CCL에는 over-engineered |
+
+**평가**: RDMA CQ는 host-facing NIC의 multi-tenant 격리에 적합.
+PE 간 단일 owner 환경에서는 불필요한 복잡성. **불채택.**
+
+#### Credit-in-Data Piggyback (v2 최적화 후보)
+
+현재 설계에서 credit return은 별도 16B packet이다. Bidirectional 통신
+패턴에서는 **reverse 방향 data flit에 credit을 합칠 수 있다.**
+
+```
+PE_A →E→ PE_B: data + sender_seq=3
+PE_B →W→ PE_A: data + sender_seq=5 + credit_ack=4  ← credit이 data에 합쳐짐
+```
+
+| 장점 | 단점 |
+|---|---|
+| Credit 전용 packet 제거 → NoC BW 절약 | Unidirectional 패턴에서는 fallback 필요 |
+| Bidirectional allreduce에서 credit latency → 0 | Flit header에 8B 추가 (overhead 미미) |
+| | Logic 복잡도 소폭 증가 |
+
+**평가**: 현재 설계의 우수한 최적화. Bidirectional allreduce에서 credit packet을
+완전 제거 가능. Standalone credit fallback도 유지. **v2로 채택 권고.**
+
+### Open HW Questions
+
+- IPCQ slot region size를 TCM의 몇 %까지 허용할 것인가? (현재 가정: ~1MB / 16MB = 6.25%)
+- Credit VC를 별도로 둘 것인가, vc_comm에 multiplexing할 것인가? (D18 참조)
+- Inter-SIP link에서의 flit format 호환성 검증 필요
+- n_slots 최대값 제한? (8 directions × 8 slots × 64KB = 4MB → TCM의 25%)
 
 ---
 
@@ -1192,29 +1646,3 @@ def neighbors(rank, world_size, neighbor_map) -> dict | None:
 - VC arbitration 모델이 first-order approximation이므로 heavy contention
   시나리오에서 실제 HW보다 약간 optimistic한 latency 결과 가능 (D8 한계)
 - VC chunk-level 인터리브로 PE_DMA 구현이 더 복잡해짐
-
----
-
-## 영향받는 파일
-
-| 파일 | 변경 |
-|------|------|
-| `topology.yaml` | pe_template에 pe_ipcq 추가, ipcq↔dma/cpu/tcm edge 추가 |
-| `components.yaml` | pe_ipcq_v1 등록 |
-| `src/kernbench/topology/builder.py` | PE 내부 edge에 ipcq 체인 추가 |
-| `src/kernbench/components/builtin/pe_ipcq.py` | 신규 |
-| `src/kernbench/components/builtin/pe_dma.py` | VC 추가, IpcqDmaToken 처리 |
-| `src/kernbench/common/pe_commands.py` | IpcqSendCmd, IpcqRecvCmd, IpcqDmaToken 정의 |
-| `src/kernbench/triton_emu/tl_context.py` | tl.send / tl.recv API |
-| `src/kernbench/runtime_api/distributed.py` | ccl.yaml 로드, init 시 IPCQ install (eager) |
-| `src/kernbench/runtime_api/kernel.py` | IpcqInitMsg (sideband) 정의 |
-| `src/kernbench/ccl/__init__.py` | 신규 — CCL 패키지 |
-| `src/kernbench/ccl/topologies.py` | 신규 — builtin topology generators (ring_1d, mesh_2d, tree_binary 등), `resolve_topology()` |
-| `src/kernbench/ccl/helpers.py` | 신규 — 알고리즘 작성 헬퍼 (chunked, ring_step 등) |
-| `src/kernbench/ccl/testing.py` | 신규 — mock CCL runtime (`run_kernel_in_mock`) |
-| `ccl.yaml` | 신규 — 알고리즘 metadata + IPCQ default 설정 |
-| `src/kernbench/ccl/algorithms/ring_allreduce.py` | 신규 — 첫 알고리즘 예제 |
-| `tests/test_pe_ipcq.py` | 신규 — PE_IPCQ 단위 테스트 |
-| `tests/test_pe_dma_vc.py` | 신규 — PE_DMA virtual channel 테스트 |
-| `tests/test_ipcq_e2e.py` | 신규 — send/recv end-to-end 테스트 |
-| `tests/test_ccl_topologies.py` | 신규 — builtin topology generator 단위 테스트 |

docs/adr-ko/ADR-0024-par-sip-tp-launcher.md

@@ -0,0 +1,236 @@
+# ADR-0024: SIP-level Launcher — rank = SIP
+
+## Status
+
+Accepted
+
+## Context
+
+### 목표
+
+`torch.distributed` collective 호출의 참여 단위(rank)를 **SIP**(device)
+경계에 맞춘다. 실제 PyTorch DDP/TP 스크립트와 **호스트 레벨에서 구분 없이**
+읽히는 bench 코드를 목표로 한다.
+
+real PyTorch와 비교:
+
+| 차원 | real PyTorch | KernBench |
+| --- | --- | --- |
+| 프로세스 모델 | N개 프로세스, 각 1 GPU | 1 프로세스, N greenlet, 각 1 SIP |
+| `get_rank()` | `RANK` env var | greenlet-local 레지스트리 |
+| `get_world_size()` | `WORLD_SIZE` env var | topology의 SIP 수 |
+| `torch.cuda.set_device(r)` (real) / `torch.ahbm.set_device(r)` (KernBench) | rank → GPU | rank → SIP |
+| `mp.spawn` | OS 프로세스 fork | greenlet fan-out |
+
+### 풀어야 할 문제
+
+1. **공개 API에서 rank = SIP** — bench worker가 PE 개념을 알지 않도록.
+2. **Greenlet-local rank/device tracking** — 1-프로세스 모델 안에서 각
+   worker greenlet이 자기 rank / 자기 SIP를 정확히 식별.
+3. **Tensor placement = structural (sip, cube, pe)** — rank가 SIP이면
+   기본 텐서 배치도 구조적 좌표로 표현되어야 함.
+
+### Non-problem (이 ADR 밖)
+
+- IPCQ direction addressing → ADR-0025
+- `DPPolicy.sip`/`num_sips` 제거 → ADR-0026
+- Megatron-style TP → ADR-0027
+- DTensor → ADR-0028 (future)
+- Worker scheduling / `mp.spawn` / collective drain / exception cleanup
+  → ADR-0027 D0/D1
+- Collective algorithm 구현 (intercube_allreduce, SFR config) → ADR-0032
+
+## Decision
+
+### D1. rank = SIP (world_size 해석)
+
+```python
+def _resolve_world_size(self) -> int:
+    if "world_size" in self._merged:
+        return int(self._merged["world_size"])
+    defaults = self._cfg_all.get("defaults", {})
+    if "world_size" in defaults:
+        return int(defaults["world_size"])
+    spec = self.ctx.spec or {}
+    return int(spec.get("system", {}).get("sips", {}).get("count", 1))
+```
+
+우선순위: 알고리즘 override > defaults override > SIP count. `ccl.yaml`
+override는 legacy "rank = PE" 테스트 경로로 유지.
+
+### D2. Greenlet-local rank registry (+ debug warning)
+
+```python
+class DistributedContext:
+    def __init__(self):
+        self._backend = None
+        self._rank_by_greenlet: dict = {}
+
+    def _bind_rank(self, g, rank: int) -> None:
+        self._rank_by_greenlet[g] = int(rank)
+
+    def get_rank(self) -> int:
+        self._ensure_initialized()
+        from greenlet import getcurrent
+        g = getcurrent()
+        if g not in self._rank_by_greenlet:
+            if os.environ.get("KERNBENCH_DEBUG"):
+                warnings.warn(
+                    "get_rank() called outside a bound greenlet — returning 0. "
+                    "Likely a bug unless running single-driver."
+                )
+            return 0
+        return int(self._rank_by_greenlet[g])
+```
+
+### D3. `torch.ahbm.set_device(rank)` — SIP 바인딩
+
+KernBench 백엔드 이름은 `ahbm` (ADR-0023). Real PyTorch는
+`torch.cuda.set_device(r)`이지만 우리는 CUDA가 아니므로 honestly-named
+namespace를 사용한다.
+
+```python
+class _AhbmNamespace:
+    """torch.ahbm — per-greenlet SIP device binding.
+
+    Real-PyTorch parity idiom: ``torch.cuda.set_device(rank)``. Since
+    KernBench's backend is 'ahbm' (not CUDA), we expose the equivalent
+    API under ``torch.ahbm`` to avoid pretending to be a CUDA runtime.
+    """
+
+    def __init__(self):
+        self._device_by_greenlet: dict = {}
+
+    def set_device(self, device: int) -> None:
+        from greenlet import getcurrent
+        self._device_by_greenlet[getcurrent()] = int(device)
+
+    def current_device(self) -> int | None:
+        from greenlet import getcurrent
+        return self._device_by_greenlet.get(getcurrent())
+
+# Attached to RuntimeContext as `self.ahbm = _AhbmNamespace()`.
+# Bench code: `torch.ahbm.set_device(rank)` mirrors `torch.cuda.set_device`.
+```
+
+**PyTorch 2.x style 병행 지원**: 최신 PyTorch는 device-agnostic한
+`torch.accelerator` 네임스페이스를 지향 (`torch.accelerator.set_device_index(r)`,
+`torch.accelerator.current_device_index()`). Device vendor에 종속되지 않는
+코드를 쓰려는 사용자를 위해 KernBench도 이 표면을 병행 지원한다.
+
+```python
+class _AcceleratorNamespace:
+    """torch.accelerator — device-agnostic API (PyTorch 2.x style).
+
+    Aliases torch.ahbm for bench code that prefers device-neutral idiom:
+        torch.accelerator.set_device_index(rank)
+        torch.accelerator.current_device_index()
+    """
+
+    def __init__(self, ahbm: _AhbmNamespace):
+        self._ahbm = ahbm
+
+    def set_device_index(self, device: int) -> None:
+        self._ahbm.set_device(device)
+
+    def current_device_index(self) -> int | None:
+        return self._ahbm.current_device()
+
+# RuntimeContext
+self.ahbm = _AhbmNamespace()
+self.accelerator = _AcceleratorNamespace(self.ahbm)   # alias
+```
+
+Bench 작성자는 다음 중 하나를 선택 — 둘 다 내부적으로 같은 레지스트리를 보유:
+
+```python
+torch.ahbm.set_device(rank)                   # KernBench-native, explicit backend
+torch.accelerator.set_device_index(rank)      # PyTorch 2.x device-agnostic
+```
+
+### D4. Tensor placement = structural (sip, cube, pe) 좌표
+
+`resolve_dp_policy`가 `target_sip`을 직접 받아 구조적 좌표로 placement 생성.
+세부는 ADR-0026.
+
+```python
+# RuntimeContext._create_tensor
+current_sip = self.ahbm.current_device()          # (D3 naming)
+if current_sip is None:
+    current_sip = 0  # single-driver fallback (D2와 일관)
+placement = resolve_dp_policy(
+    dp, shape=shape_2d, itemsize=itemsize,
+    num_pe=eff_num_pe, num_cubes=eff_num_cubes,
+    target_sip=current_sip,
+)
+```
+
+Post-hoc `pe_index` shifting 없음 — ShardSpec이 `(sip, cube, pe)` 구조적
+좌표를 직접 보유. ShardSpec 상세는 ADR-0026.
+
+### D5. SIP 그리드 크기 — 명시적 `sips.w/h` 해석
+
+2D inter-SIP topology (`torus_2d`, `mesh_2d_no_wrap`)의 SIP 그리드 형태
+(width × height)는 `system.sips.w` / `system.sips.h`에서 해석한다. D1이
+`sips.count`로 `world_size`를 해석하는 것과 같은 방식이다. 우선순위:
+명시적 `w/h` (`w*h == count` 검증) > 정사각 fallback
+(`w/h` 미지정 시에만 `round(sqrt(count))²`) > error.
+
+```python
+sips = spec.get("system", {}).get("sips", {})
+if sip_topo == "ring_1d":
+    w, h = 0, 0                          # 1D sentinel (no grid)
+elif sips.get("w") is not None and sips.get("h") is not None:
+    w, h = int(sips["w"]), int(sips["h"])
+    if w * h != n_sips:
+        raise ValueError(f"sip layout {w}x{h} != sips.count ({n_sips})")
+else:
+    side = int(round(math.sqrt(n_sips)))
+    if side * side != n_sips:
+        raise ValueError("non-square sips.count requires explicit sips.w/h")
+    w, h = side, side
+```
+
+이로써 2D SIP 그리드가 완전 정사각이어야 한다는 기존 가정을 제거한다:
+6-SIP `torus_2d` / `mesh_2d_no_wrap`은 이제 `w: 3, h: 2`(또는 `2x3`)로
+표현 가능하다. 도출된 `(w, h)`는 알고리즘의 inter-SIP exchange로 전달된다
+(ADR-0032 D5에서 소비). 이전 코드 경로는 ring이 아닌 모든 topology에서
+`round(sqrt(count))²`를 조용히 취해 잘못된 그리드(예: 6 SIP에 2×2)를
+만들었다. fail-loud fallback을 갖춘 명시적 `w/h` 경로가 이를 대체한다.
+
+---
+
+## Dependencies
+
+- **ADR-0023** (IPCQ): backend `ahbm` namespace의 기원.
+- **ADR-0026** (DPPolicy intra-device): D4의 `resolve_dp_policy` 시그니처와
+  ShardSpec의 구조적 좌표 표현.
+- **ADR-0027** (Megatron TP + scheduler): worker scheduling, `mp.spawn`,
+  collective drain, exception cleanup의 구현 기준.
+
+---
+
+## Non-goals
+
+- **IPCQ protocol 수정**: ADR-0023 유지.
+- **DPPolicy 필드 정리**: ADR-0026.
+- **Megatron-style TP**: ADR-0027.
+- **Worker scheduling / spawn / drain / exception cleanup**: ADR-0027 D0/D1.
+- **Collective algorithm 구현**: ADR-0032.
+- **Multi-node (프로세스 간)**: 단일 프로세스.
+
+---
+
+## Consequences
+
+### Positive
+
+- **Bench = real PyTorch DDP** (공개 API 관점).
+- **Greenlet-local rank**: 1-프로세스 모델에서 cross-rank correctness 가능.
+- **Structural placement 좌표**: ADR-0026 / ADR-0027 / ADR-0032의 다른 ADR이
+  `(sip, cube, pe)` 3튜플 위에서 일관되게 동작.
+
+### Neutral
+
+- IPCQ PE-level protocol (ADR-0023) 불변.
+- IO_CPU 역할 불변 (기존 transit 그대로).

docs/adr-ko/ADR-0025-algo-ipcq-direction-addressing.md

@@ -2,7 +2,7 @@
 
 ## Status
 
-Proposed (Revision 2 — Address-based matching; peer_direction field dropped)
+Accepted (Revision 2 — Address-based matching; peer_direction field dropped)
 
 ## Context
 
@@ -13,34 +13,6 @@ topology / dict-order에 의존하지 않고 **주소 기반**으로 일관되
 2-rank bidirectional ring (또는 여러 direction이 동일 peer를 가리키는
 topology 일반)에서 정확히 동작하도록 한다.
 
-### 현재 상태 (ADR-0023 D9 구현)
-
-`src/kernbench/components/builtin/pe_ipcq.py` — `_handle_meta_arrival`:
-
-```python
-def _handle_meta_arrival(self, msg: IpcqMetaArrival) -> None:
-    token = msg.token
-    sender_key = (token.src_sip, token.src_cube, token.src_pe)
-    for d, qp in self._queue_pairs.items():
-        p = qp["peer"]
-        if (p.sip, p.cube, p.pe) == sender_key:
-            qp["peer_head_cache"] = max(qp["peer_head_cache"], token.sender_seq + 1)
-            # ... wake recv waiters ...
-            return
-```
-
-`_credit_worker`도 동일한 "sender-coord-first-match" 패턴.
-
-`src/kernbench/ccl/install.py` — `reverse_direction`:
-
-```python
-def reverse_direction(my_rank: int, peer_rank: int) -> str | None:
-    for d, target in neighbor_table[peer_rank].items():
-        if target == my_rank:
-            return d
-    return None
-```
-
 ### 드러난 버그 — 2-rank bidirectional ring
 
 `ring_1d(rank, world_size=2)` → `{"E": 1, "W": 1}` (rank 0). 양쪽 방향이 같은 peer.
@@ -89,7 +61,14 @@ direction_idx × bytes_per_direction). 따라서:
 `src/kernbench/ccl/install.py`:
 
 ```python
-_OPPOSITE_DIR = {"E": "W", "W": "E", "N": "S", "S": "N"}
+# Extended in ADR-0032 with global_* pairs for inter-SIP directions,
+# which were introduced by configure_sfr_intercube_multisip to keep
+# intercube (N/S/E/W) and inter-SIP (global_N/S/E/W) namespaces disjoint.
+_OPPOSITE_DIR = {
+    "E": "W", "W": "E", "N": "S", "S": "N",
+    "global_E": "global_W", "global_W": "global_E",
+    "global_N": "global_S", "global_S": "global_N",
+}
 
 def reverse_direction(my_rank: int, peer_rank: int, my_dir: str) -> str | None:
     """Find peer's direction that reciprocates my_dir→peer_rank.
@@ -282,51 +261,6 @@ for plan in plans:
 
 ---
 
-## Test strategy
-
-### T1. Unit — `reverse_direction` opposite-preference
-
-`tests/test_ccl_install.py` (확장):
-- Ring ws=2: `reverse_direction(0, 1, "E")` → "W", `reverse_direction(0, 1, "W")` → "E"
-- Ring ws=4: `reverse_direction(0, 1, "E")` → "W" (자연스러운 opposite)
-- Mesh 2×2: `reverse_direction(r, peer, "N")` → "S", "E" ↔ "W"
-- Tree binary: opposite 없는 direction (parent) → fallback 경로
-- Non-symmetric topology: opposite가 peer에 없고 다른 direction만 있는 경우
-
-### T2. Runtime — `_handle_meta_arrival` dst_addr 매칭
-
-`tests/test_pe_ipcq.py` (확장):
-- 2-rank pair install 후, E direction dst_addr로 meta arrival → E의 `peer_head_cache`
-  증가 (W는 불변)
-- W direction dst_addr로 meta arrival → W의 `peer_head_cache` 증가
-- 잘못된 dst_addr (어느 rx range에도 속하지 않음) → 에러 또는 silent drop
-  (결정 후 명시)
-
-### T3. Credit — `dst_rx_base_pa` 매칭
-
-`tests/test_pe_ipcq.py` (확장):
-- E direction send 후 peer가 consume → credit에 자기 W의 `my_rx_base_pa`
-  담아 송신 → sender의 E direction `peer_tail_cache` 증가
-- W direction도 동일
-
-### T4. E2E — 2-rank bidirectional ring
-
-`tests/test_ipcq_e2e.py`:
-- 2-rank ring_1d로 tl.send(E) + tl.recv(W) pattern이 양방향으로 작동
-- ADR-0024의 `test_ccl_allreduce_matrix.py`에서 ring at ws=2가 통과
-
-### T5. Install invariant — rx_base range disjointness
-
-`tests/test_ccl_install_plan.py` (확장):
-- I3.1 검증: `build_install_plans` 결과에서 모든 qp의 rx_base range가 disjoint
-
-### T6. 회귀
-
-- 기존 ws≥3 ring / mesh / tree 테스트 그대로 통과
-- `test_pe_ipcq`, `test_ipcq_e2e` 기존 케이스 회귀
-
----
-
 ## Consequences
 
 ### Positive
@@ -347,19 +281,3 @@ for plan in plans:
 
 - IPCQ protocol의 semantic layer (sender가 dst_addr 계산, receiver가 수신)는
   불변.
-
----
-
-## Affected files
-
-| File | Change |
-|------|--------|
-| `src/kernbench/ccl/install.py` | D1: `reverse_direction`에 `my_dir` 인자 추가, opposite-preference |
-| `src/kernbench/components/builtin/pe_ipcq.py` | D2: `_handle_meta_arrival` dst_addr 매칭 / D3: `_credit_worker` dst_rx_base_pa 매칭 / `_delayed_credit_send`가 `dst_rx_base_pa` 필드 채움 |
-| `src/kernbench/common/ipcq_types.py` | D3: `IpcqCreditMetadata`에 `dst_rx_base_pa` 필드 추가 |
-| `src/kernbench/ccl/install_plan.py` (ADR-0024 신규) | D6: I3.1 invariant 검증 (optional) |
-| `docs/adr/ADR-0023-ipcq-pe-collective.md` | Reference note: runtime 매칭 방식이 ADR-0025에서 바뀜 |
-| `tests/test_ccl_install.py` | T1 |
-| `tests/test_pe_ipcq.py` | T2, T3 |
-| `tests/test_ipcq_e2e.py` | T4 |
-| `tests/test_ccl_install_plan.py` | T5 |

docs/adr-ko/ADR-0026-par-dppolicy-intra-device.md

@@ -13,53 +13,6 @@ intra-device 추상화로 명확화한다. SIP 간 분산(TP)은 별도 레이
 (ADR-0024의 `torch.ahbm.set_device(rank)` 또는 ADR-0027의 Megatron parallel
 layers가 담당).
 
-### 현재 상태
-
-`src/kernbench/policy/placement/dp.py`:
-
-```python
-@dataclass(frozen=True)
-class DPPolicy:
-    sip: Literal["replicate", "column_wise", "row_wise"] = "replicate"
-    cube: Literal["replicate", "column_wise", "row_wise"] = "replicate"
-    pe: Literal["replicate", "column_wise", "row_wise"] = "replicate"
-    num_pes: int | None = None
-    num_cubes: int | None = None
-    num_sips: int | None = None    # ← 제거 대상
-```
-
-`sip` / `num_sips` 필드는 텐서를 SIP 경계 **너머**로 분산하는 경로를 제공함.
-이는:
-
-- **ADR-0024의 launcher 모델과 충돌**: ADR-0024는 "rank = SIP = 1 worker per SIP"
-  모델. 각 worker가 자기 SIP에 텐서를 생성. 텐서가 여러 SIP에 걸치는 경우는
-  Megatron-style TP가 개별 primitive로 처리해야 함.
-- **사용자 의도와 불일치**: "DPPolicy는 한 디바이스 내에서 PE들로 분산하는 방법"
-  (사용자 진술).
-- **개념 혼동**: `DPPolicy.sip="column_wise"`는 실제로 **TP**. 이름이 DP인데
-  하는 일은 TP → 신규 사용자에게 혼란.
-
-### 영향받는 call site (rollback 시점 grep 결과)
-
-**생성 사이트** (`DPPolicy(sip=...` 또는 `num_sips=...`):
-- `tests/test_runtime_api_tensor.py`
-- `benches/ccl_allreduce.py` (ADR-0024 scope 내에서 이미 개편됨)
-- `tests/test_va_offset.py`
-- `benches/va_offset_verify.py`
-- `tests/test_sip_parallel.py`
-
-**참조 사이트** (`dp.sip`, `policy.sip`, `num_sips` 등):
-- `src/kernbench/runtime_api/context.py` (`_create_tensor`, `launch`)
-- `src/kernbench/components/builtin/pe_cpu.py`
-- `src/kernbench/components/legacy/builtin/pe_cpu.py`
-- `src/kernbench/policy/placement/dp.py` (구현 자체)
-- `tests/test_tensor.py`, `test_ipcq_types.py`
-
-**핵심 테스트**: `test_sip_parallel.py`는 이름 그대로 "SIP 병렬성을 DPPolicy로
-표현하는" 테스트. 이 ADR 이후 **새 launcher 모델로 재작성** 필요.
-
----
-
 ## Decision
 
 ### D1. `DPPolicy`에서 `sip` + `num_sips` 필드 제거
@@ -70,7 +23,7 @@ class DPPolicy:
     """Intra-device (cube × PE) data-parallel policy.
 
     SIP-level placement is controlled by ``torch.ahbm.set_device(rank)``
-    (ADR-0024 D10) and, for model-level TP, by Megatron-style parallel
+    (ADR-0024 D3) and, for model-level TP, by Megatron-style parallel
     layers (ADR-0027). DPPolicy does not cross SIP boundaries.
     """
     cube: Literal["replicate", "column_wise", "row_wise"] = "replicate"
@@ -84,7 +37,7 @@ class DPPolicy:
 ### D2. `ShardSpec` — structural (sip, cube, pe) 좌표, `pe_index` 완전 제거
 
 현재 `ShardSpec.pe_index`는 **global flat index** (`sip × cubes × pes + cube ×
-pes + pe`). 이는 ADR-0024 D11이 "abstraction leakage"로 지적한 형태.
+pes + pe`). 이는 ADR-0024 D4이 "abstraction leakage"로 지적한 형태.
 
 본 ADR에서 ShardSpec을 **structural 좌표로 재정의**하고, `pe_index`는
 property로도 **남기지 않는다**:
@@ -120,7 +73,7 @@ class ShardSpec:
 
 ### D3. `resolve_dp_policy`가 `target_sip`을 받아 structural 좌표 생성
 
-ADR-0024 D11의 계약 구현. Post-hoc shifting 없음.
+ADR-0024 D4의 계약 구현. Post-hoc shifting 없음.
 
 ```python
 # src/kernbench/policy/placement/dp.py (after)
@@ -182,14 +135,14 @@ def resolve_dp_policy(
 
 ### D4. `_create_tensor` — 구조적 좌표로 직접 placement
 
-ADR-0024 D11 연속선. Post-hoc shifting 제거, 구조적 좌표를 `resolve_dp_policy`
+ADR-0024 D4 연속선. Post-hoc shifting 제거, 구조적 좌표를 `resolve_dp_policy`
 호출 시점에 직접 지정.
 
 ```python
 # context.py _create_tensor (after)
 current_sip = self.ahbm.current_device()
 if current_sip is None:
-    # Single-driver fallback (ADR-0024 D9와 일관).
+    # Single-driver fallback (ADR-0024 D2와 일관).
     # Launcher 기반 코드가 set_device()를 빼먹으면 조용히 SIP 0에 박히는
     # 문제가 있음 → debug mode에서 경고.
     if os.environ.get("KERNBENCH_DEBUG"):
@@ -258,66 +211,6 @@ for sip_id in sip_range:
 권고. `PEIdentity` 값객체는 명시적 타입 장점은 있지만 boilerplate가 크고 현재
 allocator dict의 유일한 key라 오버엔지니어링. Tuple 유지.
 
-### D6. Migration — 기존 call site
-
-**(A) `DPPolicy(sip=..., num_sips=..., ...)` 사용하던 코드**:
-
-- `DPPolicy(sip="column_wise", cube=..., pe=...)` 패턴 → **해당 bench를 ADR-0024
-  launcher로 재작성**. worker가 `set_device(rank)`로 SIP 선택, DPPolicy는
-  cube/PE만.
-- `DPPolicy(sip="replicate", num_sips=1, ...)` 패턴 → `DPPolicy(cube=..., pe=...)`로
-  축소 (필드가 사라지니 자연스럽게).
-
-**(B) `dp.sip`, `dp.num_sips` 읽던 코드**:
-
-- 제거. `launch()`의 `_compute_local_shape`에서 `dp.sip` 분기 삭제.
-- `pe_cpu.py`가 `dp.sip`을 참조하던 곳도 정리.
-
-**(C) `ShardSpec.pe_index`를 사용하던 코드 — 전부 수정 필요**:
-
-- `.pe_index` 접근은 이제 `AttributeError` 발생 → 모든 call site 수정 필수.
-- Allocator lookup: `allocators[spec.pe_index]` →
-  `allocators[(spec.sip, spec.cube, spec.pe)]`
-- Flat integer가 꼭 필요한 국소 문맥: `spec.sip * N_CUBES * N_PE + spec.cube *
-  N_PE + spec.pe` 명시적 계산. **국소 변수로만 사용하고 공개 API에 노출하지
-  않는다**.
-
-**구현 착수 전 grep audit 체크리스트**:
-
-1. **Property 참조**:
-   - `\.pe_index\b` — 필드/property 접근 모두 (regex)
-   - `pe_index=` — 생성 시점의 키워드 인자
-   - `pe_index:` — dataclass 필드 선언
-2. **Allocator / dict indexing**:
-   - `allocators\[` — dict lookup 패턴. `allocators[spec.pe_index]` 같은
-     것이 걸리는지
-   - `_allocators\[` — 같은 패턴 (prefix _)
-3. **Flat index 수동 계산 블록**:
-   - `flat_idx =`
-   - `pe_index =` (좌변)
-   - `* pes_per_cube +` (전형적 flat 계산 패턴)
-   - `* self._num_cubes \* self._pes_per_cube` (global flat 계산)
-4. **Serialization / logging**:
-   - `asdict(.*shard` — dataclass 직렬화 시 `pe_index` 자동 포함 여부
-   - `repr(.*ShardSpec` — 로그 포맷에서 의존하는지
-   - JSON/YAML 저장 포맷에서 `pe_index` 키 사용 여부
-5. **Tests asserting integer PE identity**:
-   - `assert .*pe_index` — 정수 동일성 주장
-   - `spec.pe_index ==` — 비교 (SIP-local 의미로 변하면 테스트가 깨질 수 있음)
-
-각 match마다 "이 호출자가 global flat / SIP-local / 내부 lookup 중 무엇을
-기대했나"를 판단한 뒤 구조적 좌표로 교체.
-
-**(D) `test_sip_parallel.py`**:
-
-- 이름 유지, 내용은 ADR-0024의 multi-greenlet launcher 기반 재작성.
-- "SIP 병렬성 = rank 별 worker × 각자 DPPolicy" 로 검증.
-
-**(E) `test_va_offset.py`, `benches/va_offset_verify.py`**:
-
-- `num_sips=1`만 쓰는 경우가 대부분. 단순히 필드 제거.
-- SIP offset 테스트가 핵심이면 `set_device(rank)` + 구조적 좌표 관찰로 이식.
-
 ### D7. 하위 호환 — 불가 (cleanup ADR)
 
 이 ADR은 **breaking change**.
@@ -331,17 +224,6 @@ KernBench는 사내 프로젝트로 call site가 한정되어 있어 한 번에
 **Silent drift 차단**이 property 완전 제거의 주된 이점: global flat을 기대한
 코드가 SIP-local 결과를 받아 조용히 잘못된 인덱싱을 할 가능성 제거.
 
-### D8. 문서 업데이트
-
-- `ADR-0008` (tensor deploy) — DPPolicy 의미 갱신 note, ShardSpec 구조적 좌표
-  전환 명시
-- DPPolicy docstring에 "intra-device only" 명시 (D1 코드 스니펫의 docstring)
-- ShardSpec docstring에 **structural coordinates `(sip, cube, pe)`를 직접
-  사용하며, `pe_index`는 더 이상 제공되지 않음**을 명시 (D2)
-- `docs/ccl-author-guide` 등 튜토리얼에서 `sip=...` 예시 제거
-
----
-
 ## Dependencies
 
 - **ADR-0024** (launcher): `set_device(rank)` 및 current-device scoping이
@@ -378,56 +260,6 @@ KernBench는 사내 프로젝트로 call site가 한정되어 있어 한 번에
 
 ---
 
-## Test strategy
-
-### T1. 단위 테스트 갱신
-
-- `tests/test_tensor.py`, `tests/test_ipcq_types.py`, `tests/test_runtime_api_tensor.py`
-  — DPPolicy 생성자 인자 정리, ShardSpec 구조적 좌표 검증
-- `tests/test_va_offset.py` — `num_sips=1` 제거 후 동작 유지
-
-### T2. `resolve_dp_policy` 구조적 좌표 반환
-
-`tests/test_dp_policy.py` (new 또는 확장):
-- `resolve_dp_policy(dp, ..., target_sip=1)` 결과의 모든 ShardSpec이 `sip=1`
-- 각 spec의 `(cube, pe)`가 local (0..num_cubes-1, 0..num_pe-1)
-- 같은 topology에서 `target_sip=0`과 `target_sip=1` 결과가 sip 필드만 다름
-
-### T3. `test_sip_parallel.py` 재작성
-
-SIP 병렬성 검증을 launcher 기반으로:
-
-```python
-def test_sip_parallel_via_launcher(topology):
-    ...
-    def worker(rank, ws, torch):
-        torch.ahbm.set_device(rank)
-        t = torch.zeros((1, 128), dtype="f16",
-                         dp=DPPolicy(cube="column_wise", pe="column_wise"))
-        # verify shard.sip == rank (structural coord)
-
-    spawn(worker, nprocs=n_sips, ...)
-```
-
-### T4. Allocator key migration
-
-`tests/test_allocator_structural_key.py` (new 또는 기존 확장):
-- `PEMemAllocator` dict이 `(sip, cube, pe)` tuple key로 작동
-- `deploy_tensor`가 구조적 좌표로 allocator lookup
-- `_free_tensor`도 동일
-
-### T5. E2E 회귀
-
-ADR-0024의 `test_ccl_allreduce_matrix.py` 그대로 통과.
-
-### T6. 오류 검증
-
-- `DPPolicy(sip="column_wise")` 호출 → `TypeError`. 테스트로 명시.
-- `DPPolicy(num_sips=2)` 호출 → `TypeError`.
-- `spec.pe_index` 접근 → `AttributeError` (property 완전 제거 검증).
-
----
-
 ## Consequences
 
 ### Positive
@@ -435,7 +267,7 @@ ADR-0024의 `test_ccl_allreduce_matrix.py` 그대로 통과.
 - **개념 분리 명확**: DPPolicy = intra-device, TP = inter-device.
 - **API 단순화**: DPPolicy 생성자 필드 ~33% 축소.
 - **Structural 좌표 일관성**: ShardSpec이 `(sip, cube, pe)` 튜플로 표현 →
-  abstraction leakage 해소 (ADR-0024 D11 계약 충족).
+  abstraction leakage 해소 (ADR-0024 D4 계약 충족).
 - **`pe_index` 의미 명확**: SIP-local이 단일 해석. Global flat이 필요하면 명시.
 - **Launcher 모델 일관성**: ADR-0024의 "1 worker per SIP" 모델이 유일한 SIP
   경계 제어 메커니즘.
@@ -454,23 +286,3 @@ ADR-0024의 `test_ccl_allreduce_matrix.py` 그대로 통과.
 ### Neutral
 
 - 기존 `cube` / `pe` 필드 의미 불변.
-
----
-
-## Affected files
-
-| File | Change |
-|------|--------|
-| `src/kernbench/policy/placement/dp.py` | D1: `sip`/`num_sips` 제거 / D2: `ShardSpec`에 `sip`/`cube`/`pe` structural fields 추가, **`pe_index` property 제거** / D3: `resolve_dp_policy`에 `target_sip`, SIP-level 루프 제거 / 내부 resolver가 반환하는 shard 타입 이름도 `local_pe`로 명확화 (이름 충돌 방지) |
-| `src/kernbench/runtime_api/context.py` | D4: `_create_tensor` `target_sip` 전달 / D5: `_ensure_allocators` dict key → `(sip, cube, pe)` tuple / `launch`의 `dp.sip` 분기 제거 |
-| `src/kernbench/runtime_api/tensor.py` | D5: `deploy_tensor`가 구조적 좌표로 allocator lookup |
-| `src/kernbench/components/builtin/pe_cpu.py` | D6: `dp.sip` 참조 제거 |
-| `src/kernbench/components/legacy/builtin/pe_cpu.py` | D6: 동일 |
-| `benches/ccl_allreduce.py` | ADR-0024 scope에서 이미 처리 |
-| `benches/va_offset_verify.py` | D6: `num_sips=1` 제거 |
-| `tests/test_runtime_api_tensor.py` | D6 |
-| `tests/test_va_offset.py` | D6 |
-| `tests/test_tensor.py`, `test_ipcq_types.py` | D6 |
-| `tests/test_sip_parallel.py` | T3: launcher 기반 재작성 |
-| `tests/test_dp_policy.py` (new 또는 확장) | T2 |
-| `tests/test_allocator_structural_key.py` (new) | T4 |

docs/adr-ko/ADR-0027-par-megatron-tp.md

@@ -2,9 +2,7 @@
 
 ## Status
 
-Proposed (Revision 7 — resume invariant / main-context wait 비재귀 invariant /
-global barrier over-serialization tradeoff / TP forward yield-safety 명시,
-2026-04-14)
+Accepted
 
 ## Context
 
@@ -19,20 +17,6 @@ Megatron-style을 선택한 이유:
 - NVIDIA Megatron / DeepSpeed가 확립한 인더스트리 표준.
 - DTensor는 선언적이라 디자인 공간이 더 크다 → 단계적.
 
-### 현재 상태
-
-- KernBench는 TP가 없음. 기존 `DPPolicy.sip="column_wise"` 경로는 ADR-0026에서
-  제거됨. rank = SIP launcher (ADR-0024) 위에 TP primitive를 얹는다.
-- ADR-0024 Phase B에서 **worker-greenlet env.run 재진입 버그**가 드러남:
-  worker가 `ctx.wait(h)` (tensor 생성 시 MmuMapMsg 등)를 호출하면 `env.run`이
-  worker 컨텍스트에서 돌고, 이때 spawn되는 kernel greenlet의 `_parent`가
-  worker가 되어 orphan 발생. `ring_default_ws` strict xfail의 근본 원인.
-- `dist.all_reduce`는 이미 `_defer_wait=True` + worker yield 패턴으로 이 문제를
-  피함 ([distributed.py:119-134](src/kernbench/runtime_api/distributed.py#L119-L134)).
-- TP layer의 forward는 매번 `torch.launch("gemm", ...)`를 호출하고, 그 뒤에
-  `dist.all_reduce`가 따라오는 패턴이 반복됨. worker-wait 문제를 **반드시**
-  해결하지 않으면 TP 샘플이 첫 실행에서 실패.
-
 ### TP primitive 스펙 (Megatron-LM 참조)
 
 - **ColumnParallelLinear**: weight의 **column(out_features)** 축을 TP ranks에
@@ -180,9 +164,9 @@ while alive:
    - 구현이 이를 **감지**할 필요는 없다 (타임아웃/steps-since-yield 카운터
      등). 이는 user contract이며 위반 시 증상은 "simulation hang"이다.
    - **Future extension**: non-collective 긴 계산 경로가 자주 나오면
-     ADR-0024 D13의 `torch.distributed.cooperative_yield()` primitive (명시적
-     no-op yield)를 도입할 수 있다. 현 ADR 범위 밖. Breaking change 아님 —
-     필요 시 추가하면 됨.
+     명시적 `torch.distributed.cooperative_yield()` primitive (no-op yield)를
+     도입할 수 있다. 현 ADR 범위 밖. Breaking change 아님 — 필요 시 추가하면
+     됨.
    - Round 내에서는 alive worker 전체가 한 번씩 `switch`를 받는다. 단일 round
      안에서 한 worker가 여러 번 wait를 호출해도 그 turn 안에서 순차적으로
      enqueue된 뒤 scheduler drain 한 번에 일괄 처리 (FIFO).
@@ -197,7 +181,7 @@ while alive:
    - **두 큐는 서로 다른 dependency source**: worker wait은 worker가 직접
      `submit + wait` 쌍으로 만들어낸 handle (tensor deploy, MmuMap 등). collective
      큐는 `dist.all_reduce`가 내부적으로 enqueue한 kernel launch handle이며
-     worker는 이걸 직접 wait하지 않는다 (ADR-0024 D7).
+     worker는 이걸 직접 wait하지 않는다 (D0.5의 두 큐 drain 모델 참조).
    - **Correctness 관점 독립**: collective는 worker 관점에선 "이미 submit된
      후 yield한" 상태. 그 완료 타이밍은 worker의 다음 action 시점 이전이기만
      하면 됨. worker wait 큐와의 순서 dependency 없음.
@@ -220,7 +204,7 @@ while alive:
      index로 두거나 append 전 `h not in pending_set` 검사) 가능. correctness
      를 바꾸지 않는 최적화로 분류.
 
-4. **Exception propagation + sibling cleanup (ADR-0024 D13 방식 채택)**.
+4. **Exception propagation + sibling cleanup**.
    worker greenlet이 raise하면 `g.switch()`가 main으로 예외를 전달한다.
    scheduler loop은 즉시 중단되고 다음 cleanup을 **명시적으로** 수행:
 
@@ -595,7 +579,7 @@ TP layer의 weight/output 표현에서 두 개념을 명확히 분리한다:
 
 | 개념 | 결정 주체 | 범위 |
 |---|---|---|
-| **TP shard ownership** (어느 rank가 weight의 어떤 slice를 소유하는가) | greenlet-local rank + `torch.ahbm.set_device(rank)` (ADR-0024 D9/D10) | **cross-rank, cross-SIP** |
+| **TP shard ownership** (어느 rank가 weight의 어떤 slice를 소유하는가) | greenlet-local rank + `torch.ahbm.set_device(rank)` (ADR-0024 D2/D3) | **cross-rank, cross-SIP** |
 | **Intra-rank placement** (소유된 slice를 rank 내부에서 cube × PE로 어떻게 분산하는가) | `DPPolicy(cube=..., pe=...)` (ADR-0026) | **한 rank 내부 (SIP 경계 안)** |
 
 따라서 `ColumnParallelLinear`가 `(in_features, out_features // ws)` shape로
@@ -839,40 +823,11 @@ strict-xfail 케이스를 본 ADR 구현 이후 **PASS**로 전환하는 것을
 
 ## Dependencies
 
-- **ADR-0024** (launcher): rank = SIP, greenlet-local rank, `dist.all_reduce`,
-  `torch.ahbm.set_device(rank)`. 본 ADR의 D0/D1이 이 인프라를 확장.
+- **ADR-0024** (launcher): rank = SIP, greenlet-local rank,
+  `torch.ahbm.set_device(rank)`.
 - **ADR-0026** (DPPolicy intra-device): weight tensor의 per-rank slice 표현.
 - **ADR-0023 / ADR-0025** (IPCQ): `dist.all_reduce` 구현의 기반.
 
-### Supersedes (partial)
-
-ADR-0024의 다음 섹션은 **미구현 상태의 설계**이며, 본 ADR이 더 단순한 모델로
-대체한다:
-
-- **ADR-0024 D7 (`_CollectiveBarrier.submit_and_drain`)** — epoch 기반 last-
-  arriver-drains 패턴. 문제: last arriver가 **worker 컨텍스트에서** `ctx.wait`을
-  호출해 env.run을 drive → D0.2가 막으려는 orphan 원인을 재현한다. 본 ADR의
-  **D0.4 two-queue drain** (worker가 모두 yield한 뒤 main이 drain)이 동일한
-  "모든 rank가 submit 완료 전까지 어떤 rank의 collective도 진행되지 않음"
-  invariant를 **worker-safe하게** 제공한다. `_CollectiveBarrier` 클래스는
-  구현하지 않는다.
-- **ADR-0024 D12/D13 (`spawn_workers` skeleton)** — signature / scheduler
-  loop / exception handling 설계. 본 ADR의 **D1**이 real-PyTorch API와 일치하는
-  signature (`spawn(fn, args, nprocs)`)로 재정의하며, D0 scheduler drain을 단일
-  위치에서 수행한다. ADR-0024 D13의 exception cleanup (siblings
-  `throw(SystemExit)` + `SpawnException` 래핑)은 본 ADR에 그대로 흡수
-  (D0.4-(4) 참조).
-
-현 구현은 ADR-0024의 D7/D12/D13 어느 것도 landing하지 않았으므로 supersede에
-따른 마이그레이션 비용은 없음. 향후 `docs/adr/ADR-0024`에 "superseded by
-ADR-0027 D0/D1" 주석만 추가하면 정합.
-
-**Source of truth (normative, 구현자 대상)**: worker scheduling / collective
-drain / spawn / exception cleanup의 구현 기준은 **ADR-0027 D0/D1이다**. 구현
-시 ADR-0024 D7/D12/D13의 pseudocode / contract / signature를 참고하지 말 것 —
-두 ADR이 다른 결론을 낼 때는 항상 ADR-0027이 우선한다. 리뷰어도 이 원칙으로
-PR을 심사.
-
 ---
 
 ## Non-goals
@@ -907,155 +862,6 @@ PR을 심사.
 
 ---
 
-## Test strategy
-
-### T1. Unit — `tests/test_tp_parallel_state.py` (신규)
-
-- `initialize_model_parallel(ws)`가 world_size와 일치하는 경우만 통과.
-- `get_tensor_model_parallel_rank()`가 greenlet-local rank 반환 (ADR-0024 D9
-  회귀).
-- 미초기화 상태에서 `get_tensor_model_parallel_world_size()`가 적절히 실패.
-
-### T2. Unit — `tests/test_tp_layers.py` (신규)
-
-**Shape / structural checks**:
-
-- `ColumnParallelLinear(in=256, out=512).weight.shape` per-rank가 `(256, 512/ws)`.
-- `RowParallelLinear(in=512, out=256).weight.shape` per-rank가 `(512/ws, 256)`.
-- `ColumnParallelLinear.forward(x)`의 출력 텐서 shape이 `(M, K/ws)`.
-
-**Numerical correctness (weight ≠ zero)**: 단순 shape assert는 대수적 오류를
-놓치므로, 결정론적 non-zero 입력/weight으로 실제 연산 결과 검증:
-
-- **T2.a (ColumnParallel, deterministic)**: weight를 per-rank identity
-  (또는 `(i, j) → i + rank * k_local + j` 같은 결정론적 패턴)으로 초기화
-  (`tensor.copy_`). 입력 `x`를 상수 벡터로 둔 뒤 forward. 각 rank의 출력이
-  **기대치 `x @ W_rank_local`와 rtol/atol 1e-2 이내로 일치** (gemm kernel의
-  fp16 round-off 고려).
-- **T2.b (RowParallel, reduced output equality — primary)**: 모든 rank의
-  forward 결과가 동일 전역 행렬 곱 `concat([x_0..x_{ws-1}]) @ concat([W_0..
-  W_{ws-1}])`과 일치하는지 검증. rank-별 `y.numpy()` 비교로 (i) all-reduce 후
-  elementwise equality와 (ii) 기대치(host-side numpy로 계산) 일치 **둘 다**
-  assert. observable-only 검증 — internal hook 불필요.
-
-  *Optional implementation note*: partial-sum 단계를 더 세밀히 관찰하고 싶으면
-  `_pending_collective_handles` enqueue 직전 intercept hook을 쓸 수 있으나,
-  이는 내부 구현 detail에 결합되므로 ADR 수준의 test contract는 T2.b의
-  observable equality만 요구한다.
-- **T2.c (rank-identity after all_reduce)**: 모든 rank의 `y.numpy()`이 elementwise
-  identical (mean뿐 아니라 full array equality, rtol 1e-2).
-
-**기존 weak assertion 금지**: `output mean이 identical` 같은 aggregate-only
-검증은 silently 깨지기 쉽기에 **main assertion으로 쓰지 말 것** — 보조
-sanity로만 사용.
-
-### T3. Worker-wait 일반화 + orphan regression — `tests/test_worker_wait_drain.py` (신규)
-
-본 테스트의 핵심 목적은 queue 동작이 아니라 **ADR-0024 Phase B orphan
-regression의 직접 방지**이다. 다음을 assert:
-
-- **T3.a**: Worker가 `ctx.wait(h)`을 호출하면 `_pending_worker_waits`에
-  handle이 enqueue되고 main이 drain하기 전까지 worker는 resume되지 않는다.
-- **T3.b**: `_drain_pending` 직후 worker가 resume되고 handle은 `_completed`
-  상태.
-- **T3.c**: Multi-worker에서 모든 worker가 같은 drain 지점에서 resume.
-- **T3.d (orphan invariant, 핵심)**: Worker 함수가 `torch.launch(...)`를
-  호출한 뒤, SimPy engine이 실제로 돌기 시작하는 시점에 **kernel greenlet의
-  `_parent`는 main greenlet**이다. 테스트는 `kernel_runner.run`을 monkey-patch
-  하거나 `KernelRunner._parent` capture 시점에 assertion hook을 걸어 이
-  invariant를 직접 검증.
-- **T3.e (symptom regression)**: D0 없이는 T3.d와 등가인 GreenletExit 실패가
-  재현되어야 함 (historical failure mode 문서화 — 실제 테스트는 D0 도입 후
-  skip 또는 xfail 처리).
-- **T3.f (idempotency)**: 같은 handle을 `ctx.wait(h)`로 두 번 호출해도
-  `engine.wait`은 한 번만 불린다 (D0.4-(3)).
-- **T3.g (exception propagation)**: Worker가 `wait` 호출 후 raise하면 main
-  scheduler loop이 즉시 중단되고 예외가 위로 전파. 남은 `_pending_worker_waits`는
-  drain되지 않는다 (D0.4-(4)).
-
-### T4. `torch.multiprocessing.spawn` — `tests/test_mp_spawn.py` (신규)
-
-- `spawn(fn, args, nprocs)`이 nprocs 개의 greenlet을 생성하고 각각 rank로 bind.
-- 모든 worker 완료 후 return.
-- 기존 bench `ccl_allreduce.py`의 hand-rolled loop을 `mp.spawn`으로 교체해도
-  matrix 회귀 통과.
-
-### T5. Host-read barrier — `tests/test_host_read_barrier.py` (신규)
-
-D0.5 contract를 직접 검증:
-
-- **T5.a**: Worker가 `launch → tensor.numpy()`를 연속 호출하면 barrier가 동작,
-  numpy 결과는 kernel 완료 후 값 (post-drain).
-- **T5.b**: `launch → tensor.shape` (metadata)는 barrier 발동 안 함 (pending
-  queue 그대로 유지).
-- **T5.c**: Pending 큐가 비어 있는 상태의 `numpy()` 호출은 yield 없이 즉시
-  read (불필요한 context switch 방지).
-- **T5.d**: `__getitem__`, `data` 역시 T5.a와 동일한 barrier 발동.
-- **T5.e**: Collective pending (all_reduce) 진행 중 상태에서 `numpy()` 호출 시
-  collective drain까지 기다린 뒤 read.
-- **T5.f (copy_ write barrier)**: target tensor에 미완료 pending handle이
-  있는 상태에서 `target.copy_(source)` 호출 시, write 전에 drain 발동.
-  주입한 host source가 drain-이후 상태에 덮어써지는지 확인 (stale-overwrite
-  없음).
-- **T5.g (closed-set via registry)**: barrier entry-point의 closed-set은
-  **명시적 registry** (예: `tensor.py` 상단의 `_HOST_READ_BARRIERS = frozenset
-  ({"numpy", "data", "__getitem__", "__repr__", "copy_"})`)로 유지한다.
-  테스트는:
-  1. registry에 나열된 각 entry-point에 **실제 barrier 주입이 되어 있는지**
-     (invocation 시 pending queue를 확인하고 yield 경로를 거치는지) 관찰.
-  2. 새 host-read semantic API 추가는 code review에서 registry 업데이트를
-     의무화 (CODEOWNERS / review checklist로 운영).
-
-  **Non-goal**: Python introspection (method 시그니처, docstring 분석 등)으로
-  barrier-부재 API를 자동 탐지하는 것은 정밀도 문제로 ADR scope 밖. registry
-  + review 접근으로 충분.
-
-### T6. E2E — `tests/test_tp_mlp.py` (신규)
-
-2-layer MLP (ColumnParallel → RowParallel) forward:
-
-**Structural / liveness**:
-
-- `ws = SIP count` (topology.yaml 기준 current 2) 모델로 실행 완료.
-- **Deadlock 없음**: scheduler loop이 유한 시간 내 종료 (pytest-timeout 등).
-- **Completion trace**: 각 `launch` 및 `all_reduce`가 `ctx._traces`에 entry
-  남김 (count = 예상 layer 수).
-
-**Numerical correctness (필수)**:
-
-- **T6.a (zero-weight sanity)**: weight 전부 0 → 출력 전부 0. 파이프라인이
-  돌긴 하는지 확인용 smoke test. **이것만으로는 불충분 — T6.b/T6.c와 함께
-  채택**.
-- **T6.b (deterministic pattern)**: 모든 weight를 결정론적 non-zero pattern
-  (예: all 0.01, 또는 per-rank identity에서 파생된 값)으로 `copy_`. 입력도
-  상수. 기대 출력을 host-side numpy로 계산한 뒤 각 rank의 `y.numpy()`와 rtol
-  1e-2로 비교.
-- **T6.c (rank-consistency post all-reduce)**: RowParallel의 all-reduce
-  이후 **모든 rank의 output이 elementwise identical** (T2.c와 동일 기준).
-  단순 mean 일치가 아니라 full array equality.
-- **T6.d (shape contract)**: ColumnParallel 출력이 `(B, D_hidden / ws)`,
-  RowParallel 출력이 `(B, D_out)`.
-
-### T7. 회귀 — `ring_default_ws` xfail 해제
-
-- `tests/test_ccl_allreduce_matrix.py::test_ccl_allreduce_matrix[ring_default_ws]`의
-  `@pytest.mark.xfail(strict=True)` 제거 → **PASS**여야 함.
-- Acceptance criteria (observable):
-  - **Deadlock 없음**: bench가 유한 시간 내 종료.
-  - **GreenletExit 없음**: stderr/log에 GreenletExit trace 없음.
-  - **Rank 0 산출**: `ring_allreduce_tcm (ws=2): 2 OK` 문자열이 출력.
-  - **Completion trace**: `all_reduce` trace entry 존재.
-  - **Numerical**: 각 rank의 입력 `r+1`에 대한 sum(1..ws)=3 결과를 tolerance
-    1e-1 이내로 달성.
-
-### T8. 회귀 — 기존 전체 test suite
-
-- ADR-0026까지 통과하던 모든 test가 그대로 통과 (523 passed + 1 xfail).
-- Phase 2 완료 기준: 524 passed (xfail 해제 포함) + 0 xfail + 위 T1~T7 신규
-  테스트 전부 통과.
-
----
-
 ## Consequences
 
 ### Positive
@@ -1080,29 +886,3 @@ D0.5 contract를 직접 검증:
 
 - ADR-0024/0026 기반 위에 순수한 상위 레이어 추가. Hardware simulation
   stack에 영향 없음 (D0 제외).
-
----
-
-## Affected files
-
-| File | Change |
-|------|--------|
-| `src/kernbench/runtime_api/context.py` | D0.1/D0.2: `_pending_worker_waits` + `ctx.wait`의 worker fork, D1.3: `self.multiprocessing` namespace attach |
-| `src/kernbench/runtime_api/multiprocessing.py` | 신규 (D1): `_MultiprocessingNamespace.spawn` + `_drain_pending` + `SpawnException` |
-| `src/kernbench/runtime_api/distributed.py` | `_pending_collective_handles` 타입 annotation 보강 (`list[tuple[RequestHandle, int, dict]]`); spawn exception cleanup에서 clear 호출 지점 노출 |
-| `src/kernbench/runtime_api/tensor.py` | D0.5 barrier 주입: `numpy`, `__getitem__`, `data`, `__repr__`, `copy_` (source read + target write) |
-| `src/kernbench/tp/__init__.py` | 신규: public API re-export |
-| `src/kernbench/tp/parallel_state.py` | 신규: D3 |
-| `src/kernbench/tp/layers.py` | 신규: D4/D5 |
-| `src/kernbench/tp/primitives.py` | 신규: D6 |
-| `src/kernbench/tp/kernels.py` | 신규: TP layer용 `_gemm_kernel` (bench 복제) |
-| `src/kernbench/tp/mappings.py` | 신규 stub (backward TODO) |
-| `benches/tp_mlp.py` | 신규 샘플 (D7) |
-| `benches/ccl_allreduce.py` | hand-rolled loop → `torch.multiprocessing.spawn`으로 교체 (D1.4) |
-| `tests/test_tp_parallel_state.py` | 신규 (T1) |
-| `tests/test_tp_layers.py` | 신규 (T2) |
-| `tests/test_worker_wait_drain.py` | 신규 (T3): orphan invariant 직접 검증 포함 |
-| `tests/test_mp_spawn.py` | 신규 (T4) |
-| `tests/test_host_read_barrier.py` | 신규 (T5): D0.5 host-read barrier contract |
-| `tests/test_tp_mlp.py` | 신규 (T6) |
-| `tests/test_ccl_allreduce_matrix.py` | `ring_default_ws` xfail 제거 (T7) |

docs/adr-ko/ADR-0032-algo-intercube-allreduce.md

@@ -0,0 +1,279 @@
+# ADR-0032: 큐브 간 All-Reduce — pe0 큐브-메시 리듀스 + 다중-SIP 교환
+
+## Status
+
+Accepted (supersedes ADR-0029).
+
+## Context
+
+### 목표
+
+토폴로지 계층을 활용하는 단일 all-reduce 알고리즘을 정의한다: 각 SIP
+내부의 큐브 메시(큐브 간) + SIP 간 교환. 단일 커널, 단일 SFR 구성
+경로이며 `topology.yaml`과 `ccl.yaml`로 구동된다.
+
+### ADR-0029(계층적 3-레벨)를 대체하는 이유
+
+ADR-0029는 시스템의 모든 PE가 참여하는 3-레벨(큐브 내 → 큐브 간 →
+SIP 간) 알고리즘을 제안했다. 실제로는 텐서가 큐브 내 PE 단위가 아니라
+**큐브 단위로 샤딩되는** 일반적 워크로드 패턴과 맞지 않으면서, 큐브 내
+PE-PE stage 복잡성(양방향 reduce + 체인 브로드캐스트)을 추가한다.
+
+또한 계층적 설계는 다음을 요구했다:
+- PE별 이웃 그래프 설치 (`_build_pe_installs` 다중 레벨)
+- 다중 레벨 토폴로지 스키마 (`hierarchical_3level`)
+- `all_pes` 매퍼 + `multi_pe_sip_local` 검증자 인프라
+
+아래의 큐브 간 알고리즘은 이 모든 것을 제거한다: **4×4 큐브 메시 위에서
+pe0만의 same-lane 큐브 간 reduce**, 그 다음 루트 큐브에서 SIP 간 교환,
+그 다음 다시 브로드캐스트. 더 단순한 커널, 더 단순한 와이어링,
+일반적인 큐브당 DP 워크로드에 대해 동일한 대역폭 특성을 갖는다.
+
+### 현재 상태
+
+- `src/kernbench/ccl/algorithms/lrab_hierarchical_allreduce.py` — 커널
+- `src/kernbench/ccl/sfr_config.py` — `configure_sfr_intercube_multisip`
+- `src/kernbench/runtime_api/distributed.py` — `AhbmCCLBackend`가
+  `init_process_group` 시점에 자동으로 와이어링한다.
+- 기존 `ring_allreduce`, `mesh_allreduce`, `tree_allreduce`,
+  `hierarchical_allreduce` 모듈과 그 테스트는 **제거됨**.
+
+---
+
+## Decision
+
+### D1. 알고리즘 구조 — 5단계 (center-root, 양방향)
+
+루트 큐브는 큐브 메시의 기하학적 **중심**에 위치한다:
+
+```
+root_col  = cube_w // 2
+root_row  = cube_h // 2
+root_cube = root_row * cube_w + root_col   # 중심; 4×4 메시에서 10
+```
+
+각 reduce/broadcast 단계는 이 중심을 향해 **양방향으로** 수렴/발산하여,
+corner-root 워크 대비 SIP 내부 임계 경로를 절반으로 줄인다 (4×4 메시:
+reduce 4홉 + broadcast 4홉 vs SE-코너 루트의 6+6).
+
+각 SIP에 대해 (`mp.spawn`으로 동시에 launch):
+
+```
+Phase 1 — col == root_col에서 수렴하는 Row reduce (큐브 메시, pe0만):
+    좌측 절반(col < root_col)은 W→E로, 우측 절반(col > root_col)은
+    E→W로 진행; root_col 큐브가 양쪽을 병합 → row sum 보유.
+
+Phase 2 — col == root_col에서 row == root_row로 수렴하는 Col reduce:
+    위쪽(row < root_row)은 N→S로, 아래쪽(row > root_row)은 S→N로 진행;
+    루트 큐브가 양쪽을 병합 → 전체 SIP sum 보유.
+
+Phase 3 — cube_id == root_cube에서 SIP 간 교환 (pe0만):
+    Ring / torus-2d row+col ring / mesh-2d chain reduce+broadcast —
+    sip_topo_kind(topology.yaml의 sips.topology)로 선택.
+
+Phase 4 — col == root_col에서 root_row로부터 바깥쪽으로 Col 브로드캐스트.
+
+Phase 5 — root_col로부터 바깥쪽으로 큐브 메시 전반에 Row 브로드캐스트.
+```
+
+모든 단계가 끝나면 모든 큐브의 pe0이 전역 sum을 보유한다.
+
+**단일 큐브 fast-path**: `cube_w == cube_h == 1`(rank당 큐브 하나, 일반적인
+TP 케이스)인 경우 SIP 내부 reduce/broadcast 단계를 건너뛰고 곧바로
+Phase 3 SIP 간 교환으로 진행한다.
+
+커널은 `sip_topo_kind ∈ {0, 1, 2}`(ring_1d, torus_2d, mesh_2d_no_wrap)로
+파라미터화된 단일 함수이다. Phase 1-2와 4-5는 토폴로지 전반에서 동일하며,
+phase 3만 분기한다. 헬퍼 함수 `_inter_sip_ring`, `_inter_sip_torus_2d`,
+`_inter_sip_mesh_2d`가 세 가지 교환 패턴을 인코딩한다.
+
+### D2. 텐서 레이아웃 (rank = SIP, 워커별)
+
+ADR-0024에 따라 프로세스 그룹 레벨에서 rank = SIP이다. 각 워커가
+자신의 큐브-메시 전체 텐서를 할당한다:
+
+```python
+dp = DPPolicy(cube="row_wise", pe="replicate", num_cubes=16, num_pes=1)
+tensor = torch.zeros((n_cubes, n_elem), dtype="f16", dp=dp)
+```
+
+샤드 레이아웃: SIP당 16개 샤드, 큐브별 pe0에 하나씩. 커널은 각 큐브의
+샤드를 `pe_addr = t_ptr + cube_id * n_elem * 2`로 주소 지정한다.
+
+### D3. SFR / IPCQ 와이어링 — `configure_sfr_intercube_multisip`
+
+ADR-0024의 rank-to-2-PE 설치를 대체한다. 어느 큐브가 루트인지 또는 어느
+SIP 토폴로지가 선택되었는지와 무관하게 **모든 SIP의 모든 큐브의 pe0**에
+대해 PE_IPCQ 이웃 테이블을 와이어링한다. 이를 통해 커널이 런타임에 루트
+큐브를 선출할 수 있고, 재와이어링 없이 토폴로지 전환을 지원한다.
+
+| Level | Direction labels | Scope |
+|---|---|---|
+| SIP 내부 큐브 간 | N / S / E / W | 모든 큐브의 pe0 → 메시 이웃의 pe0 (랩어라운드 없음) |
+| SIP 간 (모든 큐브) | global_E / global_W / global_N / global_S | sip A의 큐브 c의 pe0 → `sips.topology`에 따른 피어 SIP의 큐브 c의 pe0 |
+
+SIP 간 방향은 `global_*` 접두사를 사용하여 큐브 간 방향과 네임스페이스를
+분리한다. ADR-0025의 `_OPPOSITE_DIR`은 `global_E ↔ global_W` 및
+`global_N ↔ global_S`로 확장되어, 2-SIP 양방향 ring에 대한 역방향
+리졸버가 올바르게 처리되도록 한다.
+
+내부적으로 이 함수는 다음 인자로 `install_ipcq`를 호출한다:
+- `world_size = n_sips × n_cubes`
+- `rank_to_pe = [(sip, cube, 0) for sip in range(n_sips) for cube in range(n_cubes)]`
+- 위 매핑을 생성하는 클로저로 캡처된 `neighbors()` 함수.
+
+이 `world_size`는 IPCQ 와이어링 내부적이며 프로세스-그룹 rank로 유출되지
+않는다.
+
+### D4. SIP 토폴로지 — `topology.yaml`에서
+
+```yaml
+system:
+  sips:
+    count: 2
+    topology: ring_1d       # or torus_2d, mesh_2d_no_wrap
+```
+
+- `ring_1d`: n_sips-1 라운드의 `send global_E / recv global_W`.
+- `torus_2d`: `w × h` 랩핑 메시. `global_E/W`에서 row ring, 이어서
+  `global_S/N`에서 col ring.
+- `mesh_2d_no_wrap`: 랩어라운드 없는 `w × h` 메시. 차원별 chain
+  reduce + 브로드캐스트.
+
+2D 그리드 크기 `(w, h)`는 `system.sips.w/h`에서 온다 (ADR-0024 D5).
+정사각 fallback (`round(sqrt(n_sips))²`)은 `w/h`가 생략된 경우에만
+적용되므로, 직사각형 그리드(예: 6 SIP을 `3×2`로)는 명시적 `w/h`로
+지원된다.
+
+### D5. 프로세스-그룹 통합 — `AhbmCCLBackend`
+
+`init_process_group` 시점에 백엔드는:
+
+1. `ccl.yaml` + `topology.yaml`을 로드한다.
+2. `system.sips.topology`로부터 알고리즘 모듈의 `TOPO_NAME_TO_KIND`를
+   통해 `sip_topo_kind`를 도출하고, `sip_topo_w, sip_topo_h`는
+   `system.sips.w/h`에서 정사각 fallback과 함께 도출한다 (ADR-0024 D5).
+3. `configure_sfr_intercube_multisip(engine, spec, cfg)`를 호출한다 —
+   일회성 SFR 와이어링, NCCL 커뮤니케이터 생성을 모방한다.
+
+각 `dist.all_reduce(tensor)` 호출 시:
+
+1. `cfg["module"]`로부터 `kernel_fn`을 해석한다.
+2. `kernel_args(world_size, n_elem)`로부터 인자
+   `(n_elem, cube_w, cube_h, n_sips)`를 구성한다.
+3. `(sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)`를 추가하며,
+   여기서 `sip_rank`는 현재 greenlet에 바인딩된 rank이다.
+4. `_defer_wait=True`로 launch; 모든 워커가 제출한 후 메인 스케줄러가
+   pending 핸들을 드레인한다 (ADR-0027 D0.4).
+
+### D6. 구성 스키마
+
+`ccl.yaml`:
+
+```yaml
+defaults:
+  algorithm: lrab_hierarchical_allreduce
+  buffer_kind: tcm
+  ...
+
+algorithms:
+  lrab_hierarchical_allreduce:
+    module: kernbench.ccl.algorithms.lrab_hierarchical_allreduce
+    topology: none
+    buffer_kind: tcm
+    n_elem: 8
+    root_cube: 15   # 현재 사용되지 않음 — 커널이 루트를 기하학적 중심으로
+                    # 동적으로 선출한다 (D1 참조). 향후 명시적 루트 override /
+                    # 런타임 선출 훅을 위한 placeholder로 유지한다.
+```
+
+`topology.yaml`:
+
+```yaml
+system:
+  sips:
+    count: 2
+    topology: ring_1d
+sip:
+  cube_mesh: { w: 4, h: 4 }
+```
+
+### D7. 알고리즘 모듈 계약
+
+`cfg["module"]`로 로드되는 모듈은 다음을 export해야 한다:
+
+| Name | Purpose |
+|---|---|
+| `kernel` | callable, 시그니처 `(t_ptr, n_elem, cube_w, cube_h, n_sips, sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h, tl)` |
+| `kernel_args(world_size, n_elem) -> tuple` | 처음 4개의 scalar 인자(텐서별) 반환 |
+| `TOPO_NAME_TO_KIND: dict[str, int]` | `system.sips.topology` 이름을 커널 분기 코드로 매핑 |
+| `SIP_TOPO_RING`, `SIP_TOPO_TORUS`, `SIP_TOPO_MESH` | 정수 상수 (0, 1, 2) |
+
+---
+
+## Dependencies
+
+- **ADR-0023**: IPCQ 프로토콜 (이웃 테이블, 송수신, credit 반환).
+- **ADR-0024**: rank = SIP launcher, `mp.spawn`, greenlet-로컬 rank.
+- **ADR-0025**: 주소 기반 IPCQ 방향 매칭; `global_*` 쌍으로 확장된
+  `_OPPOSITE_DIR`.
+- **ADR-0027**: 메인 스케줄러에서의 worker-wait / 집합 통신 pending
+  드레인.
+
+## Non-goals
+
+- **PE별 allreduce** (큐브 내 PE-PE reduce). 범위 밖 — 본 알고리즘의
+  워크로드는 큐브당 DP이다.
+- **정사각 그리드 fallback은 `n_sips = k²`를 요구**: 직사각형 SIP
+  그리드(정사각형이 아닌 메시/토러스)는 지원되지만, `system.sips.w/h`를
+  명시적으로 줄 때만 가능하다 (ADR-0024 D5). `w/h` 생략 시 2D 토폴로지는
+  정사각 그리드로 fallback하며 여전히 `n_sips = k²`를 요구한다.
+- **파이프라인 청크**: 큐브당 단일 타일, 아직 파이프라이닝 없음.
+- **루트 큐브의 런타임 선출**: 커널은 현재 SIP 내부 임계 경로를
+  최소화하기 위해 기하학적 중심인
+  `root_cube = (mesh_h // 2) * mesh_w + (mesh_w // 2)`을 사용한다. SFR
+  와이어링이 모든 큐브를 커버하므로, 필요해질 때 다른 루트를 런타임에
+  선출하는 것은 순수 커널 변경이다.
+
+---
+
+## Consequences
+
+### Positive
+
+- **단일 커널, 단일 설치 경로**로 all-reduce를 처리 — 제거된 네 개의
+  모듈(`ring`, `mesh`, `tree`, `hierarchical`)을 대체한다.
+- **토폴로지 무관 커널**: ring / torus / mesh를 정수 파라미터 하나로
+  선택, 커널 중복 없음.
+- **`dist.all_reduce`를 통한 자동화**: 벤치 레벨이나 사용자 레벨의
+  알고리즘 선택 불필요; end-to-end 구성 기반.
+- **완전한 SFR 와이어링**: 모든 SIP의 모든 큐브가 SIP 간 링크를 보유 —
+  향후 동적 루트 큐브 선출을 지원한다.
+
+### Negative
+
+- **PE별 샤딩된 텐서에 부적합**: 큐브 하나 내부에서 8개 PE에 걸쳐
+  샤딩되는 TP-레이어 스타일 텐서는 본 커널로 주소 지정할 수 없다. 이러한
+  워크로드에는 별도의 큐브 내 all-reduce 경로가 필요하다 (아직 구현되지
+  않음).
+- **`configure_sfr_intercube_multisip`는 항상 모든 pe0을 와이어링**:
+  주어진 실행이 부분집합(예: 1 SIP, ring만)만 필요하더라도. 설치 비용은
+  작지만 영(zero)은 아니다.
+
+---
+
+## Affected files
+
+| File | Change |
+|---|---|
+| `src/kernbench/ccl/algorithms/lrab_hierarchical_allreduce.py` (신규) | 커널 + `_inter_sip_*` 헬퍼 + `TOPO_NAME_TO_KIND` |
+| `src/kernbench/ccl/sfr_config.py` (신규) | `configure_sfr_intercube_multisip` |
+| `src/kernbench/ccl/topologies.py` | `torus_2d`, `mesh_2d_no_wrap` 추가 |
+| `src/kernbench/ccl/install.py` | `_OPPOSITE_DIR`을 `global_*` 쌍으로 확장 |
+| `src/kernbench/runtime_api/distributed.py` | `AhbmCCLBackend`가 `configure_sfr_intercube_multisip` 사용 + sip_rank/topo 인자 추가 |
+| `ccl.yaml` | 단일 `lrab_hierarchical_allreduce` 항목 |
+| `topology.yaml` | `system.sips.topology` 추가 |
+| `benches/ccl_allreduce.py` | Row-wise 큐브-메시 텐서 레이아웃 |
+| `tests/sccl/` (테스트 패키지) | 구성 기반 ring/torus/mesh 정확성 + 전체 `dist.all_reduce` 경로 + latency/buffer-kind 스윕 (평가 하니스 — ADR-0043) |
+| `tests/test_intercube_sfr_config.py` | SFR 와이어링 검증 |
+| 제거 | `ring_allreduce.py`, `mesh_allreduce.py`, `tree_allreduce.py`, `hierarchical_allreduce.py`, `hello_send.py`, `testing.py` 및 그 테스트 |

docs/adr-ko/ADR-0033-lat-latency-model-assumptions.md

@@ -0,0 +1,152 @@
+# ADR-0033 — 레이턴시 모델: 가정 및 알려진 단순화
+
+## Status
+
+Accepted
+
+## Context
+
+이 시뮬레이터는 분석적·이벤트 기반 성능 모델이지, 사이클 정확(cycle-accurate)
+시뮬레이터나 RTL 수준 시뮬레이터가 아니다. 실제 HW의 많은 효과들이 설계상
+근사되거나 생략되었다. 모델 전체를 감사·리뷰할 수 있도록 유지하기 위해,
+본 ADR은 그런 가정들을 한 곳에 통합한다. 개별 컴포넌트 ADR(ADR-0015,
+ADR-0017, ADR-0004)들이 *메커니즘*을 정의하고, 본 문서는 *충실도의 한계*를
+정의한다.
+
+## Decisions
+
+### D1. 정밀하게 모델링되는 것
+
+- **방향 에지별 BW 점유** (`available_at`을 통한 FIFO 직렬화) —
+  ADR-0015 D2.
+- **컴포넌트별 스위칭/오버헤드 레이턴시** (`overhead_ns` attr).
+- **HBM pseudo-channel별 병렬성**: 주소 기반 PC 선택을 동반한
+  stateless `pc_avail[N]` 배열로 (ADR-0034 D3). 버스트 granularity는 조정 가능
+  (`burst_bytes`, 기본 256B). 각 PC의 `available_at`은 read와 write가 공유한다
+  (실제 HW의 명령 버스가 PC별로 공유되기 때문).
+- **HBM 방향 전환 페널티 메커니즘**: PC별 last-direction 추적 +
+  설정 가능한 `switch_penalty_ns`. 기본값 0 — D2 참조.
+- **와이어 청크 스트리밍 (Phase 2c)**: 각 와이어는 payload가 있는
+  Transaction을 `flit_bytes` 단위의 `Flit` 객체로 분해한다(기본 = HBM
+  `burst_bytes` = 256B). 와이어는 각 flit을 `prop_ns + flit_nbytes/bw_gbs`
+  이후에 개별적으로 방출하므로 링크의 대역폭이 실제 HW의 wormhole 시맨틱대로
+  flit 도착률을 조절한다.
+- **방향 에지별로 분리된 Store** (Phase 2c 핵심 수정): 와이어는
+  `src.out_ports[dst]`와 `dst.in_ports[src]` 사이의 *유일한* 통로이다.
+  이전에는 둘이 동일한 `simpy.Store`로 별칭되어 있었다. 와이어가 청크화된
+  flit을 되돌려 넣을 때 목적지의 `fan_in`이 와이어가 대역폭 지연을 적용하기
+  전에 그것을 끌어가, flit의 절반이 병목을 우회할 수 있었다.
+- **Flit 인지 pass-through** (`TransitComponent`, `HbmCtrlComponent`):
+  각 flit을 직렬로 전달하며 트랜잭션 오버헤드는 첫 flit 도착 시점에 한 번만
+  적용된다(헤더 디코드 모델). 이후의 flit들은 추가 지연 없이 파이프라인을
+  통과한다. 다중 hop 경로 전반에서 wormhole이 자연스럽게 발현된다.
+- **HBM CTRL의 flit별 PC commit**: HBM CTRL에 도착하는 각 flit은
+  `max(env.now, pc_avail[pc]) + chunk_time`에 PC commit을 스케줄하며,
+  `is_last` flit이 마지막 PC commit을 기다린 후 `txn.done`을 신호한다.
+- **Flit 비인지 컴포넌트(기본)는 ``_fan_in``에서 flit을 재조립**하여
+  레거시 `_forward_txn` 경로가 실행되도록 한다. 이는 아직 flit 인지
+  처리로 마이그레이션되지 않은 컴포넌트(예: `MCpuComponent`,
+  `IoCpuComponent`의 sub-txn 생성기)에 대한 하위 호환성을 보존한다. 그런
+  컴포넌트들은 *leg 경계마다 한 번* 재조립하며, hop마다는 아니다 —
+  flit 인지 라우터 체인을 통한 다중 hop wormhole 타이밍이 보존된다.
+
+### D2. 근사됨 (알려진 방향성 오차와 함께)
+
+| 효과 | 실제 HW | 본 모델 | 오차 방향 |
+|--------|---------|-----------|----------------|
+| 라우터 출력 포트 중재 | Round-robin / weighted | 와이어 에지 FIFO + 직렬 워커 | 사이클당 한 txn일 때 공정; multi-stream 공유는 flit 수준에서 모델링 안 됨 |
+| HBM 스케줄러 / 쓰기 버퍼 | FR-FCFS + watermark drain | FIFO, 재정렬 없음 | 교번이 조밀한 혼합 R/W에 대해 비관적 — 기본 `switch_penalty_ns = 0`은 이상적 스케줄러가 amortize한다고 가정 |
+| Flit ↔ burst granularity | 32B flit < 256B burst | `flit_bytes = burst_bytes = 256B` | sub-flit 미세 타이밍 노이즈; 매우 작은 와이어 중재 윈도우에서만 영향 |
+| 와이어 수준 RR 공정성 | 공유 링크에서 사이클별 multi-flow 중재 | 에지마다 단일 직렬 와이어 프로세스 | 주어진 에지에 한 트랜잭션만 in-flight일 때만 공정. 동일 에지에서 동시 멀티 스트림 트래픽은 FIFO 순서로 직렬화됨 |
+
+### D3. 무시됨 (범위 외)
+
+- 뱅크 수준의 row buffer 충돌 페널티 (충돌 없음 가정 — 최적 케이스;
+  모델은 PC 내부에 뱅크별 상태를 갖지 않으므로 동일 뱅크 재사용을 감지할 수 없다).
+- HBM tRP / tRCD / tFAW / tRC 타이밍 제약 (정상 상태의
+  `burst_time = burst_bytes / pc_bw_gbs`에 흡수).
+- 리프레시, ECC, 열 throttling, 전력 게이팅.
+- 클럭 도메인 교차, PLL lock 시간.
+- 하위 버퍼 점유로 인한 상위 backpressure (입력 포트는 unbounded
+  `simpy.Store`를 사용).
+- 라우터에서의 sub-flit 사이클 수준 중재 (flit granularity가 본 모델의
+  최소 단위).
+
+### D4. 워크로드 민감도
+
+위 단순화들이 결과에 의미 있게 영향을 미치는 워크로드:
+
+- **무작위 scatter/gather**: 뱅크 충돌 무시 → 모델이 낙관적.
+- **혼합 R/W가 강한 워크로드** (예: GEMM 바이어스 누적): HBM 스케줄러
+  부재. 기본 `switch_penalty_ns = 0`은 이상적 amortization을 가정;
+  0이 아닌 값은 교번당 비관적 비용을 모델링.
+- **고동시성 (한 링크에 활성 흐름 >10개)**: HoL blocking과 VC 제한이
+  모델링되지 않음 → 모델이 낙관적.
+- **매우 작은(sub-flit) 트랜잭션**: flit 양자화 노이즈.
+- **단일 와이어상의 동시 multi-flow**: 와이어는 flit 수준에서 직렬
+  FIFO이므로 단일 에지 내에서의 흐름별 공정성은 모델링되지 않는다.
+  Pre-edge 병합(여러 source가 라우터에 도착하여 동일한 downstream
+  와이어로 전달되는 경우)은 flit 인지 라우터의 직렬 워커를 통해 올바르게
+  모델링된다.
+
+### D5. 검증 정책
+
+D4의 워크로드에 대해 절대값 결론을 내리기 전에 실제 HW나 사이클 정확
+시뮬레이터와 cross-check 할 것. 모델은 모델링된 영역 내에서의 **상대적
+비교**에 대해서는 여전히 정확하다.
+
+### D6. 향후 작업
+
+참고: 라우터에서의 multi-stream 병합은 올바르게 모델링되고 있다 — 각
+in_port가 자신의 fan_in 프로세스를 가지며 모두 공유 인박스로 push하고,
+라우터 워커가 인박스 FIFO 순서로 전달한다. 서로 다른 상위 스트림의 flit들이
+flit granularity에서 자연스럽게 인터리브된다. 아래 항목들은 별개의 관심사이며,
+예상되는 워크로드 영향 순으로 정렬되어 있다.
+
+**영향이 큼 (워크로드 정확도 격차)**:
+
+- [ ] PC 내의 **뱅크 수준 충돌 모델링** (`track_banks: true`로 opt-in).
+  현재는 동일 뱅크 재사용이 없다고 가정; 무작위 scatter/gather 워크로드는
+  이 부분에서 낙관적이다.
+- [ ] write buffer + watermark drain을 동반한 **HBM 스케줄러** (설계
+  논의에서의 Tier 2). 기본 `switch_penalty_ns=0`은 이상적 amortization의
+  stand-in; 버스티한 혼합 R/W 워크로드는 명시적 모델링으로부터 이득을 본다.
+- [ ] 유한한 컴포넌트 버퍼에 대한 **Backpressure** 모델링. 버퍼 점유가
+  상위 stall을 유발하는 고동시성/지속적 포화 상황에서 중요.
+- [ ] **청크 스트리밍과 op_log 통합**: 현재 op_log는 청크화되지 않는
+  PE 내부 명령 메시지(DmaReadCmd, DmaWriteCmd, GemmCmd, MathCmd)에 대해
+  발화한다. 통합은 flit 인지 컴포넌트들이 트랜잭션당 op_log start/end
+  hook(첫 flit에 start, is_last에 end)을 함께 방출하도록 요구한다.
+
+**영향이 작음 (학술적 / 특정 use case)**:
+
+- [ ] **사이클 정확 라우터 중재 정책** (우선순위·age를 동반한 RR, iSLIP).
+  FIFO 인박스는 스트림 간 flit 도착 시간이 약간씩 다를 때 이미 근사적으로
+  공정하다(유사한 비율의 워크로드에서 흔한 경우). 실질적 영향은 (a)
+  우선순위/QoS 모델링, (b) 지속적 포화에서의 스트림별 tail latency 분석에서만
+  나타난다. makespan이나 평균 레이턴시 연구에는 결정적이지 않음.
+- [ ] 더 미세한 와이어 중재 사이클을 위한 **Sub-flit (32B) granularity**.
+  본 모델의 `flit_bytes`는 burst(256B)와 같지만, 실제 HW는 32B flit마다
+  중재한다. 대부분 워크로드에서는 영향이 작다(작은 메시지에 대한 sub-flit
+  타이밍 노이즈).
+
+## Consequences
+
+- 모든 모델 충실도 질문에 대한 단일 리뷰 지점. 레이턴시를 건드리는 향후
+  모든 PR은 본 문서의 해당 절을 갱신해야 한다.
+- 워크로드별 규모 오차 envelope이 명시적이다.
+- 빌더측 `pc_bw_gbs = hbm_to_router_bw_gbs / num_pcs` 유도가
+  yaml의 수동 일관성에 의존하지 않고 코드 내에서 ADR-0017 D8의 불변성을
+  강제한다.
+- 와이어 전송 시간은 터미널의 `drain_ns` 주입을 통해서가 아니라
+  병목 링크 통과당 한 번 부과된다(Phase 2c flit별 타이밍). 단일 트랜잭션은
+  `drain + commit_time + small_overheads`에 도달; 다중 hop은 wormhole
+  파이프라이닝을 보존; multi-stream 병합은 공유 와이어의 FIFO에서 올바르게
+  직렬화된다.
+
+## Cross-references
+
+- ADR-0015 — 컴포넌트 / 포트 / 와이어 모델.
+- ADR-0017 — 큐브 NOC 아키텍처 및 HBM 연결성.
+- ADR-0004 — 메모리 시맨틱, 로컬 HBM.
+- ADR-0034 — HBM 컨트롤러 내부 설계.

docs/adr-ko/ADR-0034-dev-hbm-controller-internal-design.md

@@ -0,0 +1,263 @@
+# ADR-0034: HBM 컨트롤러 내부 설계
+
+## Status
+
+Accepted
+
+## Context
+
+`HbmCtrlComponent`는 큐브 NOC의 말단(leaf)에 위치하는 PE별 HBM
+파티션 엔드포인트이다. 토폴로지 노드
+`sip{S}.cube{C}.hbm_ctrl.pe{idx}` 아래에 PE마다 하나의 인스턴스가
+생성되며 해당 PE의 라우터에 연결된다 (ADR-0017 D4). 본 컴포넌트는
+의사 채널(PC, pseudo-channel)별 스케줄링, 버스트 단위 커밋 타이밍,
+주소 기반 PC 선택, 그리고 응답을 요청자에게 되돌리는 라우팅을
+모델링한다.
+
+본 ADR은 현재 구현된 컴포넌트를 문서화한다. ADR-0017 D4/D8은 HBM CTRL이
+*어디에* 부착되는지와 *어떤* 집계 대역폭을 제공해야 하는지를 정의한다.
+ADR-0033 D1/D2는 HBM 모델링의 *어떤 정밀도(fidelity)*가 범위에 포함되는지를
+정의한다. 본 ADR은 그 둘 사이의 공백 — 인스턴스별 내부 스케줄링 모델을
+채운다.
+
+## Decision
+
+### D1. 역할
+
+`HbmCtrlComponent`는 PE별 HBM 파티션 엔드포인트이다. PE당 하나의
+인스턴스(큐브당 기본 8개, `cube.memory_map.hbm_slices_per_cube`로 설정)가
+`cube_mesh.yaml`의 `peX.hbm` 부착 목록을 통해 해당 PE의 라우터에 연결된다
+(ADR-0017 D4). 기본 n:1 채널 매핑(ADR-0017 D8)에서는 인스턴스가
+`channels_per_pe`개의 의사 채널을 하나의 엔드포인트로 집계한다.
+
+본 컴포넌트는 다음을 모델링한다:
+
+- PC별 스케줄링(D2) 및 R/W 명령 버스 공유.
+- 주소 기반 PC 선택(D3).
+- 버스트 단위 커밋 타이밍(D4).
+- Flit 인지 per-flit PC 커밋 및 비동기 finalize(D5, D6).
+- 읽기 데이터 드레인(drain)을 위한 명령 전용 Transaction 처리(D7).
+- 요청자에게 되돌리는 응답 라우팅(D8).
+
+다음은 모델링하지 않는다:
+
+- Bank 수준의 row-buffer 충돌, refresh, ECC, 열 스로틀링
+  (ADR-0033 D3).
+- 자신의 라우터 엣지를 넘어가는 PE 간 HBM 경합(라우터 메시가 처리 —
+  ADR-0017 D3).
+- 1:1 채널 모드(ADR-0017 D8 향후 작업).
+
+### D2. PC별 스케줄링 모델
+
+`start()`에서 초기화되는 인스턴스별 상태:
+
+- `_pc_avail: list[float]` — 각 PC가 다음에 자유로워지는 가장 빠른
+  시뮬레이션 시각; 길이 `num_pcs`, 초기값 0.0.
+- `_pc_last_dir: list["R"|"W"|None]` — 각 PC의 마지막 커밋 방향, 스위치
+  페널티 감지에 사용(D4); 초기값 `None`.
+
+`num_pcs`와 `burst_bytes`는 각각 양의 2의 거듭제곱이어야 주소 기반 PC
+선택(D3)이 시프트와 마스크로 축약된다.
+
+읽기와 쓰기 요청은 PC별로 동일한 `_pc_avail` 슬롯을 공유한다 — 실제 HW에서
+PC별 명령 버스는 읽기와 쓰기 트래픽이 공유하므로, PC k에 쓰기를 발행하면
+PC k에 대한 후속 읽기가 정확히 버스트 시간만큼 블록된다.
+
+요청의 방향 `dir`은 요청 타입으로부터 추론된다:
+
+- `MemoryWriteMsg` → `"W"`.
+- `is_write=True`인 `PeDmaMsg` → `"W"`.
+- 그 외 전부(`MemoryReadMsg`, 읽기 `PeDmaMsg`) → `"R"`.
+
+### D3. 주소 기반 PC 선택
+
+접근에 대한 PC 인덱스는 접근 주소로부터 시프트와 마스크로 도출된다:
+
+```text
+pc_shift = log2(burst_bytes)         # 기본값 8  (burst=256B)
+pc_mask  = num_pcs - 1               # 기본값 7  (8 PCs)
+pc       = (address >> pc_shift) & pc_mask
+```
+
+대안적인 `(burst_bytes, num_pcs)` 쌍과의 정합성을 유지하기 위해
+`start()`에서 토폴로지 설정으로부터 한 번 계산된다. 정규 기본값
+`(256, 8)`에서는 PC 선택 필드가 HBM 바이트 오프셋의 비트 `[10:8]`에
+배치된다: 비트 `[7:0]`은 버스트 내부(같은 PC), 비트 `[10:8]`은 3비트
+PC 인덱스, 비트 `[36:11]`은 PC 슬라이스 내부의 row/bank/column이다
+(`phyaddr.py` 주석 참조).
+
+주소 기반 스트라이핑은 — 주소를 보지 않는 전역 라운드로빈과 달리 —
+오프셋이 분리된 동시 전송들에 대해 PC 병렬성을 보존한다: 각 전송의
+버스트는 자신의 바이트 주소가 함의하는 PC 집합 위에 결정론적으로
+떨어지므로, 분리된 영역에 접근하는 멀티 PE 워크로드가 단일 PC에서
+충돌하지 않는다.
+
+### D4. 버스트 단위 시간 및 PC 커밋 타이밍
+
+단일 PC 커밋에 걸리는 시간:
+
+```text
+chunk_time = burst_bytes / pc_bw_gbs    # ns
+```
+
+- `burst_bytes`(기본 256)는 flit 크기와 일치하는 버스트 단위이다
+  (ADR-0033 D1).
+- `pc_bw_gbs`는 **빌더에서 도출**된다:
+  `hbm_to_router_bw_gbs / num_pcs` (`topology/builder.py`). 이는 PE당
+  집계 대역폭이 라우터-HBM 링크 대역폭과 같아야 한다는 ADR-0017 D8의
+  불변식을 강제한다.
+
+방향 `dir`로 PC `pc`에 도착한 접근에 대한 PC별 커밋 스케줄링:
+
+```text
+switch_cost = switch_penalty_ns
+              if pc_last_dir[pc] not in (None, dir) else 0
+start  = max(env.now, pc_avail[pc]) + switch_cost
+finish = start + chunk_time
+pc_avail[pc]    = finish
+pc_last_dir[pc] = dir
+```
+
+기본 `switch_penalty_ns = 0` — 이상적인 HBM 스케줄러가 R/W 스위칭
+비용을 분할 상환한다는 Tier 0 가정(ADR-0033 D2). 0이 아닌 값은
+교차마다 발생하는 비관적 비용을 모델링한다.
+
+### D5. Flit 인지 per-flit PC 커밋 (주 경로)
+
+`_handle_flit`이 주 워커 경로이다. 각 도착 `Flit`에 대해:
+
+1. 트랜잭션의 **첫 번째** flit인 경우(`tid = id(txn)`가 `_txn_state`에
+   없는 경우):
+   - `run(env, nbytes)`를 통해 `overhead_ns`를 한 번 적용 — 헤더 디코드
+     모델, first-flit overhead 패턴(ADR-0033 D1).
+   - `_txn_state[tid] = {"last_finish": env.now}`로 초기화.
+2. `pc = _pc_for_address(flit.address)`를 계산(D3).
+3. 요청 방향(D2)을 사용하여 PC별 스케줄(D4)을 적용.
+4. `state["last_finish"] = max(state["last_finish"], finish)`로 갱신.
+5. `flit.is_last`이면: `_txn_state[tid]`를 pop하고 `_finalize_txn`을
+   spawn(D6).
+
+per-flit 주소 인지 커밋이 분리된 HBM 오프셋으로 향하는 동시 멀티 PE
+트래픽이 서로 다른 PC를 통해 병렬로 파이프라인되도록 하는 메커니즘이다.
+
+### D6. 트랜잭션별 비동기 finalize
+
+트랜잭션의 마지막 flit이 스케줄링되고 나면, finalize는 별도로 spawn된
+프로세스에서 실행된다:
+
+```python
+def _finalize_txn(env, txn, last_finish):
+    wait = last_finish - env.now
+    if wait > 0:
+        yield env.timeout(wait)
+    yield from _send_response(env, txn)
+```
+
+`_handle_flit`은 이를 `env.process(...)`로 spawn한 뒤 즉시 반환하므로,
+마지막 PC 커밋이 드레인되는 동안에도 워커는 다음 inbox 메시지를 집어들
+수 있다.
+
+이 분리가 없다면 — 즉 워커 자신이 `yield env.timeout(wait)`를 한다면 —
+서로 다른 PC에 떨어지는 주소를 가진 동시 단일 flit 트랜잭션들도 결국
+워커 내부에서 각각 `chunk_time`만큼 직렬화되어, D3와 D5가 노출하려고
+설계한 PC 병렬성을 숨겨버린다.
+
+### D7. 명령 전용 트랜잭션을 위한 non-flit 폴백
+
+`_handle_txn`은 inbox가 `Flit`이 아닌 `Transaction`을 전달할 때 실행된다.
+이는 와이어가 flit으로 분할하지 않는 명령 전용 요청에 대한 경로로 —
+대표적으로 명령 트랜잭션이 `nbytes=0`을 운반하는 `MemoryReadMsg`가
+해당한다(데이터 드레인은 HBM CTRL 후처리에서 모델링되며, 인바운드
+flit으로 모델링되지 않는다).
+
+절차:
+
+1. `work_bytes = txn.nbytes if txn.nbytes > 0 else int(request.nbytes or 0)`
+   — 읽기 명령의 경우 작업량은 요청으로 결정된다.
+2. `work_bytes > 0`이면 `n_chunks = ceil(work_bytes / burst_bytes)`,
+   아니면 0.
+3. 둘 다 > 0일 때 `chunk_interval = drain_ns / n_chunks` — 청크는
+   `drain/n_chunks` ns 간격으로 시간상에 스케줄링되어 병목 링크의 데이터
+   도착 속도를 모델링한다(ADR-0033 D1 청크 루프 드레인).
+4. `overhead_ns`를 위해 `run(env, txn.nbytes)`를 한 번 적용.
+5. 각 청크 `i`에 대해 `chunk_interval` ns만큼 진행한 뒤
+   `pc = _pc_for_address(base_address + i * burst_bytes)`로 D4 스케줄을
+   적용.
+6. 모든 청크 스케줄링 후 `last_finish - env.now`만큼 대기한 다음
+   `_send_response`를 호출.
+
+`_handle_txn`은 `_handle_flit`과 동일한 `_pc_avail` / `_pc_last_dir`
+상태를 공유한다 — 두 경로에 걸쳐 PC 스케줄링의 단일 진실 원천이 정확히
+하나만 존재한다.
+
+### D8. 응답 라우팅
+
+`_send_response`는 요청 타입과 경로 형상에 따라 디스패치한다:
+
+| 경우 | 트리거 | 응답 |
+| --- | --- | --- |
+| PE_DMA | `isinstance(txn.request, PeDmaMsg)` | 신규 역방향 경로 Transaction(`is_response=True`, `nbytes=0`), 동일한 `done` |
+| Bypass — Memory Read | `"m_cpu" not in any(txn.path)` AND `MemoryReadMsg` | `nbytes=request.nbytes`(데이터 반환)인 역방향 경로 Transaction |
+| Bypass — Memory Write | `"m_cpu" not in any(txn.path)` AND not Memory Read | `txn.done.succeed()` (쓰기는 로컬에서 완료) |
+| 기본 | 그 외 | 역방향 경로상의 신규 `ResponseMsg(correlation_id, request_id, src_cube, src_pe, success=True)` |
+
+"bypass" 분류는 ADR-0015 D4에서 정의된 Memory R/W 패브릭 경로(PCIE_EP →
+io_noc → ucie → 큐브 라우터 → hbm_ctrl, M_CPU 미경유)와 일치한다.
+PE_DMA 케이스는 내부 루프 DMA를 빠르게 유지하기 위한 전용 역방향 경로이다
+(PE_DMA 읽기/쓰기는 ResponseMsg 봉투를 합성하지 않는다).
+
+모든 역방향 경로 케이스에서, 응답 Transaction은
+`out_ports[reverse_path[1]]` — 기록된 정방향 경로를 따라 되돌아가는 첫
+홉 — 에 put된다. `reverse_path`의 엔트리가 2개 미만이면(축퇴된 경로),
+원래의 `txn.done`이 직접 시그널된다.
+
+### D9. 설정 가능한 속성
+
+| 속성 | 기본값 | 출처 | 비고 |
+| --- | --- | --- | --- |
+| `num_pcs` | 8 | 토폴로지 큐브 `hbm_ctrl.attrs` | 2의 거듭제곱이어야 함 |
+| `pc_bw_gbs` | 32.0 | 빌더 도출: `hbm_to_router_bw_gbs / num_pcs` | ADR-0017 D8 불변식 강제 |
+| `burst_bytes` | 256 | 토폴로지 attrs | 2의 거듭제곱이어야 함; `flit_bytes`와 동일(ADR-0033 D1) |
+| `switch_penalty_ns` | 0.0 | 토폴로지 attrs | Tier 0 기본값; 0이 아니면 비관적 R/W 스위칭 모델링 |
+| `efficiency` | 1.0 | 토폴로지 attrs | 빌더 시점에 `hbm_to_router_bw_gbs`에 적용(라우터 엣지 BW 스케일링만) |
+| `overhead_ns` | 0.0 | 토폴로지 attrs | First-flit 디코드 오버헤드(D5) |
+
+`pc_bw_gbs`는 yaml 측 중복 없이 PE당 집계 대역폭을 라우터-HBM 링크
+대역폭과 일치시키기 위해 직접 설정되지 않고 `topology/builder.py`에서
+도출된다.
+
+## Consequences
+
+### Positive
+
+- 주소 기반 PC 선택은 주소를 보지 않는 라운드로빈이 무너뜨릴 멀티 스트림
+  HBM 병렬성을 보존한다 — 분리된 HBM 영역을 갖는 멀티 PE 워크로드에서
+  중요하다.
+- Flit 인지 경로(D5) + 비동기 finalize(D6)는 웜홀 파이프라이닝을
+  보존하며, 연속적인 단일 flit 트랜잭션에 대해 PC 병렬성을 노출한다.
+- PC 스케줄링의 단일 진실 원천(D4 메커니즘이 D5 flit 경로와 D7 청크 루프
+  경로 모두에서 사용됨).
+- 빌더 도출 `pc_bw_gbs`가 yaml 규율이 아닌 코드에서 ADR-0017 D8을
+  강제한다.
+
+### Negative
+
+- PC 내부의 bank 수준 충돌 모델링이 없음; bank/row-buffer 재사용에
+  주소-무관(ADR-0033 D3).
+- HBM 스케줄러 없음(FR-FCFS / write-buffer / watermark drain); PC당 고정
+  FIFO. 버스티한 혼합 R/W는 `switch_penalty_ns`로 근사화된다
+  (ADR-0033 D2).
+- `_txn_state`는 `id(txn)`로 키를 잡는 일반 dict이다; 동시 트랜잭션마다
+  in-flight 상태가 누적되며 `is_last` 시에만 제거된다. 현재 워크로드에는
+  충분하다.
+
+## Links
+
+- ADR-0001 (물리 주소 레이아웃 — PC 비트 필드 주석)
+- ADR-0015 D4 (Memory R/W 패브릭 경로 — bypass 응답 케이스)
+- ADR-0017 D4 (PE별 HBM 파티셔닝 — PE 라우터로의 부착)
+- ADR-0017 D8 (HBM 채널 매핑 모드 — 본 ADR이 구현하는 n:1 집계)
+- ADR-0017 D9 (AddressResolver — `hbm_ctrl.pe{pe_id}` 엔드포인트 해석)
+- ADR-0033 D1 (정확한 모델링 — PC별 병렬성, 스위치 페널티, flit 인지
+  PC 커밋, first-flit 오버헤드, 청크 루프 드레인)
+- ADR-0033 D2 (스위치 페널티 기본값 0 — 이상적 스케줄러의 분할 상환)

docs/adr-ko/ADR-0035-dev-m-cpu-and-m-cpu-dma-component-model.md

@@ -0,0 +1,273 @@
+# ADR-0035: M_CPU 및 M_CPU.DMA 컴포넌트 모델
+
+## Status
+
+Accepted
+
+## Context
+
+M_CPU는 큐브 수준의 명령 프로세서이다. IO_CPU로부터(또는 엔진이
+Memory R/W를 폴백으로 M_CPU를 거쳐 라우팅할 때 PCIE_EP로부터) 명령을
+수신하여 자신의 큐브 내 PE들로 팬아웃하고, PE별 응답을 단일 ResponseMsg로
+집계하여 역방향 경로를 통해 IO_CPU로 되돌려 보낸다.
+
+M_CPU.DMA는 Memory R/W 팬아웃을 처리하는 큐브 수준의 DMA 채널 쌍이다.
+ADR-0015 D5에 따라 별도의 토폴로지 노드가 **아니다** — `MCpuComponent`의
+내부 상태로서 존재한다.
+
+본 ADR은 위의 책임들을 실현하는 M_CPU 컴포넌트 구현을 문서화한다. 여기에는
+세 가지 구별되는 팬아웃 경로(Memory R/W, Kernel Launch, MMU Map/Unmap),
+M_CPU.DMA 자원 모델, 그리고 응답 집계 계약이 포함된다.
+
+## Decision
+
+### D1. 역할
+
+M_CPU는 세 가지 책임을 갖는다:
+
+1. **Transit 포워딩** — 종단 홉이 아닐 때(예: 역방향 응답 경로 PE →
+   M_CPU → IO_CPU), 사전 계산된 경로의 `next_hop`으로 Transaction을
+   전달한다.
+2. **종단 홉에서의 멀티 PE 팬아웃** — 요청 타입에 따라 세 팬아웃 경로
+   중 하나로 디스패치한다(D2).
+3. **응답 집계** — PE별 응답을 수집하여 역방향 경로를 통해 단일 집계
+   ResponseMsg를 IO_CPU로 되돌려 보낸다.
+
+호출당(`run()`): 들어오는 Transaction마다 `overhead_ns`를 한 번 적용한다.
+
+M_CPU는 다음을 하지 **않는다**:
+
+- 라우팅 결정 — 경로는 라우터에 의해 사전 계산된다(ADR-0002).
+- PE 내부 실행 처리 — PE_CPU / PE_SCHEDULER / 엔진들이 담당(ADR-0014).
+- 주소 디코드 — `ctx.resolver.resolve(pa)`가 PE별 `hbm_ctrl.pe{X}`를
+  직접 반환한다(ADR-0017 D9).
+- 텐서 또는 커널 의미 해석 — 팬아웃 디스패치는 Python isinstance
+  체크만으로 이루어진다.
+
+### D2. 요청 타입으로 디스패치되는 세 가지 팬아웃 경로
+
+종단 홉에서 워커는 요청 타입에 따라 디스패치한다:
+
+```python
+elif self.ctx is not None and txn.request is not None:
+    if isinstance(txn.request, KernelLaunchMsg):
+        env.process(self._kernel_launch_fanout(env, txn))
+    elif isinstance(txn.request, (MmuMapMsg, MmuUnmapMsg)):
+        env.process(self._mmu_msg_fanout(env, txn))
+    else:
+        env.process(self._dma_fanout(env, txn))
+```
+
+각 경로는 서로 다른 라우터 메서드를 사용한다:
+
+- `_dma_fanout`은 `ctx.router.find_mcpu_dma_path()`를 사용 — PE 파이프라인
+  노드를 우회하는 M_CPU 전용 DMA 경로.
+- `_kernel_launch_fanout`은 `ctx.router.find_node_path()`를 사용 — PE_CPU로
+  향하는 범용 NOC 명령 경로.
+- `_mmu_msg_fanout`은 `ctx.router.find_node_path()`를 사용 — PE_MMU로
+  향하는 NOC 명령 경로.
+
+### D3. M_CPU.DMA 내부 서브 컴포넌트 (ADR-0015 D5)
+
+`MCpuComponent.start()`는 두 개의 SimPy 자원을 초기화한다:
+
+```python
+self._dma_write = simpy.Resource(env, capacity=1)  # MemoryWriteMsg
+self._dma_read  = simpy.Resource(env, capacity=1)  # MemoryReadMsg
+```
+
+특성:
+
+- **토폴로지 노드가 아님** — 전적으로 `MCpuComponent` 내부에서 관리됨;
+  `topology.yaml`이나 컴파일된 그래프에 나타나지 않는다.
+- **독립된 읽기/쓰기 채널** — 동시 in-flight Memory R/W가 허용된다.
+- **채널당 capacity=1**은 본 M_CPU에서 동시 in-flight Memory R/W 요청의
+  **디스패치 단계**(`yield self.out_ports[...].put(...)`)를 직렬화한다.
+  실제 패브릭 전송 시간은 컴포넌트 사이의 와이어 프로세스(ADR-0015 D2)와
+  종단 홉의 `drain_ns`로 모델링되며, DMA 자원은 전송 지속 시간을
+  게이팅하지 않는다.
+
+자원 선택은 요청 타입에 기반한다:
+
+```python
+dma_res = self._dma_write if isinstance(request, MemoryWriteMsg) else self._dma_read
+```
+
+### D4. 비종단 홉에서의 transit 포워딩
+
+`txn.next_hop`이 None이 아닐 때 — 전형적으로 역방향 응답 경로(PE →
+M_CPU → IO_CPU)에서 — 워커는 정상적으로 전달한다:
+
+```python
+if next_hop:
+    yield self.out_ports[next_hop].put(txn.advance())
+```
+
+팬아웃 분기는 종단 홉에서만 발화한다. 따라서 동일한 컴포넌트가 정방향
+명령 디스패치 역할과 역방향 응답 중계 역할을 모두 수행한다.
+
+### D5. DMA 팬아웃 (`_dma_fanout` — Memory R/W)
+
+종단 홉에서 각 Memory R/W 요청에 대해:
+
+1. `_resolve_dma_destinations(request)`가 요청의 PA로부터
+   `ctx.resolver.resolve(PhysAddr.decode(pa))`를 통해 도출된 PE별
+   `hbm_ctrl.pe{X}`를 반환한다(ADR-0017 D9).
+2. 각 목적지에 대해:
+   - `with dma_res.request() as req`를 통해 적절한 DMA 자원(`_dma_write`
+     또는 `_dma_read`)을 획득.
+   - `ctx.router.find_mcpu_dma_path()`로 경로를 해석.
+   - `drain_ns = ctx.compute_drain_ns(path, nbytes)`를 계산.
+   - `drain_ns`를 운반하는 서브 Transaction을 생성하여 `path[1]`로
+     디스패치.
+3. 목적지들에 걸쳐 `max_drain_ns`를 추적하고, 모든 응답 도착 후
+   `txn.result_data["xfer_ns"]`로 기록한다.
+4. PE별 응답이 모두 수집된 후(D8), IO_CPU로 되돌아가는 역방향 명령
+   경로로 집계 ResponseMsg를 전송한다.
+
+PA 디코드 폴백(`f"{cube_prefix}.hbm_ctrl"`)은 레거시 데드 코드이다 —
+ADR-0017 D4의 PE별 파티셔닝 이후로 그러한 노드는 존재하지 않는다.
+방어적으로 남겨두었으나 실제 목적지로 라우팅되지는 않는다.
+
+### D6. Kernel launch 팬아웃 (`_kernel_launch_fanout`)
+
+종단 홉에서 `KernelLaunchMsg`에 대해:
+
+1. `_resolve_pe_ids(target_pe)` → 본 큐브 내 PE id 리스트.
+2. 각 PE에 대해: `ctx.router.find_node_path()`를 통해
+   `f"{cube_prefix}.pe{pe_id}.pe_cpu"`로의 경로를 찾음.
+3. **`target_start_ns` 처리**(ADR-0009 D5):
+   - 요청에 이미 `target_start_ns`가 실려 있으면(IO_CPU가
+     ADR-0036 D3에 따라 스탬프함): **변경 없이 통과**.
+   - 없으면(단위 테스트에서의 직접 M_CPU 런치):
+     `env.now + max(PE별 leg 레이턴시)`로 큐브별 배리어를 계산하고
+     `dataclasses.replace`로 스탬프.
+4. `nbytes=0`인 서브 Transaction으로 디스패치(커널 런치는 제어 메시지;
+   nbytes=0 유지는 팬아웃을 공유 first-hop 패브릭 BW에서 떼어내며,
+   ADR-0036 D4를 미러링).
+5. PE별 응답이 모두 도착한 후(D8), 각 서브 Transaction의 `result_data`로부터
+   PE별 메트릭을 부모 트랜잭션으로 집계한다:
+
+   ```python
+   txn.result_data["pe_exec_ns"]  = max(existing, max(pe_exec_values))
+   txn.result_data["dma_ns"]      = max(existing, max(dma_values))
+   txn.result_data["compute_ns"]  = max(existing, max(compute_values))
+   ```
+
+   기존 값과의 max 병합이 중요한 이유는 크로스 큐브 IO_CPU 팬아웃이
+   동일한 부모 `result_data`를 공유하기 때문이다; 병합을 통해 한 큐브가
+   다른 큐브의 메트릭을 덮어쓰는 일을 방지한다.
+6. IO_CPU로 되돌아가는 역방향 경로로 집계 ResponseMsg를 전송.
+
+### D7. MMU map/unmap 팬아웃 (`_mmu_msg_fanout`)
+
+종단 홉에서 `MmuMapMsg` / `MmuUnmapMsg`에 대해:
+
+1. `_resolve_pe_ids(target_pe)` → PE id들.
+2. 각 PE에 대해: `find_node_path()`를 통해
+   `f"{cube_prefix}.pe{pe_id}.pe_mmu"`로의 경로를 찾음.
+3. `nbytes=0`인 서브 Transaction으로 디스패치.
+4. PE_MMU는 종단 노드이다 — ResponseMsg를 되돌려 보내지 **않는다**.
+   대신 서브 Transaction 자체의 `sub_done` 이벤트가 완료 시그널 역할을
+   한다.
+5. 모든 `sub_done` 이벤트를 인라인으로 기다림(`_pending` 카운터를 사용
+   **하지 않음** — D8은 응답을 동반하는 팬아웃 전용).
+6. IO_CPU로 되돌아가는 역방향 경로로 집계 ResponseMsg를 전송.
+
+### D8. 응답 집계 (`_pending` + `_parent_txns`)
+
+DMA 및 kernel-launch 팬아웃(역방향 경로로 도착하는 PE별 ResponseMsg를
+예상함)에 대해:
+
+```python
+self._pending: dict[str, tuple[int, int, simpy.Event]] = {}
+self._parent_txns: dict[str, Any] = {}
+```
+
+- 디스패치 시: `(expected, received=0, all_done)`을 등록하고 부모
+  트랜잭션을 기억.
+- `_worker`는 `is_response=True`로 응답을 인식하여 `_collect_response`로
+  라우팅하며, `_collect_response`는 `received`를 증가시키고 `received >=
+  expected`일 때 `all_done`을 시그널한다.
+- `yield all_done` 후, 팬아웃 경로는 집계 ResponseMsg를 구성한다:
+
+  ```python
+  resp_msg = ResponseMsg(
+      correlation_id=request.correlation_id,
+      request_id=request.request_id,
+      src_cube=cube_id,
+      src_pe=-1,             # -1 = M_CPU 집계, 단일 PE가 아님
+      success=True,          # 실패 의미는 구현되어 있지 않음
+  )
+  ```
+
+- 응답 Transaction은 `list(reversed(txn.path))`를 따라 IO_CPU로
+  되돌아간다.
+
+MMU 팬아웃(D7)은 PE_MMU가 종단이므로 더 단순한 `sub_done` 이벤트의
+인라인 리스트를 사용한다 — 가로챌 ResponseMsg 경로가 없다.
+
+### D9. 헬퍼와 설정 가능한 속성
+
+`_resolve_pe_ids(target_pe)`:
+
+- `int` → `[target_pe]`
+- `tuple[int, ...]` → `list(target_pe)`
+- `"all"` → `range(n_slices)`, 여기서 `n_slices`는 큐브
+  `memory_map.hbm_slices_per_cube`(기본 8)에서 가져온다.
+
+Kernel-launch 및 MMU 팬아웃 경로에서 사용된다.
+
+인스턴스별 레이턴시를 결정하는 단일 설정 가능 속성:
+
+| 사이트 | impl 이름 | overhead_ns |
+| --- | --- | --- |
+| 큐브 `m_cpu` | `builtin.m_cpu` | 5.0 |
+
+Transaction마다 `run()`에서 한 번 적용 — M_CPU에서의 명령 해석 및
+디스패치 결정 시간을 모델링한다.
+
+## Consequences
+
+### Positive
+
+- 세 가지 팬아웃 경로가 요청 타입에 의해 명확히 분리됨 — 새로운 요청
+  종류 추가는 isinstance 분기 한 줄과 팬아웃 메서드 하나로 가능.
+- M_CPU.DMA 채널은 독립적이며(읽기/쓰기가 동시 실행됨) capacity=1에서
+  디스패치 단계만 직렬화된다.
+- Transit 대 종단 동작이 단일 `if next_hop` 체크이므로, 동일한 컴포넌트가
+  역할 중복 없이 정방향 디스패치와 역방향 응답 중계를 처리한다.
+- `target_start_ns` 통과(D6)는 IO_CPU가 수립한 크로스 큐브 배리어
+  (ADR-0036 D3)를 보존하며, 폴백 계산은 직접 M_CPU 단위 테스트가 계속
+  동작하도록 한다.
+- 부모 `result_data`의 기존 값에 대한 PE별 메트릭의 `max` 병합은 동일한
+  부모를 공유하는 크로스 큐브 IO_CPU 팬아웃에 견고하다.
+
+### Negative
+
+- 부분 실패 의미가 없음 — 누락된 PE별 응답은 부모 `all_done`을 무기한
+  스톨시킨다. 시뮬레이션 용도로는 수용 가능하나 프로덕션 스타일의
+  엔드포인트로는 적합하지 않다.
+- `_resolve_dma_destinations`의 큐브 전역 hbm_ctrl 폴백은 데드 코드이다
+  (ADR-0017 D4 이후 그런 노드는 존재하지 않음). 방어적으로 남겨두었으나
+  혼동을 유발하므로 후속 정리가 권장된다.
+- DMA 자원 직렬화는 디스패치에만 적용된다(언바운드 store에서 `put`
+  호출은 즉시적). capacity=1 채널은 "본 M_CPU에서 동시에 in-flight인
+  요청은 하나"를 모델링하며 "전송 지속 시간 직렬화"를 모델링하지 않는다
+  — 실제 전송 병렬성은 와이어 프로세스(ADR-0015 D2)와 `drain_ns`를
+  참조해야 한다.
+
+## Links
+
+- ADR-0009 D3 (M_CPU 팬아웃 및 집계 완료 의미)
+- ADR-0009 D5 (`target_start_ns` — 존재 시 변경 없이 통과; 부재 시
+  큐브별 배리어로 계산)
+- ADR-0011 D-VA3 (MmuMapMsg 패브릭 경로에 M_CPU가 PE 팬아웃 지점으로
+  포함됨)
+- ADR-0014 D4 (DMA 엔진 capacity=1; M_CPU.DMA가 큐브 수준에서 동일한
+  계약을 미러링)
+- ADR-0015 D5 (M_CPU.DMA는 M_CPU의 내부 서브 컴포넌트이며 토폴로지
+  노드가 아님)
+- ADR-0017 D9 (AddressResolver가 PE별 `hbm_ctrl.pe{X}`를 반환)
+- ADR-0036 D3 / D4 (IO_CPU가 `target_start_ns`를 스탬프; M_CPU는 변경
+  없이 통과; 팬아웃 전반에서 nbytes=0 불변식 보존)

docs/adr-ko/ADR-0036-dev-io-cpu-component-model.md

@@ -0,0 +1,205 @@
+# ADR-0036: IO_CPU 컴포넌트 모델
+
+## Status
+
+Accepted
+
+## Context
+
+IO_CPU는 시뮬레이션 그래프 내부의 IO 칩렛 호스트 대향 엔드포인트이다.
+PCIE_EP는 런타임 API로부터 호스트 메시지를 수신하여 io_noc를 통해
+라우팅한다; 명령을 동반하는 요청(KernelLaunch, MmuMap/Unmap)의 경우
+io_noc는 IO_CPU로 전달하며, IO_CPU는 다음을 수행한다:
+
+- 요청을 큐브별 M_CPU로 팬아웃.
+- 큐브별 응답을 단일 호스트 가시 완료로 집계.
+- 커널 런치의 경우, 타깃이 된 모든 큐브의 모든 PE가 동일한 시뮬레이션
+  시각에 커널 본체 실행을 시작하도록 전역 `target_start_ns` 배리어를
+  스탬프함(ADR-0009 D5).
+
+Memory R/W 트래픽은 ADR-0015 D4 / ADR-0016 D3에 따라 IO_CPU를 우회한다;
+따라서 본 컴포넌트는 정상 동작에서 명령 평면 트래픽만을 처리한다.
+
+본 ADR은 위의 책임을 실현하는 IO_CPU 컴포넌트 구현을 문서화한다.
+
+## Decision
+
+### D1. 역할
+
+IO_CPU는 IO 칩렛의 호스트 대향 엔드포인트이다. 두 가지 주요 책임을
+갖는다:
+
+1. **멀티 큐브 팬아웃** — KernelLaunchMsg / MmuMapMsg / MmuUnmapMsg를
+   큐브별 M_CPU로 분배.
+2. **응답 집계** — 큐브별 ResponseMsg를 수집하고, 타깃이 된 모든 큐브가
+   응답한 후 부모 `txn.done`을 시그널.
+
+세 번째이자 더 좁은 책임은 KernelLaunchMsg에만 적용된다:
+**`target_start_ns` 전역 배리어 스탬핑**(D3).
+
+본 컴포넌트는 다음을 하지 **않는다**:
+
+- 라우팅 결정 — 경로는 라우터에 의해 사전 계산된다(ADR-0002).
+- 텐서 또는 커널 내부 디코드 — 그러한 관심사는 M_CPU / PE_CPU / 엔진에
+  속한다.
+- PE 수준 팬아웃 처리 — M_CPU가 큐브 내에서 팬아웃한다(ADR-0009 D3).
+- Memory R/W 데이터 경로 처리 — ADR-0015 D4와 ADR-0016 D3에 따라
+  IO_CPU를 우회한다(`_resolve_cube_targets` 내의 Memory R/W 해석 코드는
+  방어적 폴백으로만 존재).
+
+호출당(`run()`): 들어오는 Transaction마다 설정된 `overhead_ns`를 한 번
+적용한다(D8).
+
+### D2. 정방향 경로 — 멀티 큐브 팬아웃
+
+응답이 아닌 Transaction이 도착하면, 워커는:
+
+1. `run()`을 통해 `overhead_ns`를 지불.
+2. `_resolve_cube_targets`를 호출하여 요청으로부터 `(sip, cube)` 타깃
+   리스트를 도출(D5).
+3. 각 타깃에 대해:
+   - `ctx.resolver.find_m_cpu(sip, cube)`를 통해 M_CPU 노드 id를 해석.
+   - `ctx.router.find_node_path(io_cpu, m_cpu)`를 통해 경로를 해석.
+   - `path`가 채워진 큐브별 서브 Transaction을 생성하여 `path[1]`
+     (io_noc의 첫 홉)으로 전달.
+4. 집계 상태 등록: `_pending[request_id] = (expected, received=0,
+   parent_done)`.
+
+### D3. KernelLaunch `target_start_ns` 전역 배리어 (ADR-0009 D5)
+
+IO_CPU는 `target_start_ns`의 정규 스탬퍼이다. 요청이
+`KernelLaunchMsg`일 때, IO_CPU는 타깃이 된 모든 큐브의 모든 PE를 포괄하는
+단일 전역 배리어를 계산한다:
+
+```text
+for (sip, cube) in cube_targets:
+    leg1 = compute_path_latency_ns(io_cpu → m_cpu(sip, cube), nbytes=0)
+    for pe_id in target_pe_ids:
+        leg2 = compute_path_latency_ns(m_cpu → pe_cpu(sip, cube, pe_id),
+                                       nbytes=0)
+        latency = leg1 + leg2 - io_overhead_ns - m_overhead_ns
+        global_max = max(global_max, latency)
+
+target_start_ns = env.now + global_max
+```
+
+이후 요청은 (`dataclasses.replace`를 통해) 교체되어 스탬프된 값이 팬아웃
+전반에 전파된다.
+
+두 가지 오버헤드 보정:
+
+- `io_overhead_ns`는 차감되는데, IO_CPU가 본 메서드 실행 전에 `run()`에서
+  이미 지불했기 때문이다.
+- `m_overhead_ns`는 한 번 차감되는데, 경로 레이턴시에서 leg1의 종단점인
+  동시에 leg2의 시작점으로 두 번 등장하지만 M_CPU는 런타임에 단 한 번만
+  지불하기 때문이다.
+
+모든 다운스트림 PE_CPU는 커널 본체 실행을 시작하기 전 `target_start_ns`
+까지 yield한다; 이를 통해 개별 디스패치 경로가 얼마나 오래 걸렸는지와
+무관하게 모든 PE가 동일한 시뮬레이션 시각에 시작한다.
+
+### D4. KernelLaunch 서브 Transaction은 `nbytes=0`을 운반
+
+KernelLaunchMsg의 큐브별 서브 Transaction은 부모 `txn.nbytes`를 무시하고
+`nbytes=0`을 강제한다:
+
+- 커널 런치는 제어 메시지이다; 데이터 패브릭 수준에서 페이로드 크기는
+  무관하다.
+- `nbytes > 0`이면 모든 큐브별 서브 트랜잭션이 io_noc의 공유 first-hop
+  패브릭 BW를 점유한다. 16개 큐브에서는 이로 인해 팬아웃이 직렬화되어
+  먼 M_CPU들이 `target_start_ns`를 지나치게 되고 D3 불변식이 깨진다.
+
+KernelLaunch가 아닌 서브 Transaction은 `txn.nbytes`를 보존한다(실제
+페이로드 크기를 운반하는 방어적 Memory R/W 폴백 경로에만 관련됨).
+
+### D5. 요청 타입별 큐브 타깃 해석
+
+`_resolve_cube_targets`는 요청 타입에 따라 디스패치한다:
+
+| 요청 타입 | `(sip, cube)`의 출처 | `target_cubes="all"` 의미 |
+| --- | --- | --- |
+| `MemoryWriteMsg` | `dst_sip`, `dst_cube` (또는 `PhysAddr.decode(dst_pa).die_id` 폴백) | PA 디코드로 도출되는 단일 큐브 |
+| `MemoryReadMsg` | `src_sip`, `src_cube` (또는 `PhysAddr.decode(src_pa).die_id` 폴백) | PA 디코드로 도출되는 단일 큐브 |
+| `KernelLaunchMsg` | `shard.sip == my_sip`으로 필터링된 텐서 샤드 | 이 SIP 위에서 샤드를 소유하는 모든 큐브 |
+| `MmuMapMsg` / `MmuUnmapMsg` | 본 SIP로 필터링된 `target_cubes` 리스트 | 스펙으로부터 `range(cubes_per_sip)` |
+
+각 IO_CPU 인스턴스는 자기 SIP 내에서만 팬아웃한다 — `_my_sip()`이
+노드 id에서 SIP id를 파싱한다(예: `sip0.io0.io_cpu` → 0).
+
+Memory R/W 행은 방어적 완전성을 위해 존재한다; 엔진의 정상 경로는
+Memory R/W를 `_process_memory_direct()` / `find_memory_path()`로
+라우팅하여 IO_CPU를 완전히 우회한다(ADR-0015 D4 / ADR-0016 D3).
+
+### D6. 응답 집계
+
+`_pending: dict[request_id → (expected, received, parent_done)]`:
+
+- 디스패치 시: `(len(cube_targets), 0, txn.done)`을 등록.
+- `_worker`는 `is_response=True`로 응답을 인식하여 `_collect_response`로
+  라우팅한다.
+- `_collect_response`는 `received`를 증가시키며, `received >= expected`가
+  되면 `parent_done.succeed()`를 호출하고 엔트리를 `_pending`에서
+  제거한다.
+
+이는 단순한 요청별 카운터이다. 큐브별 정체성 추적이나 부분 실패 처리는
+없다 — 누락된 응답은 부모 done을 무기한 스톨시킨다. 프로덕션 스타일의
+실패 경로는 현재 시뮬레이터 모델의 범위 밖이다.
+
+### D7. `target_pe` 해석 헬퍼
+
+`_resolve_pe_ids(target_pe)`:
+
+- `int` → `[target_pe]`.
+- `tuple[int, ...]` → `list(target_pe)`.
+- `"all"` → `range(n_slices)`, 여기서 `n_slices`는 큐브
+  `memory_map.hbm_slices_per_cube`(기본 8)에서 가져온다.
+
+D3의 배리어 계산에서 큐브별로 모든 PE 타깃을 열거하는 데 사용된다.
+
+### D8. 설정 가능한 `overhead_ns`
+
+단일 속성이 인스턴스별 레이턴시를 결정한다:
+
+| 사이트 | impl 이름 | overhead_ns |
+| --- | --- | --- |
+| IO 칩렛 `io_cpu` | `builtin.io_cpu` | 10.0 |
+
+Transaction마다 `run()`에서 한 번 적용된다. IO_CPU에서의 명령 해석 및
+디스패치 결정 시간을 모델링한다.
+
+## Consequences
+
+### Positive
+
+- 크로스 큐브 및 크로스 SIP 커널 런치가 단일 전역 배리어를 공유한다
+  (D3 + D4) — 시작 시각의 큐브별 분기가 없다.
+- nbytes=0 불변식이 팬아웃을 공유 first-hop 패브릭 BW로부터 떼어내,
+  대규모(16 큐브)에서도 배리어의 정확도를 보존한다.
+- 단일 카운터를 통한 응답 집계 → 최소 상태, 결정론적 완료 순서.
+- SIP별 스코핑(`_my_sip()`)이 서로 다른 SIP의 IO_CPU들을 깨끗이
+  독립시킨다.
+
+### Negative
+
+- 부분 실패 의미가 없음 — 누락된 큐브별 응답은 부모를 무기한
+  스톨시킨다. 시뮬레이션 용도로는 충분하나 프로덕션 스타일의
+  엔드포인트로는 적합하지 않다.
+- `_pending`은 일반 dict이다; in-flight 요청이 상태로 누적된다. 현재
+  벤치마크 워크로드(미해결 런치가 적음)에는 허용 가능하나, 원리적으로는
+  무한하다.
+- `_resolve_cube_targets`의 Memory R/W 해석 분기는 정상 엔진 경로에서
+  데드 코드이다. 방어적으로 남겨두었으나 우회 경로가 변경되면 드리프트
+  위험을 초래한다.
+
+## Links
+
+- ADR-0002 (라우팅 거리 — 경로 계산)
+- ADR-0009 D1 (커널 런치는 IO_CPU에 대한 엔드포인트 요청)
+- ADR-0009 D3 (M_CPU는 큐브 내에서 팬아웃; IO_CPU는 큐브 사이에서 팬아웃)
+- ADR-0009 D5 (IO_CPU에서의 target_start_ns 정규 스탬핑)
+- ADR-0011 D-VA3 (MmuMapMsg가 큐브 팬아웃을 위해 IO_CPU를 경유)
+- ADR-0012 (호스트 ↔ IO_CPU 메시지 스키마)
+- ADR-0015 D4 (Memory R/W는 IO_CPU 우회; 커널 런치는 IO_CPU 경유)
+- ADR-0016 D1 (IO 칩렛 io_noc — IO_CPU가 여기 부착됨)
+- ADR-0016 D3 (Memory R/W 경로가 IO_CPU 우회)
+- ADR-0016 D4 (명령 해석을 위한 IO_CPU 경유 커널 런치 경로)

docs/adr-ko/ADR-0037-dev-forwarding-component.md

@@ -0,0 +1,185 @@
+# ADR-0037: Forwarding 컴포넌트 (forwarding_v1)
+
+## Status
+
+Accepted
+
+## Context
+
+시뮬레이션 그래프에는 순전히 패브릭 통과를 모델링하기 위해 존재하는 노드
+위치들이 많다 — NOC 메시 라우터, 스위치, UCIe 프로토콜 엔드포인트, IO
+칩렛 io_noc, transit 큐브. 이들은 공통 패턴을 공유한다: 메시지를 수신하고,
+컴포넌트별 오버헤드(헤더 디코드 + 라우팅 결정 시간을 모델링)를 적용하며,
+사전 계산된 경로를 따라 다음 홉으로 전달한다.
+
+본 ADR은 이러한 transit 노드에 대한 계약을 정의한다: 웜홀 cut-through
+의미로 flit 인지 포워딩을 처리하는 단일 컴포넌트 타입(`TransitComponent`)이며,
+각 인스턴스가 수행하는 개념적 역할에 따라 여러 impl 이름 아래에 사용된다.
+
+## Decision
+
+### D1. 역할
+
+Forwarding 컴포넌트(`TransitComponent` 클래스)는 시뮬레이션 그래프의
+**상태 없는 transit 노드**이다. 메시지가 물리적으로 통과하지만 의미론적
+처리는 일어나지 않는 모든 패브릭 위치를 모델링한다.
+
+통과당 컴포넌트는:
+
+1. `in_port`에서 들어오는 Transaction 또는 Flit을 읽는다.
+2. 설정된 컴포넌트별 오버헤드(`overhead_ns`)를 적용한다. 멀티 flit
+   페이로드라도 **Transaction당 한 번** 적용된다(D2 참조).
+3. Transaction의 사전 계산된 `path`를 따라 다음 홉을 조회한다.
+4. 해당 `out_port`로 전달한다; 종단 노드(다음 홉 없음)에서는 `is_last`
+   flit이 도착하면 `txn.done`을 시그널한다.
+
+본 컴포넌트는 다음을 하지 **않는다**:
+
+- 라우팅 결정 — 경로는 라우터에 의해 사전 계산된다(ADR-0002 /
+  ADR-0017 D2). Forwarding은 홉별 단계만 실행한다.
+- 와이어 전파나 대역폭 점유 모델링 — 컴포넌트 사이의 별도 와이어
+  프로세스가 처리한다(ADR-0015 D2).
+- 주소 해석 — AddressResolver가 담당한다(ADR-0017 D9).
+- 완료 집계 — 종단 엔드포인트(IO_CPU, M_CPU, HBM_CTRL)가 담당한다.
+
+### D2. First-flit 오버헤드 모델 (헤더 디코드)
+
+Transaction별 `overhead_ns`는 첫 flit 도착 시 **정확히 한 번** 적용된다:
+
+- `_txn_decoded: set[int]`이 본 노드에서 이미 오버헤드를 지불한
+  Transaction들을 추적한다.
+- 어떤 Transaction의 첫 flit 도착 시: `yield self.run(env, msg.txn.nbytes)`
+  — 오버헤드를 지불한다.
+- 동일 Transaction의 후속 flit들은 오버헤드를 건너뛰고 추가 지연 없이
+  파이프라인 통과한다.
+- `is_last` flit 시: Transaction을 `_txn_decoded`에서 제거한다.
+
+이는 실제 HW의 동작 — 헤더 디코드와 라우팅 결정이 첫 flit에서 한 번
+일어나고, 이후 페이로드 flit들은 같은 경로로 스트리밍되는(웜홀
+cut-through) — 을 모델링한다. 멀티 홉 파이프라이닝은 자연스럽게
+발현된다 — 각 홉이 자신의 first-flit 오버헤드를 추가하지만, 첫 flit
+이후의 flit들은 이미 첫 flit이 통과한 어떤 홉에서도 오버헤드를 다시
+지불하지 않는다.
+
+### D3. 직렬 워커 포워딩 (순서 보존)
+
+본 컴포넌트의 워커는 `_inbox`에서 flit을 소비하여 도착 순서대로 직렬
+포워딩하는 단일 SimPy 프로세스이다. 컴포넌트는 flit마다
+`env.process(...)`를 spawn하지 **않는다**.
+
+근거: 첫 flit이 `overhead_ns`에서 yield하는 동안 후속 flit이 병렬
+프로세스에서 실행되면, 후속 flit이 첫 flit을 추월할 수 있다. 이는 순서가
+어긋난 전달을 낳고, `is_last` flit이 첫 flit보다 먼저 목적지에 도착하게
+하여 — 트랜잭션의 완료 의미와 다운스트림의 flit 인덱스 기반 처리 모두를
+손상시킨다.
+
+### D4. 경로 기반 next-hop 라우팅
+
+라우팅은 Forwarding 컴포넌트의 관심사가 **아니다**. Transaction은 라우터에
+의해 사전 계산된 `path`(ADR-0002 / ADR-0017 D2)와 함께 도착한다.
+컴포넌트는 단지 자신의 경로상 위치를 찾아 `path[index + 1]`로 전달한다:
+
+```python
+def _next_hop_in_path(self, txn):
+    my_id = self.node.id
+    path = txn.path
+    for i, n in enumerate(path):
+        if n == my_id and i + 1 < len(path):
+            return path[i + 1]
+    return None
+```
+
+`next_hop`이 발견되고 `out_ports`에 존재하면 flit이 전달된다. 그렇지
+않으면(종단 노드) `is_last` flit이 도착할 때 `txn.done.succeed()`가
+호출된다.
+
+### D5. Flit 인지 모드와 Non-Flit 폴백
+
+`_FLIT_AWARE = True`는 본 컴포넌트가 베이스 클래스의 `_fan_in` 내 flit
+재조립 로직에서 제외되도록 한다. Flit은 재조립 없이 `_inbox`에 직접
+놓이며, 이는 워커 루프(D2, D3)에서의 per-flit 처리를 가능케 한다.
+
+Non-Flit 메시지 — 0바이트 제어 Transaction이나 그 외 청크화되지 않는
+페이로드 — 는 `env.process`를 통해 베이스 클래스의 레거시 `_forward_txn`
+경로로 빠진다. 이는 flit 수준 처리의 이득이 없는 제어 평면 트래픽에
+대한 하위 호환성을 보존한다.
+
+### D6. 베이스 클래스에서의 멀티 스트림 병합
+
+라우터에서의 멀티 스트림 FIFO 병합은 Forwarding이 아닌 베이스 클래스의
+책임이다. 베이스 클래스의 `_fan_in`은 `in_port`마다 하나의 프로세스를
+spawn한다; 모두가 공유된 단일 `_inbox`에 push한다. 따라서 서로 다른
+업스트림 스트림의 flit들은 `_inbox`의 FIFO 순서로 flit 단위에서
+인터리브된다.
+
+Forwarding 워커는 단지 `_inbox`를 도착 순서대로 소비할 뿐이다 —
+공유 inbox 위의 공정 FIFO로 라우터별 멀티 플로우 중재를 올바르게
+모델링한다.
+
+### D7. 여러 impl 이름 아래의 단일 구현
+
+단일 `TransitComponent` 클래스가 `components.yaml`에서 네 가지 impl
+이름으로 등록된다:
+
+- `builtin.forwarding` — 범용 forwarding (예: `io_noc`, `noc_router`,
+  UCIe conn 브리지)
+- `builtin.switch` — 트레이 수준 스위치
+- `builtin.noc` — 큐브 수준 NOC 패브릭(레거시 싱글톤; 현재 NOC
+  라우터는 `builtin.forwarding`을 사용)
+- `builtin.ucie` — UCIe 프로토콜 엔드포인트
+
+네 별칭 모두 동일한 동작을 갖는 동일한 클래스를 인스턴스화한다.
+인스턴스별 차별화는 `attrs.overhead_ns`에만 존재한다. 별도 impl 이름이
+존재하는 것은 가독성을 위한 의도 태그이자, 하위 호환을 깨지 않고 향후
+분기를 허용하기 위함이다.
+
+### D8. 설정 가능한 `overhead_ns`
+
+단일 속성이 인스턴스별 레이턴시를 결정한다:
+
+| 사용 사이트 | impl 이름 | overhead_ns |
+| --- | --- | --- |
+| 트레이 수준 스위치 | `builtin.switch` | 5.0 |
+| 큐브 NOC 라우터 | `builtin.forwarding` | 2.0 |
+| IO 칩렛 io_noc | `builtin.forwarding` | 0.0 |
+| UCIe 프로토콜 엔드포인트(`ucie-{N,S,E,W}`) | `builtin.ucie` | 8.0 |
+| UCIe conn 브리지(`ucie-{PORT}.conn{N}`) | `builtin.forwarding` | 0.0 |
+
+기본값은 0.0이다. 속성은 매 `run()` 호출에서 읽히므로 동적 재설정이
+가능하나 현재는 사용되지 않는다.
+
+## Consequences
+
+### Positive
+
+- 단일 클래스가 시뮬레이션 그래프의 모든 transit 노드 역할을 처리한다
+  — 개체 수가 많은 컴포넌트 타입에 대한 최소 코드 표면.
+- Flit 인지 처리 + 직렬 워커는 per-flit 프로세스 오버헤드 없이 멀티 홉
+  경로 전반에 걸쳐 웜홀 의미를 보존한다.
+- `overhead_ns`만이 유일한 인스턴스별 튜너블이다; 라우팅, 대역폭, 주소
+  해석은 자체 컴포넌트/모듈에서 깨끗이 분리되어 있다.
+- 멀티 스트림 병합이 베이스 클래스 구조에서 자연스럽게 발현된다; 라우터
+  전용 로직이 공정 FIFO 중재를 중복 구현하지 않는다.
+- Non-Flit 폴백 경로는 모든 메시지를 flit 프레임워크로 강제하지 않고도
+  제어 평면 트래픽이 계속 동작하도록 한다.
+
+### Negative
+
+- 단일 클래스가 사용 사이트의 의도를 `attrs.overhead_ns` 설정 안에
+  숨긴다; 어떤 impl 이름이 어떤 동작 클래스로 매핑되는지 보려면 독자가
+  `topology.yaml` + `components.yaml`을 참조해야 한다.
+- per-flit 직렬 워커는 `overhead_ns`가 크고 같은 라우터에 다수의 동시
+  트랜잭션이 도착할 때 병목이 된다; 현재 값(0–8 ns)에서는 무시할 만한
+  수준이다.
+
+## Links
+
+- ADR-0002 (라우팅 거리 — 경로 계산)
+- ADR-0015 D1 (컴포넌트 포트 모델)
+- ADR-0015 D2 (와이어 프로세스 — 본 컴포넌트와 별개의 BW + 전파)
+- ADR-0015 D6 (Transit 큐브 forwarding 패턴)
+- ADR-0016 D1 (IO 칩렛 io_noc — 본 컴포넌트 사용)
+- ADR-0017 D1 (큐브 NOC 라우터 — 본 컴포넌트 사용)
+- ADR-0017 D6 (UCIe 분해 — `ucie-{PORT}` 인스턴스가 본 컴포넌트 사용)
+- ADR-0033 D1 (Flit 인지 통과, first-flit 오버헤드, 멀티 스트림 병합
+  의미)

docs/adr-ko/ADR-0038-dev-pcie-ep-component-model.md

@@ -0,0 +1,133 @@
+# ADR-0038: PCIE_EP Component Model
+
+## Status
+
+Accepted (2026-05-20).
+
+ADR-0035 (M_CPU), ADR-0036 (IO_CPU), ADR-0037 (Forwarding)
+와 같은 결의 컴포넌트-레벨 ADR.
+
+## First action (제일 처음에 하는 일)
+
+`_inbox`에서 Transaction을 한 건 꺼내 `_forward_txn`을 통해 `run()`을 호출하고,
+그 안에서 `node.attrs["overhead_ns"]` 만큼 `env.timeout()`으로 PCIe 프로토콜
+처리 지연을 적용한다. 그 이후 시점부터는 일반 `ComponentBase` 워커가 정의한
+forwarding 규약을 따른다 (다음 hop이 있으면 `out_ports[next_hop].put(...)`,
+아니면 `drain_ns`를 소비하고 `txn.done.succeed()`).
+
+즉, **PCIE_EP의 첫 번째 일은 "PCIe 프로토콜 오버헤드를 시간으로 표현하는 것"**
+하나뿐이고, 라우팅·페이로드 변환·MMIO 디코딩 같은 부가 의사결정은 하지 않는다.
+
+## Context
+
+PCIE_EP는 토폴로지 그래프에서 **호스트와 디바이스 사이의 단방향 경계 포인트**
+역할을 한다. 빌더 (`topology/builder.py`)는 SIP마다 IO chiplet 인스턴스를
+생성하고 그 안에 `pcie_ep`, `io_cpu`, `io_noc`을 둔 뒤, 외부 호스트 측의 cross-SIP
+switch와 `pcie_ep` 사이에 양방향 엣지를 깐다:
+
+- `switch → pcie_ep`: host → device 트래픽 (MemoryWrite, MemoryRead, KernelLaunch).
+- `pcie_ep → switch`: device-side outbound (예: cross-SIP IPCQ 토큰).
+
+IOChiplet 내부적으로는 `pcie_ep ↔ io_noc` 양방향 엣지가 깔리고, 그 다음 hop이
+`io_cpu`나 cube 측 hbm_ctrl 경로로 분기된다 (ADR-0036 IO_CPU 모델 참고).
+라우터·리졸버는 SPEC R7이 요구하는 "PCIE_EP는 메모리 오퍼레이션을 위한
+엔드포인트"라는 계약을 이미 인지하고 있어, `find_pcie_ep(sip)`,
+`find_memory_path(pcie_ep, dst_node)` 같은 helper가 PCIE_EP를 시작점으로 한다.
+
+문제는 이 모든 의존 관계가 builder/router/resolver 쪽에는 있으나, **PCIE_EP
+자신의 내부 모델을 명시하는 ADR이 없다**는 것이다. 결과적으로:
+
+- "PCIE_EP는 어떤 latency를 모델링하나?"가 코드를 읽어야만 답이 나온다.
+- 다른 컴포넌트(IO_CPU=ADR-0036, M_CPU=ADR-0035)와의 비대칭이 발생한다.
+- 향후 PCIe link-layer 모델(예: TLP credit, retry)을 더 정교하게 만들지에 대한
+  의사결정 근거가 흩어진다.
+
+이 ADR은 현재의 **얇은 (thin) PCIE_EP 모델**을 명시적으로 못 박고, 그것이
+의도된 단순화임을 기록한다 (ADR-0033 latency model 단순화 정책과 정렬).
+
+## Decision
+
+### D1. PCIE_EP는 ComponentBase의 일반 forwarding 워커를 그대로 사용한다
+
+`PcieEpComponent`는 `ComponentBase`를 상속하며 `_worker`/`_forward_txn`을
+오버라이드하지 않는다. 따라서 모든 Transaction은 다음 순서로 처리된다:
+
+1. `_fan_in`이 들어오는 메시지(또는 Flit reassembly된 Transaction)를 `_inbox`에
+   적재한다.
+2. `_worker`가 `_inbox`에서 하나 꺼내 `env.process(self._forward_txn(env, txn))`로
+   포크한다 (per-message 파이프라이닝).
+3. `_forward_txn`이 op_log 시작 hook → `run()` 지연 → op_log 종료 hook 순서로
+   호출한다.
+4. `run()`은 단 한 줄: `yield env.timeout(overhead_ns)`.
+5. 다음 hop이 있으면 `out_ports[next_hop].put(txn.advance())`, 없으면 (terminal로
+   도착한 경우) `drain_ns`를 소비 후 `txn.done.succeed()`.
+
+### D2. PCIE_EP의 유일한 시간 모델은 `overhead_ns`다
+
+`node.attrs["overhead_ns"]`만 latency 파라미터로 인정한다. 코드 기본값은
+`0.0`이며, `topology.yaml` 의 IOChiplet `components.pcie_ep.attrs` 가 실제 값을
+지정한다 (현재 토폴로지: `overhead_ns: 5.0` ns).
+
+별도의 BW 직렬화 자원(simpy.Resource), 큐 깊이, retry 모델은 두지 않는다.
+링크-레벨 BW 직렬화는 wire-side에서 처리된다 — IOChiplet 내부는
+`pcie_ep_to_noc_bw_gbs = 256.0 GB/s` 링크, 외부는 system의 `io_ep_to_switch`
+링크 BW가 적용된다 (ADR-0015 port/wire 모델). PCIE_EP 컴포넌트 자체는 이
+BW 회계에 관여하지 않는다.
+
+### D3. PCIE_EP는 양방향 사용을 인지하지만, 방향에 따라 동작을 바꾸지 않는다
+
+토폴로지 빌더가 `switch ↔ pcie_ep` 와 `pcie_ep ↔ io_noc` 양방향 엣지를 깐다.
+따라서 PCIE_EP는:
+
+- inbound (host→device): switch에서 도착한 Transaction을 io_noc 쪽으로 다음 hop
+  계산을 통해 forward.
+- outbound (device→host): io_noc/io_cpu에서 도착한 Transaction을 switch 쪽으로
+  forward.
+
+두 경우 모두 D1의 일반 forwarding 워커가 처리하며, 컴포넌트 코드 자체는 방향을
+구분하지 않는다 (`txn.next_hop`만 따른다).
+
+### D4. PCIE_EP는 Flit-aware가 아니다 (legacy reassembly 경로)
+
+`_FLIT_AWARE`를 `True`로 두지 않는다. 따라서 `_fan_in`이 상류에서 chunkify된
+Flit들을 부모 Transaction으로 재조립하여 `_inbox`에 넣는다 (ADR-0033 Phase 2c
+점진적 rollout 정책과 정렬).
+
+PCIE_EP가 PCIe TLP-level credit 모델을 갖도록 확장될 미래에 D4를 재평가한다.
+
+### D5. PCIE_EP는 라우팅 helper의 **명명된 노드**다
+
+`policy/routing/router.py`의 `find_pcie_ep(sip, io_id="io0")`,
+`find_all_pcie_eps()`, `find_memory_path(pcie_ep, dst_node)`는 PCIE_EP를 메모리
+경로의 시작점(또는 종점)으로 간주한다. 컴포넌트 본체는 이 helper에 어떤 정보도
+제공하지 않으며, 명명 규칙(`sip{S}.{io_id}.pcie_ep`)은 토폴로지 빌더가 보장한다.
+
+## Alternatives Considered
+
+### A1. PCIe TLP-level 모델 (credit, retry, MPS 분할)
+
+기각. ADR-0033이 명시한 "현재 latency 모델은 abstract overhead + BW 직렬화로
+표현"이라는 단순화 원칙에 어긋난다. 호스트↔디바이스 protocol 정합성은 SPEC §5
+"Non-Goals"에 의해 의도적으로 out-of-scope이다.
+
+### A2. PCIE_EP에 자체 simpy.Resource로 inflight 제한 두기
+
+기각. 현재 워크로드에서 호스트 트래픽은 컨텐션 병목이 아니다. 필요해지는 시점에
+별도 ADR로 도입한다 (호환성 측면에서 D1은 그대로 두고 D2를 확장하는 형태).
+
+### A3. PCIE_EP를 IO_CPU와 합치기
+
+기각. PCIE_EP는 host-side에서 처음 만나는 protocol boundary 노드이고, IO_CPU는
+디바이스-쪽 control-plane 처리 노드다 (ADR-0036). 트래픽 fan-out·command 디코딩
+같은 의사결정 비용은 IO_CPU에 모이며, PCIE_EP는 link-edge overhead만 표현하는
+것이 의미가 있다. 합치면 두 책임이 섞여 ADR-0007 (runtime API/sim_engine 경계)
+정신에 어긋난다.
+
+## Consequences
+
+- PCIE_EP는 코드 라인이 거의 0인 채로 명시적인 모델 ADR을 갖게 된다 — 일관성
+  ↑, 유지보수 비용 ↓.
+- 향후 PCIe-level 정밀화가 필요해지면 D2/D4를 확장하는 새 ADR을 만들어
+  supersede한다.
+- `find_memory_path` 등 router helper가 PCIE_EP를 명명된 노드로 의존한다는
+  사실이 D5에서 명시되므로, 컴포넌트 ID 명명 규칙 변경 시 영향 범위가 명확해진다.

docs/adr-ko/ADR-0039-dev-pe-mmu-component-model.md

@@ -0,0 +1,194 @@
+# ADR-0039: PE_MMU Component Model — 컴포넌트 + 유틸리티 이중 역할
+
+## Status
+
+Accepted (2026-05-20).
+
+ADR-0011 (PA/VA/LA address model) 의 VA 모델에서 "PE_MMU가 VA→PA 변환"이라고만
+선언되어 있는데, **PE_MMU 컴포넌트 자신의 동작 모델**을 별도로 못 박는 ADR.
+
+## First action (제일 처음에 하는 일)
+
+생성 시점에 `node.attrs["page_size"]` (default `2 MiB`) 와
+`node.attrs["tlb_overhead_ns"]` (default `0.0`) 를 읽어 내부 `PeMMU` 객체
+(`policy.address.pe_mmu.PeMMU`) 를 단 한 번 인스턴스화한다. 이 객체가 페이지
+테이블·서브페이지 region 리스트·TLB 오버헤드의 단일 보유자(single owner)이다.
+
+런타임에서의 첫 동작은 두 갈래로 갈린다:
+
+- **컴포넌트 경로 (inbox 소비)**: `_worker`가 `_inbox`에서 Transaction을 한 건
+  꺼내, 그 `request`가 `MmuMapMsg`이면 각 엔트리에 대해
+  `self._mmu.map(va, pa, size)`를 호출하고 `txn.done.succeed()`.
+  `MmuUnmapMsg`이면 `unmap(va, size)`, 그 외 타입이면 표준 `_forward_txn`으로
+  떨군다. 즉 **MMU의 첫 일은 "map/unmap 명령을 페이지 테이블에 반영하는 것"**.
+- **유틸리티 경로 (직접 호출)**: PE_DMA / PE_GEMM 같은 동일 PE 내부 엔진이
+  `pe_mmu.mmu.translate(va)`를 직접 호출한다. 이 경로에서는 SimPy 이벤트가
+  발생하지 않으며, 호출자가 (overhead_ns > 0인 경우) 본인 process에서
+  `yield env.timeout(mmu.overhead_ns)`를 처리한다.
+
+## Context
+
+ADR-0011은 PA/VA/LA 세 가지 주소 모델을 정의하고 "VA 모델 = PE_MMU를 통한 변환"
+이라고만 합의했다. 그러나 코드 상의 `PeMmuComponent`는 두 가지 상호 보완적인
+역할을 동시에 수행한다:
+
+1. **토폴로지 그래프 상의 컴포넌트**: cube NoC에서 `MmuMapMsg` / `MmuUnmapMsg`
+   sideband 메시지를 수신하여 페이지 테이블을 갱신한다.
+2. **PE-로컬 유틸리티 객체**: 동일 PE의 PE_DMA / PE_GEMM이 latency 0으로 (혹은
+   호출자 측에서 `overhead_ns`만 부담하면서) 직접 `translate(va)`를 호출한다.
+
+이 두 역할을 모두 다루는 ADR이 없어 다음 모호함이 발생한다:
+
+- "왜 MMU 변환에 SimPy 이벤트가 안 잡히나?" (실제로는 호출자 측에서 잡고 있음)
+- 서브페이지 region 모델은 무엇이고, 왜 그 모델인가? (코드 docstring에는 있으나
+  ADR이 없음 — `project_mmu_subpage_stopgap`라는 memory note 참조만 존재)
+- map/unmap 메시지가 **누구로부터** 와서 **언제까지** 갱신되어야 하는가
+  (ordering 계약)?
+
+또한 `PeMMU.map()` 은 "later append, last-write-wins (역방향 탐색)" 의미를 갖는데,
+이것은 단순한 단일-PA 페이지 테이블 모델로는 표현 불가능한 DPPolicy의 서브페이지
+샤딩 (예: 128B 페이로드 × 4KB 페이지) 시나리오를 위해 의도적으로 추가된
+**stopgap**이다. 진짜 HW MMU와는 다른 단순화임을 ADR로 못 박을 필요가 있다.
+
+## Decision
+
+### D1. 이중 역할의 명시 — 컴포넌트와 유틸리티
+
+`PeMmuComponent`는 단일 클래스 안에서 다음 두 인터페이스를 노출한다:
+
+- 컴포넌트 인터페이스: `_inbox` 소비, `_worker` 루프 (MMU sideband 메시지 처리).
+- 유틸리티 인터페이스: `pe_mmu.mmu` 속성으로 underlying `PeMMU` 객체를 노출 —
+  PE_DMA / PE_GEMM이 이 객체를 직접 들고 `translate()`를 호출.
+
+후자는 **layer skip이 아니다**: PE 내부는 ADR-0007이 정의한 "components" 레이어
+하나 안의 sibling 관계이고, 같은 PE prefix에서 가져온 PE_MMU 객체에 대한 직접
+호출은 cross-layer가 아니다. cross-layer 위반은 runtime API / sim_engine /
+components 경계를 넘는 경우에만 적용된다.
+
+### D2. Latency 모델: `translate()`는 순수 함수, overhead는 호출자 책임
+
+`PeMMU.translate()`는 순수 함수이며 SimPy yield를 하지 않는다. 호출자(PE 엔진)
+가 변환 후 `if self._mmu.overhead_ns > 0: yield env.timeout(self._mmu.overhead_ns)`
+를 자기 process에서 발생시킨다.
+
+이유: PE 엔진의 SimPy process는 이미 자체 record_start / record_end (op_log)
+hook을 들고 있어 timing을 일관되게 잡을 수 있다. MMU가 별도의 process를 만들면
+PE 엔진의 처리 흐름을 두 갈래로 쪼개 op_log/pipeline overlap 의미가 흐려진다.
+
+#### D2.1. 현재 구현의 비대칭 — pipeline vs non-pipeline (Known asymmetry)
+
+본 ADR 작성 시점의 `pe_dma.py` 구현은 두 호출 경로에서 overhead 처리가 다르다:
+
+- **non-pipeline (`handle_command`)**: `translate()` 직후
+  `if self._mmu.overhead_ns > 0: yield env.timeout(self._mmu.overhead_ns)` 를
+  발생시킨다.
+- **pipeline (`_do_pipeline_dma`)**: `translate()` 만 호출하고 overhead timeout을
+  **생략**한다 — 함수 주석에 "same logic as non-pipeline path"라고 적혀 있으나
+  실제로는 일치하지 않는다.
+
+기본 토폴로지에서 `tlb_overhead_ns = 0.0` 이라 이 차이는 timing에 직접 드러나지
+않으나, `tlb_overhead_ns > 0` 으로 설정한 시뮬레이션에서는 pipeline 경로의
+GEMM/Math 가 non-pipeline 동일 워크로드 대비 MMU overhead 만큼 빠르게 측정된다.
+
+D2의 계약은 "**모든** 호출자가 overhead를 책임진다" 이며, pipeline 경로의 누락은
+**의도된 설계가 아니라 구현 비일관성**이다. ADR-0014 D6 (pipeline self-routing)
+이 이 overhead를 면제한다고 명시한 부분은 없다.
+
+조치 선택지(별도 Phase 1/2 제안 필요):
+
+- (a) `_do_pipeline_dma` 에서도 `if mmu.overhead_ns > 0: yield env.timeout(...)`
+  를 추가하여 D2 계약과 일치시킨다 — 권장.
+- (b) D2 계약을 "non-pipeline 경로에만 적용" 으로 좁히고, pipeline 경로의 면제를
+  ADR-0014 D6 갱신과 함께 정당화한다 — overhead 의미가 약해지므로 비권장.
+
+본 ADR은 (a) 를 권장하며, accept 전 또는 직후의 별도 작은 변경으로 이를
+교정하는 것을 가정한다.
+
+### D3. 페이지 테이블 구조 — 서브페이지 region 리스트 (stopgap)
+
+`self._table: dict[vpn, list[(start_in_page, end_in_page, pa_at_offset_zero)]]`
+구조로 한 페이지 안에 여러 disjoint region을 보유할 수 있다.
+- `map(va, pa, size)`: 페이지를 가로지르면 region들을 **append**한다.
+- `translate(va)`: VPN으로 region 리스트를 가져온 후, **역방향**으로 순회하며
+  처음 매칭되는 region을 채택 (last-write-wins).
+- `unmap(va, size)`: extent가 unmap 범위에 **완전히 포함된** region만 제거한다.
+  경계가 어긋난 부분 overlap은 그대로 남기며, 매핑 호출자는 mapping과 동일한
+  경계로 unmap할 책임을 진다.
+
+이는 진짜 HW MMU와는 다른 **시뮬레이터 stopgap**임을 ADR-0011 VA 모델 보강
+요소로 명시한다. DPPolicy 서브페이지 샤딩 시 last-write-wins overwrite로 인한
+조용한 미스라우팅을 방지하기 위함이다 (메모리 노트: project_mmu_subpage_stopgap).
+
+### D4. PageFault는 PA fallback 신호다
+
+매핑이 없는 VA로 `translate()`가 호출되면 `PageFault`가 발생한다. PE_DMA는 이
+예외를 잡아 **원본 주소를 PA로 그대로 사용**한다 (ADR-0011의 PA fallback 호환
+경로). 따라서 PageFault는 에러가 아닌 "VA 매핑 부재 시 PA로 해석한다"는 신호다.
+
+이 호환 경로는 ADR-0011이 합의한 PA-only 모드와의 후방 호환을 유지하기 위한
+의도된 동작이다.
+
+### D5. MMU sideband 메시지의 수신 계약
+
+`MmuMapMsg` / `MmuUnmapMsg`는 fabric을 통해 PE_MMU 컴포넌트의 `_inbox`로
+도달한다 (R10이 명시하는 "MMU map 설치는 fabric latency를 따른다"). 메시지
+schema는 runtime API (`runtime_api/kernel.py`) 가 정의하며, 현재 형식:
+
+- `MmuMapMsg.entries: tuple[dict, ...]` — 각 dict는 `{"va": int, "pa": int,
+  "size": int}` 키를 갖는다.
+- `MmuUnmapMsg.entries: tuple[dict, ...]` — 각 dict는 `{"va": int, "size": int}`
+  키를 갖는다.
+
+PE_MMU 측 수신 처리:
+
+1. `_worker` 가 `_inbox.get()` 에서 메시지 한 건을 꺼낸다.
+2. `hasattr(msg, "request")` 로 Transaction wrapper 인지 확인.
+3. `isinstance(msg.request, MmuMapMsg)` 이면 각 entry 에 대해
+   `self._mmu.map(va=e["va"], pa=e["pa"], size=e["size"])`.
+4. `isinstance(msg.request, MmuUnmapMsg)` 이면 각 entry 에 대해
+   `self._mmu.unmap(va=e["va"], size=e["size"])`.
+5. 둘 다 `msg.done.succeed()` 로 완료 통지.
+
+외부 호출자(runtime API 측)가 `done`을 await하면 "매핑이 디바이스에 설치된
+시점"이 SimPy 시간으로 보장된다 — 이 wait이 ADR-0011이 요구하는 "MMU map
+installation incurs measured fabric latency" 의 실현이다.
+
+이 ADR은 sideband 메시지의 **sender 와 fan-out 정책**을 정의하지 않는다 —
+그것은 runtime API 책임이다. 본 ADR은 PE_MMU 측 수신 계약만 명시한다.
+
+### D6. 비-MMU Transaction은 일반 forwarding으로 위임
+
+`_worker`가 inbox에서 꺼낸 메시지의 `request`가 `MmuMapMsg` / `MmuUnmapMsg`가
+아닌 경우 (또는 `request` 속성이 없는 경우) `_forward_txn`으로 떨군다. 이는
+미래에 PE_MMU가 cube-internal NOC 상의 통과 노드로 사용될 가능성을 차단하지
+않기 위함이다 (현재는 그런 통과 트래픽이 없으나, 토폴로지 변경에 대해 안전).
+
+## Alternatives Considered
+
+### A1. translate()를 SimPy generator로 만들기
+
+기각. D2에서 설명한 대로, PE 엔진의 op_log/pipeline overlap 의미가 흐려진다.
+호출자 측에서 timeout을 일으키는 현재 패턴이 op_log 회계와 일치한다.
+
+### A2. 서브페이지 region 리스트 대신 페이지 크기 자체를 작게 하기 (예: 128B)
+
+기각. 페이지 테이블 메모리 폭발과 cube-wide map message 크기 폭발을 초래한다.
+DPPolicy 샤딩이 128B를 요구한다 해도 그 외 대다수 매핑은 2MiB 단위이므로,
+페이지 크기를 작게 잡는 것은 평균 비용이 비대해진다.
+
+### A3. PE_MMU를 컴포넌트가 아닌 PE_CPU의 내장 헬퍼로만 두기
+
+기각. ADR-0011이 요구하는 "fabric을 통해 측정된 latency로 MMU map 설치"
+(MmuMapMsg 경로)를 표현하려면 토폴로지 그래프 상의 노드여야 한다. 또한 cube NoC
+visualizer에서 PE_MMU가 노드로 보여야 디버깅·진단이 일관된다.
+
+## Consequences
+
+- PE_MMU의 이중 역할(컴포넌트 + 유틸리티)이 ADR-level에서 정당화되어, 미래의
+  refactor 압박 (둘 중 하나로 통일하라)에 대한 논거가 생긴다.
+- 서브페이지 region 모델이 시뮬레이터 stopgap임을 ADR이 명시 — 이후 LA 모델
+  (ADR-0011) 도입 시 이 stopgap 제거 가능성을 평가하는 기준이 된다.
+- `translate()`가 yield하지 않는다는 계약이 ADR로 굳어지므로, 향후 누군가
+  "MMU에 자체 timeout을 넣자"는 제안을 할 때 D2를 근거로 거절할 수 있다.
+- PA fallback (D4) 이 정상 흐름임이 명시되어, PageFault를 에러로 오인하여
+  방어 로직을 추가하는 일을 막는다.

docs/adr-ko/ADR-0040-dev-pe-tcm-component-model.md

@@ -0,0 +1,142 @@
+# ADR-0040: PE_TCM Component Model — 듀얼 채널 BW 직렬화
+
+## Status
+
+Accepted (2026-05-20).
+
+ADR-0014 (PE Pipeline Execution Model) 가 "PE_TCM은 BW-기반 직렬화 scratchpad
+memory" 라고 언급하나 (D1), TCM 컴포넌트 자체의 정확한 동작 모델을 별도로
+명시한다.
+
+## First action (제일 처음에 하는 일)
+
+`start()`가 호출되면 즉시 두 개의 `simpy.Resource(env, capacity=1)`을 만들고
+`self._read_res` / `self._write_res`에 보관한다. 이 두 자원이 **읽기 채널**과
+**쓰기 채널**을 각각 1-in-flight로 직렬화하는 단일 결정 포인트다.
+
+런타임 첫 동작: `_worker`가 `_inbox`에서 메시지를 한 건 꺼내 타입 분기:
+
+- `TcmRequest` (`pe_fetch_store`에서 옴): `env.process(self._handle_tcm_request)`로
+  포크. 즉 **TCM의 첫 일은 "방향 (read/write)에 맞는 채널 락을 잡는 것"**.
+  락 획득 후 `bw > 0 and nbytes > 0` 이면 `delay_ns = nbytes / bw` 만큼
+  `env.timeout`, 그리고 `req.done.succeed()`.
+- 그 외 (Transaction): `env.process(self._forward_txn)`로 포크 (legacy fabric
+  통과 경로).
+
+생성 시점에 `node.attrs["read_bw_gbs"]` / `node.attrs["write_bw_gbs"]`
+(default 각 `512.0 GB/s`) 를 읽어 보관해 둔다.
+
+## Context
+
+PE 파이프라인 (ADR-0014 D1, D6) 에서 PE_TCM은 다음 두 종류의 트래픽을 받는다:
+
+1. **PE_FETCH_STORE → PE_TCM의 `TcmRequest`** — TCM ↔ Register File 전송 시,
+   PE_FETCH_STORE가 TCM의 BW로 직렬화된 access latency를 받아오기 위해 짧은
+   sideband 요청을 보낸다 (`direction = "read"` 또는 `"write"`, `nbytes`,
+   `done` 이벤트).
+2. **legacy Transaction forwarding** — 토폴로지 그래프 상에서 TCM이 통과 노드로
+   잡힐 가능성에 대비한 일반 forwarding 경로 (현재 critical path에서는 사용되지
+   않으나 보존됨).
+
+문제: ADR-0014는 "PE_TCM은 BW-기반 직렬화"라고만 언급한다. 그러나 코드에는
+명시적으로:
+
+- **읽기와 쓰기는 별도 채널이며 동시 진행 가능**, 다만 같은 방향끼리는
+  cap=1로 직렬화된다.
+- BW는 `read_bw_gbs` / `write_bw_gbs` 두 값으로 분리 설정 가능하다.
+- `delay_ns = nbytes / bw_gbs` 공식 (단위 환산: GB/s × ns ≈ B 라는 약식).
+- nbytes==0이면 BW 항을 건너뛰지만 채널 락은 잡는다.
+- `run()`은 `overhead_ns` (default 0.0) 만큼 yield 하나, 이는 legacy fabric
+  경로(Transaction forwarding)에서만 사용된다.
+
+이 모든 사항을 별도 ADR로 못 박을 필요가 있다. 특히 "왜 read/write가 분리
+채널인가" 와 "BW는 누가 결정하는가" 는 향후 누군가가 capacity=2 등으로 변경하려
+할 때 명확한 근거가 필요한 항목이다.
+
+## Decision
+
+### D1. 듀얼 채널 — read와 write는 독립 자원
+
+`_read_res = simpy.Resource(env, capacity=1)`,
+`_write_res = simpy.Resource(env, capacity=1)`.
+같은 방향의 동시 요청은 자원 큐에서 직렬화되나, 다른 방향끼리는 동시에 진행 가능.
+이는 실제 HW에서 TCM이 듀얼 포트 (read port + write port) 로 운용되는 모델과
+정합되며, GEMM 파이프라인에서 fetch(read)와 store(write)가 시간상 겹치는 정상
+케이스를 BW-직렬화 모델로 표현하기 위해 의도된 분리다.
+
+### D2. 단일 채널의 BW 모델 — `nbytes / bw_gbs`
+
+채널 락 획득 후, `nbytes > 0 and bw > 0`이면 `yield env.timeout(nbytes / bw_gbs)`.
+단위 약식은 GB/s × ns ≈ B 로, 시뮬레이터 전체에서 사용하는 BW 공식과 동일
+(ADR-0033 참고 — 시뮬레이터는 일관된 약식 단위를 사용한다).
+
+- `nbytes == 0`: BW 항은 0이지만 락은 잡혔다가 즉시 풀린다. 이 케이스가 의도된
+  이유: 빈 fetch/store를 보내는 plan generator가 PE_FETCH_STORE 측에서 `nbytes`만
+  0으로 채워 보내는 경우에도, TCM 측의 op_log / 채널 회계가 일관되게 한 번
+  소비된다.
+- `bw == 0` (config 실수): timeout 호출 자체를 skip하므로 0-time pass. 정상
+  세팅에서는 발생하지 않는다.
+
+### D3. BW는 `node.attrs`의 `read_bw_gbs` / `write_bw_gbs`로 설정
+
+기본값 `512.0 GB/s`. 토폴로지 빌더 (`topology/builder.py`) 가 `pe_template`에서
+TCM을 인스턴스화할 때 해당 attrs를 전달한다. 기본값 변경은 ADR-0014 D1 또는
+ADR-0033 latency model 측의 의사결정과 함께 가야 한다.
+
+### D4. TcmRequest의 schema는 PE_TCM이 owner다
+
+`@dataclass TcmRequest(direction: str, nbytes: int, done: simpy.Event, tag: str = "")`
+는 `components/builtin/pe_tcm.py`에 정의된다. PE_FETCH_STORE는 이 dataclass를
+import해서 생성·송신만 한다. 호출자 측이 schema를 정의하지 않는 이유:
+
+- BW 직렬화의 의미는 TCM 측 책임 — 어떤 필드가 직렬화 결정에 쓰이는가는 TCM이
+  결정한다.
+- `direction` 문자열을 `"read"` / `"write"` 둘로 좁히는 유효값 검증도 TCM 측에
+서 담당 (`_handle_tcm_request`의 if/else 분기).
+
+### D5. legacy Transaction forwarding 경로의 보존
+
+`_worker`가 `TcmRequest`가 아닌 메시지를 받으면 `_forward_txn`으로 보낸다. 이때
+`run()`의 `overhead_ns`가 적용된다. 현재 표준 PE 파이프라인에서는 TCM이
+Transaction의 통과 노드로 잡히지 않으나, fabric 토폴로지가 향후 변경될 때를
+위해 보존한다 (D1 의 사용 패턴과 직교).
+
+이 경로는 op_log 측에서 일반 Transaction 회계로 잡히며, BW 채널 락은 잡지 않는다.
+
+### D6. PE_TCM은 자체 데이터 저장소가 아니다 (timing only)
+
+TCM은 **시간만** 모델링한다. 실제 데이터 페이로드는 sim_engine의 별도
+`memory_store` (있다면) 가 보관하고, TCM 컴포넌트는 그것을 갱신하지 않는다.
+PE_FETCH_STORE도 TcmRequest를 통해 BW 지연만 받아오고 실제 register 컨텐츠는
+별도 경로로 다룬다 (ADR-0020 2-pass data execution 모델 — Phase 2에서 데이터
+처리).
+
+## Alternatives Considered
+
+### A1. 단일 채널 (capacity=2 의 read+write 공유)
+
+기각. fetch(read)와 store(write)가 시간상 겹치는 정상 케이스를 인공적으로
+직렬화하게 되어 PE 파이프라인의 BW upper bound가 잘못 모델링된다.
+
+### A2. 채널 capacity > 1 (예: 2-banked TCM)
+
+기각. 현재 HW 모델은 단일 bank 가정. 멀티-bank로 확장하고 싶다면 별도 ADR이
+필요하며, 그때 D1을 supersede한다. 지금 단계에서 capacity를 늘리면 BW upper
+bound는 그대로인데 명목상의 직렬화만 헐거워져 실제 모델 정확도 ↓.
+
+### A3. BW 공식을 `nbytes / bw + overhead_ns`로 일반화
+
+기각. `overhead_ns`는 D5의 legacy forwarding 경로에만 사용한다. fetch/store
+critical path에 추가 overhead가 필요해지면, 그것은 TCM이 아니라 PE_FETCH_STORE
+측 `run()` 또는 register-file access 모델에 두는 것이 책임 경계 측면에서 더
+적절하다.
+
+## Consequences
+
+- TCM의 BW 회계가 ADR-level에서 굳어지므로, GEMM/Math sweep의 op_log 해석 시
+  "왜 fetch와 store가 동시에 진행되었나" / "왜 같은 방향만 직렬화되나" 같은
+  질문이 빠르게 D1으로 해결된다.
+- 미래의 멀티-bank TCM이나 read/write 비대칭 BW 모델 변경 시 영향 범위가
+  명확해진다 (D1·D2·D3 중 어디를 수정하는지).
+- TCM이 데이터 저장소가 아니라는 점(D6)이 명시되어, ADR-0020 2-pass execution
+  과의 책임 경계가 견고해진다.

docs/adr-ko/ADR-0041-dev-cube-sram-component-model.md

@@ -0,0 +1,187 @@
+# ADR-0041: Cube SRAM Component Model — terminal scratchpad on cube NoC
+
+## Status
+
+Accepted (2026-05-20).
+
+ADR-0017 (Cube NOC and HBM Connectivity) 에서 SRAM이 cube NoC의 attachment로
+존재한다고만 언급되는 점을 보완하여, SRAM 컴포넌트 자체의 latency/response
+모델을 명시한다.
+
+## First action (제일 처음에 하는 일)
+
+`_worker`가 `_inbox`에서 Transaction을 한 건 꺼낸 직후 가장 먼저 하는 일은
+`yield from self.run(env, txn.nbytes)` 호출이고, 그 안에서
+`node.attrs["overhead_ns"]` (default `0.0`) 만큼 `env.timeout()`을 발생시킨다.
+
+즉, **SRAM의 첫 일은 "access overhead를 시간으로 표현하는 것"**이다.
+overhead 소비 이후에 `drain_ns` (그 Transaction에 부여된 terminal BW 직렬화 비용)
+를 yield하고, 그 다음에 reverse path로 `ResponseMsg`를 생성하여 발사한다.
+
+이는 일반 `ComponentBase._worker`와 다른 점이 있다: SRAM은 **terminal node**
+임을 알고 있어서 `_forward_txn`을 거치지 않고 자체 워커가 `run → drain →
+_send_response` 순서를 명시한다.
+
+## Context
+
+cube 토폴로지 (`topology/builder.py`) 는 cube마다 다음 명명된 노드를 만든다:
+
+- `sip{S}.cube{C}.m_cpu`
+- `sip{S}.cube{C}.sram`
+- `sip{S}.cube{C}.hbm_ctrl` (PE당 partition)
+- `sip{S}.cube{C}.pe{P}` (PE 내부 sub-component들)
+
+SRAM은 cube NoC 의 attachment 중 하나로, 가장 가까운 router에 부착된다
+(`topology/mesh_gen.py`가 placement 좌표로 nearest router 결정 후 `attach`에
+추가). 빌더는 `sram ↔ router` 양방향 엣지를 깐다 (BW: `sram_to_router_bw_gbs`,
+기본 `128.0 GB/s`).
+
+SRAM의 두 가지 핵심 역할:
+
+1. **fabric terminal**: cube NoC에서 SRAM으로 향한 메모리 access Transaction의
+   끝점. SRAM이 access overhead와 drain을 소비하고 response를 reverse path로
+   되돌린다.
+2. **IPCQ slot tier 중 하나**: ADR-0023 D9.7 가 정의한 `buffer_kind ∈ {tcm,
+   sram, hbm}` 중 `sram` 티어의 slot bw/overhead를
+   `common/ipcq_types._BUFFER_KIND_BW`에서 참조 — 현재 값 `(512.0 GB/s, 2.0 ns)`.
+   이 값은 SRAM 노드 attrs의 `overhead_ns`와는 별도이며, IPCQ slot 회계 시점에서
+   PE_DMA가 시간으로 환산한다.
+
+이 두 역할은 하나의 SRAM 컴포넌트에서 동시에 충족되는데, 별도 ADR이 없으면:
+
+- "SRAM은 어떤 latency를 모델링하나?" — fabric drain + overhead, 아니면 IPCQ
+  티어의 slot latency? — 답이 흩어진다.
+- 미래에 SRAM 크기 (`size_mb`) attr이 실제로 어떤 의미를 갖는지 불명확. 현재
+  코드는 size를 사용하지 않으며 timing만 모델링한다.
+- SRAM이 cube의 어떤 router에 붙는지 (placement-based)에 대한 의사결정 근거가
+  토폴로지 코드 안에만 있다.
+
+## Decision
+
+### D1. SRAM은 cube NoC의 terminal scratchpad 노드다
+
+`SramComponent`는 `ComponentBase`를 상속하나 `_worker`를 오버라이드해서 terminal
+의미를 직접 표현한다:
+
+```
+while True:
+    txn = yield self._inbox.get()
+    yield from self.run(env, txn.nbytes)     # overhead_ns
+    if drain_ns > 0: yield env.timeout(drain_ns)
+    yield from self._send_response(env, txn)
+```
+
+이 패턴은 SRAM이 reverse path를 알아야 하므로 일반 `_forward_txn` (다음 hop으로
+forward)이 아닌 자체 워커가 필요하다.
+
+#### D1.1. 현재 미사용 — `_worker` 오버라이드는 dormant 경로다
+
+본 ADR 작성 시점의 코드베이스에서는, **어떤 컴포넌트도 SRAM 노드로 Transaction
+을 실제로 전송하지 않는다**. 확인된 SRAM 노드 ID 참조 위치:
+
+- `policy/routing/router.py` 등 routing helper — path 조회 가능성만 보장.
+- `components/builtin/pe_dma.py::_handle_ipcq_inbound` — IPCQ slot의
+  `buffer_kind == "sram"` 일 때 `bank_node = f"{cube_prefix}.sram"` 의 *path*
+  만 조회하여 `compute_drain_ns(path, ...)` 로 환산, **로컬에서 timeout** 한다.
+  Transaction 자체는 SRAM 노드로 흘러가지 않는다 (D4 참고).
+- `tests/test_routing.py` — `find_path("sip0.cube0.pe0", "sip0.cube0.sram")`
+  로 connectivity만 검증.
+
+따라서 `_worker`/`_send_response` 오버라이드는 **dormant code path** 이다.
+삭제하지 않고 보존하는 이유:
+
+- 향후 SRAM이 실제 fabric Transaction의 종점(예: M_CPU → SRAM 명시 access)이
+  되는 토폴로지 변경 시 즉시 사용 가능.
+- ADR-0017 (Cube NOC) 가 정의한 cube-attached scratchpad 의미에서 종점 동작은
+  의미상 자연스러우므로, 의도된 placeholder 다.
+
+이 dormant 상태가 종료되는 시점은 별도 ADR(또는 본 ADR의 후속 revision)이
+명시한다.
+
+### D2. ResponseMsg 생성과 reverse path 발사
+
+`_send_response`는:
+
+1. `reverse_path = list(reversed(txn.path))`로 역방향 경로 산출.
+2. `ResponseMsg(correlation_id=txn.request.correlation_id, request_id=...,
+   src_cube=<this cube>, src_pe=-1, success=True)` 생성.
+3. `Transaction(request=resp_msg, path=reverse_path, step=0, nbytes=0,
+   done=env.event(), is_response=True)` 로 감싸 `out_ports[reverse_path[1]]` 로
+   put.
+4. reverse path가 비정상이거나 (`< 2 hops`) ctx가 없으면, fallback으로 원본
+   `txn.done.succeed()` 만 호출.
+
+`src_pe = -1`은 "SRAM은 PE-localized가 아니다"를 의미한다. `src_cube`은 노드
+ID (`sip{S}.cube{C}.sram`) 의 cube 인덱스를 파싱해 채운다.
+
+### D3. Timing 파라미터는 `overhead_ns`와 wire-side `drain_ns`로 분리
+
+- **컴포넌트 측 latency**: `node.attrs["overhead_ns"]`. 기본 토폴로지에서는 `2.0
+  ns` 정도로 세팅.
+- **링크 측 직렬화**: `drain_ns`는 Transaction이 도착 시점에 carry해 온 값으로,
+  ADR-0015 (port/wire 모델) 의 wire-side BW 직렬화 결과다. SRAM은 이를 그대로
+  yield하기만 한다.
+- `size_mb` (default `32 MiB`) attr은 현재 timing에 사용되지 않는다 — 향후
+  capacity-aware 모델이 도입되면 그때 의미를 부여한다 (별도 ADR에서).
+
+### D4. IPCQ slot 회계는 SRAM 컴포넌트가 직접 모델링하지 않는다
+
+ADR-0023 D9.7 에 따른 IPCQ slot의 SRAM-티어 write latency는 PE_DMA의
+`_handle_ipcq_inbound`가 직접 `slot_io_latency_ns("sram", nbytes)`를 호출하여
+시간을 소비한다 (그 함수는 `common/ipcq_types._BUFFER_KIND_BW["sram"]` 의 값을
+사용). 즉:
+
+- SRAM 컴포넌트가 fabric Transaction을 받아 처리할 때는 **D1·D2·D3** 만 적용.
+- IPCQ slot이 SRAM에 살 때는 PE_DMA가 IPCQ slot-write 시점에 별도로 시간을
+  지불 — 이는 SRAM 컴포넌트 코드와 무관하며, IPCQ 측 회계다.
+
+이 분리는 의도된 것: IPCQ는 fast path (sub-cycle slot bookkeeping) 라 fabric
+Transaction을 거치지 않으므로, SRAM이 IPCQ를 인지할 필요가 없다.
+
+### D5. SRAM의 cube NoC 부착 위치는 placement-driven
+
+`topology/mesh_gen.py`는 `placement.sram.pos_mm` (`topology.yaml` 기본
+`[1.5, 9.0]`)을 보고 가장 가까운 router의 `attach`에 `"sram"`을 추가한다. 빌더
+(`topology/builder.py` 의 attachment 루프)가 그 attach 정보를 보고 `sram` 노드와
+router 사이에 양방향 엣지를 깐다.
+
+이 의사결정은 SRAM 컴포넌트 코드 외부 (mesh_gen / builder) 에 있으며, 컴포넌트
+는 어느 router에 붙었는지 알 필요가 없다. 컴포넌트는 `txn.path` / `reverse_path`
+가 router를 거쳐 자신에게 도달한다는 사실만 알면 된다.
+
+### D6. SRAM은 자체 데이터 저장소가 아니다 (timing-only)
+
+ADR-0040 D6 과 같은 맥락: SRAM 컴포넌트는 시간만 모델링하며, 실제 데이터
+페이로드는 sim_engine의 `memory_store` (있을 때) 가 보관한다.
+
+## Alternatives Considered
+
+### A1. SRAM이 `_forward_txn`을 그대로 사용하고 IO_CPU / HBM_CTRL 처럼 별도 응답 노드를 두기
+
+기각. cube NoC 상에서 SRAM은 terminal이며, 응답을 받아 줄 별도 노드를 두면
+의미 없는 hop이 늘어나고 ADR-0017 의 cube NoC 단순화 정신에 어긋난다.
+
+### A2. SRAM이 BW 직렬화를 자체 resource로 모델링
+
+기각. 링크 측 BW 직렬화 (`drain_ns`) 가 이미 의미를 충분히 잡고 있다. 컴포넌트
+내부에 또 `simpy.Resource`를 두면 ADR-0015 wire-side 모델과 이중계산을 야기.
+
+### A3. SRAM이 IPCQ slot 회계를 컴포넌트 측에서 처리
+
+기각. D4에서 명시한 대로 IPCQ는 fast path며 fabric Transaction을 통과하지
+않는다. SRAM이 IPCQ를 인지하면 책임이 두 갈래로 갈라져 추론이 어려워진다.
+
+### A4. `size_mb`로 capacity-aware latency 모델
+
+기각 (현재 단계). capacity는 토폴로지 visualizer 측 라벨링 정도에만 쓰이며,
+실제 timing 영향은 아직 모델링하지 않는다. 필요해지면 별도 ADR로 도입.
+
+## Consequences
+
+- SRAM의 timing 모델이 `overhead_ns + drain_ns + ResponseMsg(reverse_path)`로
+  ADR-level에서 굳어지므로, 누군가 IPCQ slot latency를 SRAM 컴포넌트에 추가하려
+  할 때 D4를 근거로 거절할 수 있다.
+- `size_mb` 가 현재 timing-neutral 임이 명시되어 (D3), 미래의 capacity-aware
+  모델 도입 시 호환성 영향 범위가 좁다.
+- placement-driven router 부착 (D5) 이 명시되어, SRAM 좌표 이동 시 어떤 부분에
+  파급이 있는지 (`mesh_gen`만) 명확해진다.

docs/adr-ko/ADR-0042-prog-tile-plan-generators.md

@@ -0,0 +1,194 @@
+# ADR-0042: Tile Plan Generators — GEMM/Math 파이프라인 plan 빌더
+
+## Status
+
+Accepted (2026-05-20).
+
+본 ADR은 `tiling.py`가 SimPy 컴포넌트가 아니라
+**plan-generator 모듈**임을 명시한다.
+
+ADR-0014 (PE Pipeline Execution Model) 의 D6 (tile plan / self-routing) 가
+tile-plan 생성 알고리즘을 직접 정의하지 않으므로, 본 ADR이 그 비어 있는 자리를
+채운다.
+
+## First action (제일 처음에 하는 일)
+
+`generate_gemm_plan(M, K, N, tile_m, tile_k, tile_n, ..., pe_prefix, a_pinned,
+b_pinned, epilogue_specs)`이 호출되면 가장 먼저 하는 일은 **타일 수 계산과
+컴포넌트 ID 문자열 구성**이다:
+
+```
+M_tiles = max(1, ceil(M / tile_m))
+K_tiles = max(1, ceil(K / tile_k))
+N_tiles = max(1, ceil(N / tile_n))
+dma_id   = f"{pe_prefix}.pe_dma"
+fetch_id = f"{pe_prefix}.pe_fetch_store"
+gemm_id  = f"{pe_prefix}.pe_gemm"
+math_id  = f"{pe_prefix}.pe_math"
+```
+
+즉 **plan generator의 첫 일은 "타일 개수를 ceiling으로 산출하고, 이 PE의
+sub-component ID 4개를 한 번에 짜놓는 것"**이다. SimPy 이벤트나 환경 객체는
+일절 다루지 않는다 — 이 모듈은 순수 함수다.
+
+`generate_math_plan(M, N, tile_m, tile_n, ..., math_op, src_addr, dst_addr,
+pe_prefix)` 도 마찬가지로 `M_tiles`, `N_tiles` 산출과 component ID 3개
+(`dma_id`, `fetch_id`, `math_id`) 구성이 첫 일이다.
+
+## Context
+
+ADR-0014 D6은 "PE_SCHEDULER가 CompositeCmd를 받으면 TilePlan을 생성하고
+self-routing tile token을 피드한다"고만 합의했다. 그러나 코드에서는 **plan
+생성 알고리즘의 구체적 내용**이 `src/kernbench/components/builtin/tiling.py`
+모듈에 자리잡고 있고, 이 모듈은:
+
+- 컴포넌트가 아니라 **순수 함수**의 모음이다 (`generate_gemm_plan`,
+  `generate_math_plan`).
+- SimPy 환경, 큐, op_log, hook 등에 의존하지 않는다.
+- 결과로 `PipelinePlan` (dataclass) 를 돌려준다.
+
+기존 G4 분석은 `tiling.py`를 컴포넌트로 잘못 가정했으나, 실제는 PE_SCHEDULER에
+주입되는 plan-builder 함수다. 이 차이는 ADR-0014 의 D6 와 짝을 이루는 별도
+ADR로 못 박혀야 한다 — 그렇지 않으면:
+
+- "tile plan을 만드는 책임이 PE_SCHEDULER인가 별도 모듈인가" 가 모호.
+- GEMM plan과 Math plan의 stage sequence 가 일관성 있는지 (예: FETCH/STORE 위치)
+  의사결정 근거가 흩어진다.
+- `a_pinned` / `b_pinned` / `epilogue_specs` 같은 옵션이 왜 plan 단에서 분기되는지
+  근거 없음.
+
+## Decision
+
+### D1. tiling은 순수 plan-generator 모듈이며 컴포넌트가 아니다
+
+`components/builtin/tiling.py`는 ComponentBase 하위 클래스를 정의하지 않는다.
+모듈-레벨 함수 두 개만 노출한다:
+
+- `generate_gemm_plan(...) -> PipelinePlan`
+- `generate_math_plan(...) -> PipelinePlan`
+
+토폴로지 그래프에서 `tiling` 이라는 노드는 존재하지 않는다. 명명상 `builtin/`
+디렉터리에 있는 이유는 PE_SCHEDULER (ADR-0014 D6) 의 직접 helper이기 때문이며,
+의미상으로는 PE_SCHEDULER 내부 utility에 가깝다.
+
+### D2. GEMM plan의 stage 시퀀스 — `M → N → K` order
+
+각 (m, n, k) 타일에 대한 stage 시퀀스 (operand pinning과 epilogue 미적용 기본):
+
+```
+[DMA_READ(A)] → [DMA_READ(B)] → FETCH → GEMM
+                                ↑
+                                ↓
+(last k tile only)              [MATH(output_tile)]* → STORE → DMA_WRITE
+```
+
+`k_tile` epilogue는 매 K-타일마다 GEMM 직후, `output_tile` epilogue는 (m,n)당
+마지막 K-타일에서 STORE/DMA_WRITE 직전에 한 번. K-루프 누적자(accumulator) 는
+RegFile에 남아 K 타일들 사이에 STORE/DMA_WRITE가 발생하지 않는다 (last_k에서만
+출력).
+
+### D3. Operand pinning — `a_pinned` / `b_pinned`
+
+호출자가 `a_pinned=True`로 호출하면 **모든 (m, n, k) 타일에서 A DMA_READ를
+생략**한다. 의미: 호출자(예: `tl.composite`)가 사전에 `tl.load`로 A 전체를
+TCM에 한 번 적재했음을 plan generator에 알리는 신호.
+
+이 분기는 plan 단에서 결정한다 (런타임 분기 아님). 따라서 op_log 상의 stage
+record 수는 pinning에 따라 결정적으로 달라지며, sweep 분석 측 (예: gemm_sweep
+의 stage record count) 이 이 결정을 그대로 본다.
+
+### D4. Epilogue scope — `k_tile` vs `output_tile`
+
+`epilogue_specs`는 op-spec 객체의 iterable이다. 각 op 객체는 다음 속성을 갖는
+다고 가정한다:
+
+- `op.kind: str` — math op 이름 (예: `"dequant"`, `"bias"`, `"relu"`, `"scale"`).
+  stage의 `params["op_kind"]` 로 들어간다.
+- `op.scope: Scope` — `Scope.K_TILE` 또는 `Scope.OUTPUT_TILE` (`Scope` 는
+  `kernbench.common.pe_commands` 에 정의된 enum).
+- op-별 추가 필드 (예: `bias`, `scale`, `factor`) — 현재 plan generator는 사용
+  하지 않으며 런타임 (PE_MATH) 측이 소비.
+
+plan generator는 `getattr(o, "scope", None)` 기준으로 두 그룹으로 분기:
+
+- `scope == Scope.K_TILE`: 매 K-타일 GEMM 직후 MATH stage 추가.
+- `scope == Scope.OUTPUT_TILE`: (m, n)당 마지막 K-타일 STORE 직전 MATH stage
+  추가.
+
+`scope` 속성이 없거나 두 enum 어느 쪽도 아닌 op는 **plan에 포함되지 않는다**
+(`getattr(..., None) == Scope.X` 가 둘 다 False). 기본값(`output_tile`) 채택은
+**호출자(예: `tl.composite`) 측 책임**이며, plan generator는 이미 채워진 scope
+값을 보고 분기할 뿐이다 (ADR-0014 의 composite epilogue 계약과 정렬).
+
+`Scope` 임포트는 `pe_commands ← pe_types ← tiling` 의 순환 참조를 피하기 위해
+함수 내부에서 lazy import 한다. 이는 의도된 패턴이며 개선 대상이 아니다 (D1의
+"tiling은 PE_SCHEDULER의 utility" 관점에서, pe_commands에 대한 컴파일타임 의존
+이 없는 편이 모듈 경계를 깔끔히 유지함).
+
+### D5. Math plan의 stage 시퀀스 — `M → N` order
+
+각 (m, n) 타일에 대한 stage 시퀀스:
+
+```
+DMA_READ → FETCH → MATH → STORE → DMA_WRITE
+```
+
+K 차원이 없으므로 epilogue / accumulator residency 같은 개념은 적용되지 않는다.
+PE_FETCH_STORE의 register-file 회계는 GEMM plan과 동일한 방식으로 다뤄진다.
+
+### D6. plan은 데이터다 — SimPy 의존성 없음
+
+`PipelinePlan` 은 `pe_types.py`에 정의된 dataclass로, `tiles: list[TilePlan]`을
+보유. 각 `TilePlan` 은 `stages: tuple[Stage, ...]` 를 보유. plan 자체는
+immutable에 가까운 데이터 구조이며 (Stage 의 `params: dict` 만 mutable),
+SimPy 객체나 event를 갖지 않는다.
+
+런타임 시점에 PE_SCHEDULER가 plan 의 첫 stage를 보고 `TileToken`을 생성하여
+파이프라인에 피드하며, TileToken 이 `plan: TilePlan`, `stage_idx: int`,
+`params: dict` 를 들고 다닌다. self-routing은 `TileToken.advance()` 가 다음
+stage의 `params`를 캐시하는 방식으로 진행된다 (ADR-0014 D6).
+
+### D7. plan generator의 contract — pure, deterministic, idempotent
+
+같은 입력으로 두 번 호출하면 같은 PipelinePlan을 돌려준다 (`TilePlan.stages`의
+순서까지 deterministic). 이 contract는 ADR-0014 D6 의 "결정적 tile dispatch
+순서" 요구와 정렬된다.
+
+부수효과(SimPy event, file I/O, 글로벌 상태) 없음 — 테스트에서 환경 객체 없이
+호출 가능 (`tests/test_pe_pipeline.py`의 일부 케이스가 이 방식 사용).
+
+## Alternatives Considered
+
+### A1. tiling을 컴포넌트로 만들기 (e.g., PE_PLANNER)
+
+기각. plan 생성은 SimPy 시간을 소비하지 않는 결정 알고리즘이다. 컴포넌트로
+만들면 (a) inbox·자원 등 불필요한 인프라가 따라붙고, (b) PE_SCHEDULER 가
+"plan 받기" → "tile 피드" 두 단계를 분리해 받게 되어 의미 없는 hop이 생긴다.
+
+### A2. plan 생성을 PE_SCHEDULER 클래스 메서드로 옮기기
+
+기각 (현재). 모듈 분리가 (1) 테스트 용이성, (2) 다른 plan 알고리즘 (예:
+DTensor-aware plan) 도입 시 추가 함수만 정의하면 되는 확장성을 준다. 만약 향후
+plan 종류가 많아져 명시적 dispatch가 필요해지면, 그때 PE_SCHEDULER에 plan
+factory를 두는 것을 별도 ADR로 도입한다.
+
+### A3. plan을 immutable로 강제 (frozen dataclass + tuple)
+
+부분 채택. `Stage` 와 `TilePlan` 은 dataclass지만 frozen은 아니다. 이유:
+`Stage.params: dict` 가 plan generator 시점에 채워지고 런타임에서 읽히기만 한다
+(TileToken 이 advance 시 캐시할 뿐). 완전 frozen은 dict → frozendict 마이그레이션
+비용 대비 이득이 적다. 다만 plan 단계 외에는 mutation 하지 말 것을 컨벤션으로
+유지한다.
+
+## Consequences
+
+- `tiling.py`가 컴포넌트가 아니라 plan-generator 모듈임이 ADR-level에서
+  명시되어, G4 같은 미래의 "이 컴포넌트는 ADR이 없다"는 분석을 차단한다.
+- GEMM plan의 stage sequence (D2) 와 pinning/epilogue 분기 (D3·D4) 가 ADR로
+  굳어지므로, sweep 분석 (`scripts/gemm_sweep.py`)의 stage record count 해석
+  근거가 명확해진다.
+- plan generator의 pure contract (D7) 덕분에 테스트가 환경 없이 plan 검증
+  가능 — ADR-0013 (verification strategy) 의 "behavior validated by tests with
+  meaningful input cases" 정신과 정렬.
+- 향후 DTensor-aware plan, K-major plan 등 새 plan 종류 추가 시 본 ADR이
+  baseline 역할 — 새 함수만 추가하고 D1·D6·D7을 따른다.

docs/adr-ko/ADR-0043-eval-allreduce-harness.md

@@ -0,0 +1,131 @@
+# ADR-0043: Allreduce 평가 하니스 — `tests/sccl/`
+
+## Status
+
+Accepted
+
+`tests/sccl/` 평가 하니스를 문서화한다; 구현과 대조 검증 완료
+(상수, 파일 집합, 스윕 차원을 교차 확인).
+
+**ADR-0054로 개정됨**: 드라이버 코어, sweep, renderer가 `milestone-1h-ccl`
+bench(단일 home)로 이동했다; `tests/sccl/_allreduce_helpers.py`는 이제 거기서
+re-export한다(pytest 전용 param 빌더 + `_run_distributed` wrapper는 로컬
+유지). figure 테스트는 변경 없음.
+
+## Context
+
+ADR-0032는 intercube all-reduce *알고리즘*을 정의하고, ADR-0023/0024/0027은
+IPCQ 백엔드, rank=SIP launcher, `mp.spawn`을 정의한다. 그러나 어느 것도
+**allreduce를 어떻게 구동하고 특성화하는가** — 정확성 테스트, latency/
+buffer-kind 스윕, 파생 플롯 — 는 기술하지 않는다. ADR-0013(verification
+strategy)이 일반 정책이라면, 본 ADR은 구체적 allreduce 하니스를 고정하여
+작업의 "평가" 절반이 구현과 함께 문서화되도록 한다.
+
+하니스는 `tests/sccl/`(allreduce 테스트 통합 시 생성된 패키지)에 위치한다.
+이전의 평면적 `tests/test_allreduce_multidevice.py` +
+`tests/test_distributed_*` 레이아웃을 대체한다.
+
+## Decision
+
+### D1. 평가를 공개 `torch.distributed` 경로로 구동
+
+정확성과 스윕은 collective를 실제 DDP 형태 경로 —
+`init_process_group(backend="ahbm") → mp.spawn → dist.all_reduce`
+(ADR-0024/0027) — 로 실행하며, 하위 레벨 `ctx.launch`를 쓰지 않는다.
+`tests/sccl/_allreduce_helpers.py`의 공유 헬퍼
+`_run_distributed(tmp_path, monkeypatch, topo_path, corr_id, n_elem)`가
+엔진을 빌드하고 워커를 실행하고 `(engine, n_cubes)`를 반환한다.
+`monkeypatch.chdir`이 백엔드의 `load_ccl_config()`(cwd 조회)를 케이스별
+임시 `ccl.yaml`로 향하게 한다.
+
+직접 launch 레퍼런스(`run_allreduce`)는 같은 헬퍼 모듈에 유지된다 —
+distributed 테스트용이 아니라, `tests/`의 IPCQ buffer-kind / root-center
+마이크로 테스트가 import하기 때문이다.
+
+### D2. 평가 관심사별 파일 하나
+
+| 파일 | 관심사 | `torch.distributed`? |
+|---|---|---|
+| `test_allreduce_ring_torus_mesh.py` | ring_1d / torus_2d (2×3) / mesh_2d_no_wrap (2×3) 정확성 | yes |
+| `test_distributed_default_topology.py` | `topology.yaml` 그대로의 전체 경로 | yes |
+| `test_plot_latency_sweep.py` | latency 스윕 행 (n_elem × topology) | yes |
+| `test_plot_buffer_kind_sweep.py` | TCM/SRAM/HBM 스윕 행 | yes |
+| `test_plot_topology_diagram.py` | topology.png (순수 matplotlib) | no |
+| `test_plot_comparison_fsim.py` | broken-axis 모델 vs FSIM 비교 | no |
+| `test_intercube_root_center.py` | ADR-0032 center-root latency 가드 (직접 경로) | no |
+
+`_allreduce_helpers.py`는 공유 plumbing(드라이버, config writer, 스윕/
+buffer-kind 상수, 플롯 aggregator, topology-diagram + FSIM 비교 emitter)을
+보유한다. 수집되지 않는다(`test_` 접두사 없음).
+
+### D3. Latency 메트릭 — critical-path `pe_exec_ns`
+
+config별 보고 latency는 `engine._results`에 대한
+`crit_ns = max(pe_exec_ns)` — 가장 느린 rank의 PE 실행 시간 — 이다.
+모든 latency 차트에 그려지고 `summary.csv`에 기록되는 값이다.
+
+### D4. 스윕 차원
+
+- **Latency 스윕**: `n_elem ∈ {8, 32, 64, 128, 512, 1024, 2048, 4096,
+  8192, 16384, 32768, 49152}` (16 제외 — `n_cubes`와 충돌) × topology ∈
+  {ring_1d (6), torus_2d 2×3 (6), mesh_2d_no_wrap 2×3 (6)}.
+- **Buffer-kind 스윕**: `buffer_kind ∈ {tcm, sram, hbm}` × 더 작은
+  `n_elem` 그리드, torus_2d 6-SIP (3×2)에서. buffer_kind는 임시
+  `ccl.yaml`에 설정되며(백엔드가 `init_process_group` 시점에 읽음,
+  ADR-0023 D6) 적용된다.
+
+2×3 / 3×2 그리드는 명시적 `w/h` SIP 해석(ADR-0024 D5)을 행사한다.
+
+### D5. `pytest_sessionfinish` aggregator를 통한 파생 플롯
+
+스윕 테스트는 xdist 친화적이다: 각 parametrized 케이스가 staging 디렉터리에
+JSON 행 하나를 쓴다. conftest `pytest_sessionfinish` 훅(controller 노드
+전용)이 `_allreduce_helpers.py`의 aggregator를 호출한다:
+
+- `_aggregate_sweep_plots()` → topology별 PNG + `summary.csv`
+- `aggregate_buffer_kind_plot()` → TCM/SRAM/HBM 비교 PNG + csv
+
+topology-diagram 및 FSIM-비교 figure는 각자의 `test_plot_*` 테스트가
+직접 emit한다(행 staging 없음 — 각각 `topology.yaml`과 `summary.csv`의
+순수 함수). 모든 출력은 `docs/diagrams/allreduce_latency_plots/`에 떨어지며
+CLAUDE.md에 따라 **파생 아티팩트**다(ADR과 일관, Phase-2 게이트 없음).
+
+### D6. FSIM 비교 레퍼런스는 하드코딩 상수
+
+`emit_comparison_fsim_plot()`은 모델 곡선을 외부 FSIM single-device
+레퍼런스(`366 µs`) 하나와 겹쳐 그리며, 이는 리터럴로 보유된다 — 외부 데이터
+파일 없음. "measured" 시리즈는 시뮬레이터(`op_log` GEMM 카운트,
+`composite_window_ns`)에서, "theoretical" 시리즈는 손으로 도출한 해석적
+모델(ADR-0044 D5가 ADR-미검증으로 표시한 동일 모델)에서 온다.
+
+## Consequences
+
+### Positive
+
+- allreduce가 실제 DDP 스크립트와 같은 API로 평가되므로, 하니스가
+  ADR-0024/0027의 통합 테스트 역할도 겸한다.
+- figure는 매 `pytest` 실행마다 committed 데이터로 재생성된다; 수동 플롯
+  단계 없음.
+- 직사각형 그리드 스윕이 ADR-0024 D5 `w/h` 수정을 드러낸 회귀 커버리지를
+  제공했다.
+
+### Negative / limitations
+
+- 전체 latency 스윕은 기본 `pytest`에서 실행된다(~분 단위); `slow`로
+  표시되지 않는다. (ADR-0044는 GEMM 스윕을 `slow`로 표시하는 것과 대조.)
+- `test_intercube_root_center.py`는 latency *임계값* assertion(ADR-0032
+  center-root 가드)을 보유한다 — 스위트에서 유일한 절대-latency
+  assertion이며 latency 모델 변경(ADR-0033)에 민감하다.
+
+## Dependencies
+
+- **ADR-0013**: verification strategy (본 ADR이 특수화하는 일반 정책).
+- **ADR-0023 / ADR-0024 / ADR-0027**: IPCQ 백엔드, rank=SIP launcher,
+  `mp.spawn` — D1이 구동하는 경로.
+- **ADR-0032**: 평가 대상 알고리즘; D4 그리드가 그 topology 분기를 행사.
+- **ADR-0044**: 형제 격인 GEMM 평가 하니스.
+
+## Open questions
+
+- GEMM 스윕과의 일관성을 위해 latency 스윕을 `slow`로 표시할 것인가?
+- FSIM 레퍼런스를 하드코딩 상수에서 버전 관리되는 데이터 파일로 옮길 것인가?

docs/adr-ko/ADR-0044-eval-gemm-harness.md

@@ -0,0 +1,133 @@
+# ADR-0044: GEMM 평가 하니스 — `scripts/gemm_sweep.py` + `tests/gemm/`
+
+## Status
+
+Accepted
+
+GEMM 평가/특성화 하니스를 문서화한다; 구현과 대조 검증 완료
+(상수, tile 크기, figure 집합, script↔test 분할을 교차 확인). D5/D6
+caveat은 부정확이 아니라 기록된 한계다.
+
+**ADR-0054로 개정됨**: sweep + renderer가 `milestone-1h-gemm` bench(단일
+home)로 이동했다; `scripts/gemm_sweep.py`와 `tests/gemm/`는 이제 거기서
+re-export한다. D1/D2의 "데이터 생성은 수동 script / 무거운 작업은 opt-in"은
+평가-bench 패턴으로 대체된다(하나의 bench가 전부 재생성;
+`MILESTONE_FAST=1`은 committed JSON 재사용).
+
+## Context
+
+ADR-0014(PE pipeline)와 ADR-0042(tile-plan generator)는 GEMM *구현*을
+정의하고, ADR-0033은 latency 모델을 정의한다. 그러나 어느 것도 **GEMM
+성능을 어떻게 스윕하고 특성화하는가** — 타이밍 데이터를 만드는 shape/variant
+스윕과 이를 해석하는 figure — 는 기술하지 않는다. 본 ADR이 그 하니스를
+고정한다.
+
+allreduce 하니스(ADR-0043)와 달리 GEMM 스윕은 **무겁다**(24 sim 실행:
+8 shape × 3 operand-staging variant; `512` shape 하나가 2048 tile). 이
+무게가 아래 분할을 결정한다.
+
+## Decision
+
+### D1. 두 계층 분할 — 무거운 데이터 생성(script) vs. 빠른 figure(test)
+
+- **데이터 생성은 수동 script로 유지**: `scripts/gemm_sweep.py`가
+  `matmul-composite`(ADR-0042 plan)를 CLI와 동일한 `run_bench` 경로로
+  shape × variant에 걸쳐 실행하고, `result.engine.op_log`를 수확하여
+  `docs/diagrams/gemm_sweep.json`(stage별/engine별 wall-clock + occupancy
+  + record count + pe/composite window)을 쓴다.
+- **figure 렌더링은 test 생성**: `tests/gemm/`이 committed `gemm_sweep.json`을
+  읽어 matplotlib PNG를 `docs/diagrams/gemm_plots/`에 렌더링한다. 이
+  테스트는 빠르고 기본 실행된다.
+
+근거: 슬라이드덱 규모의 sim 스윕은 매 `pytest` 실행에 속하지 않지만,
+figure(저렴·결정적)는 자유롭게 재생성되고 CI로 가드되어야 한다. 이는
+CLAUDE.md의 script-vs-test 분할(무거운/수동 생성은 script; 빠른 assertion은
+test)을 반영한다.
+
+### D2. Slow regenerator 테스트가 script를 감싼다
+
+`tests/gemm/test_gemm_sweep.py`는 `@pytest.mark.slow`로 표시된다(기본
+`addopts: -m "not slow"`에서 제외). 이는 `scripts/gemm_sweep.py`를
+subprocess로 호출하여 `gemm_sweep.json`을 on-demand로 재생성한다
+(`pytest -m slow tests/gemm/test_gemm_sweep.py`). 스윕 로직은 단일
+home(script)을 가지며 테스트는 이를 감싸기만 하므로 sim 구동 코드의
+중복이 없다.
+
+### D3. Figure 집합 (3개 차트, `load_ref` variant)
+
+| 테스트 | PNG | 내용 |
+|---|---|---|
+| `test_plot_gemm_stage_breakdown.py` | `gemm_stage_breakdown.png` | stage별 engine wall-clock (DMA in / Fetch / GEMM / DMA out) |
+| `test_plot_gemm_mac_utilization.py` | `gemm_mac_utilization_measured.png` | GEMM util % + useful eff % |
+| `test_plot_gemm_mac_utilization.py` | `gemm_mac_utilization_theoretical_vs_measured.png` | theoretical vs 시뮬레이터-measured util/eff |
+
+`tests/gemm/_gemm_plot_helpers.py`가 공유 renderer를 보유한다(시리즈 로직은
+`scripts/build_overview_slides.py`의 GEMM `_render_*` 함수를 미러링하며,
+그쪽은 여전히 PPTX에 네이티브로 그린다). 수집되지 않음(`test_` 접두사
+없음). 각 `test_plot_*`는 `gemm_sweep.json`이 없으면 skip한다.
+
+### D4. Tile 크기는 데이터 기반; under-tile shape는 표시
+
+Tile 크기는 `gemm_sweep.json`(`tile_sizes`)에서 읽으며, 이는 스윕이
+`PeSchedulerComponent.TILE_M/K/N = 32/64/32` — 권위 소스 — 에서 기록한
+값이다. `M<TILE_M ∨ K<TILE_K ∨ N<TILE_N`인 shape는 차트에
+("under-tile") 표시된다. `512³` shape는 figure에서 제외된다
+(`EXCLUDED_SHAPES`).
+
+### D5. Theoretical 모델 — 상속된 상수, 아직 ADR-미검증
+
+"theoretical" 곡선은 `scripts/build_overview_slides.py`에서 그대로 복사한
+상수로 해석적 ideal-pipeline 모델을 사용한다:
+
+```
+HBM_GBS = 256.0   # GB/s        T_STAGE = 16.0 ns
+D_STAGES = 3                    BPE = 2
+```
+
+**이 값들은 아직 ADR과 대조 소싱되지 않았다.** 특히 ADR-0033의 `256`은
+`burst_bytes`(256 B)로 이 `256 GB/s`와 *다른* 양이며, ADR-0033은
+대역폭을 `pc_bw_gbs = hbm_to_router_bw_gbs / num_pcs`로 도출한다.
+`T_STAGE`/stage 수도 여기서 ADR-0014로 추적되지 않았다. 따라서 모델은
+**기존 deck script와 일관할 뿐 ADR과 검증되지 않았고**, 상수가 중복된다
+(deck + helper). 이를 조정(topology/ADR-0033/0014에서 소싱, 중복 제거)하는
+것은 보류 — Open questions 참조.
+
+### D6. 알려진 네이밍 caveat — `_measured` 차트
+
+`gemm_mac_utilization_measured.png`는 현재 *theoretical* ideal-pipeline
+수치를 그린다(footnote가 그렇게 명시). 파일명만 "measured"라고 한다. 이는
+그 내용을 시뮬레이터-measured 시리즈로 재지정할지 또는 제목을 바꿀지
+결정을 보류 중인 알려진 misnomer다.
+
+## Consequences
+
+### Positive
+
+- GEMM figure가 allreduce처럼 test 생성·CI 가드된다.
+- 무거운 스윕은 opt-in으로 유지되어 기본 테스트 실행이 빠르다.
+- 스윕 로직의 단일 소스(script)를 slow 테스트가 재사용.
+
+### Negative / limitations
+
+- theoretical 모델 상수(D5)는 미검증·중복이다.
+- `_measured` figure는 misnomer(D6).
+- `build_overview_slides.py`는 여전히 이 PNG를 임베드하지 않고
+  `gemm_sweep.json`에서 GEMM 막대를 네이티브로 그린다 — test 아티팩트를
+  소비하도록 deck를 재배선하는 작업은 미완.
+
+## Dependencies
+
+- **ADR-0013**: verification strategy.
+- **ADR-0014 / ADR-0042**: PE pipeline + tile-plan generator — 스윕이
+  측정하는 GEMM 구현; D4의 stage record count는 ADR-0042 D2/D3에서 온다.
+- **ADR-0033**: latency 모델 — D5 상수가 (아직은 아니지만) 추적되어야 할
+  소스.
+- **ADR-0043**: 형제 격인 allreduce 평가 하니스.
+
+## Open questions
+
+- D5 상수를 `topology.yaml` / ADR-0033 / ADR-0014와 대조 조정하고
+  중복 제거할 것인가(모델 파라미터의 단일 소스)?
+- D6 `_measured` 네이밍 해결(내용 재지정 vs. 제목 변경)?
+- `build_overview_slides.py`를 네이티브 막대 그리기 대신 `gemm_plots/`
+  PNG 임베드로 재배선할 것인가?

docs/adr-ko/ADR-0045-prog-bench-module-contract.md

@@ -0,0 +1,265 @@
+# ADR-0045: Bench Module Contract — registration, dispatch, and authoring
+
+## Status
+
+Accepted (2026-05-21).
+
+`src/kernbench/benches/` 패키지의 등록 메커니즘(@bench), CLI 디스패치 경로
+(`kernbench run/list`), 그리고 새 bench 모듈 작성 시 따라야 할 계약을 통합
+정의한다. ADR-0010 (CLI surface)이 `kernbench list/run` 인터페이스를 명세하나,
+**bench가 어떻게 등록되고 어떤 함수 시그너처를 따라야 하는가**는 ADR 레벨에
+없었음.
+
+**ADR-0054로 확장됨**: D5의 단일 구성 규칙에 세 번째 패턴이 추가된다 —
+*평가 bench*(예: `milestone-1h-*`)는 여러 구성을 구동하고, 구성별 자체 엔진을
+빌드하며, D4를 만족시키기 위해 sentinel 텐서를 제출한다.
+
+## First action (제일 처음에 하는 일)
+
+`kernbench.benches` 패키지가 임포트되면 `__init__.py` 가 즉시
+`_eager_import_and_audit(__path__, __name__)` 를 호출한다. 이 함수의 첫 일은
+패키지 디렉터리 안의 모든 형제 모듈을 `pkgutil.iter_modules(__path__)`로 나열한
+뒤, 다음 두 조건을 만족하지 않는 모듈을 모두 `importlib.import_module(...)`로
+**즉시 로드**하는 것이다:
+
+- 이름이 `registry` 인 경우 (인프라 자체)
+- 이름이 `_` 로 시작하는 경우 (helper 모듈)
+
+임포트 시점에 각 모듈 안의 `@bench(name=..., description=...)` 데코레이터가
+실행되어 `_PENDING` 리스트에 `(name, description, fn)` 튜플이 append 되고,
+`_REGISTERED_MODULES` 셋에 `fn.__module__` 가 추가된다.
+
+전체 임포트가 끝나면 `_audit_modules(imported, _REGISTERED_MODULES)` 가 호출되어,
+**임포트는 되었지만 @bench를 한 번도 호출하지 않은 모듈**이 있으면
+`RuntimeError("Bench module(s) missing @bench decorator: ...")` 가 즉시 발생한다.
+이 audit이 통과한 시점에 인덱스 할당은 아직 일어나지 않은 상태이며, 첫
+`list_all()` / `resolve(...)` 호출 시 `_finalize()` 가 이름 알파벳 정렬 순으로
+1-based index를 부여한다.
+
+즉, **bench 인프라의 첫 일은 "패키지 디렉터리의 모든 비-helper 모듈을 임포트
+하고, 각 모듈이 최소 한 번 @bench를 호출했는지 감사하는 것"** 이다.
+
+## Context
+
+`src/kernbench/benches/` 는 현재 8개의 bench 모듈을 보유한다 (`ccl_allreduce`,
+`gemm_single_pe`, `gpt3_qkv`, `ipcq_allreduce`, `matmul_composite`, `qkv_gemm`,
+`qkv_gemm_multi_pe`, `va_offset_verify`). 모든 bench는 다음 통합 흐름을 따른다:
+
+```
+kernbench run --topology <T> --bench <N>
+   ↓
+cli/main.py::cmd_run
+   ↓  resolve_topology(T)  + resolve(N)  + resolve_device(device_arg)
+   ↓
+runtime_api/bench_runner.py::run_bench(topology, bench_fn, device, engine_factory)
+   ↓  engine_factory(topology, device) → GraphEngine
+   ↓  RuntimeContext(engine, target_device, correlation_id, spec)
+   ↓
+bench_fn(ctx)        ← bench가 정의한 run(torch) 가 호출됨
+   ↓  ctx.empty/zeros/from_numpy/launch/distributed.* 등을 통해 submit
+   ↓
+ctx.wait_all()       ← 미완료 핸들이 있으면 drain
+   ↓
+BenchResult(completion, correlation_id, trace, traces, engine)
+```
+
+ADR-0010 은 CLI 표면만 다루고 (`run/list/probe/web`), ADR-0007 은 runtime API ↔
+sim_engine 책임 경계만 다룬다. 정작 "새 bench 파일을 추가하려면 어떤 모양으로
+써야 하는가"는 코드 컨벤션만으로 추적해야 한다. 결과적으로:
+
+- @bench 데코레이터의 호출 규약 (kebab-case 이름, non-empty description)이
+  코드에만 존재.
+- bench 함수 시그너처 (`def run(torch)`) 가 사실상 컨벤션인데, CLI 디스패치 측이
+  `spec.run` 을 호출한다는 사실로 강제되고 있음.
+- 신규 bench 추가자가 "helper 모듈은 `_` 접두로 분리해야 한다"는 것을 audit
+  RuntimeError를 받아본 뒤에야 학습.
+- single-device 컨벤션 (CLAUDE.md Part 2 CLI Semantics)이 bench 작성자 관점에서
+  어디까지 적용되는지 (CCL 멀티-SIP bench는 예외인가?) 명확하지 않음.
+
+이 ADR이 이런 모호함을 한 곳에 정리한다.
+
+## Decision
+
+### D1. @bench 데코레이터 계약
+
+```python
+from kernbench.benches.registry import bench
+
+@bench(name="my-bench", description="Short, complete-sentence description.")
+def run(torch):
+    ...
+```
+
+- `name`: kebab-case 문자열. 정규식 `^[a-z][a-z0-9]*(-[a-z0-9]+)*$` 통과 필요.
+  소문자/숫자/대시만 허용; 밑줄(`_`) 금지; 알파벳으로 시작.
+- `description`: non-empty 문자열 (strip 후 길이 > 0). CLI `list` 출력에 그대로
+  표시됨.
+- 데코레이터는 **fn을 변형 없이 반환**한다 — 즉 직접 호출도 가능. 부수효과로
+  `_PENDING` 에 등록만 추가한다.
+
+위 두 규칙 위반은 즉시 `ValueError`. duplicate name은 `_finalize()` 시점에
+`RuntimeError("duplicate bench name: ...")` 로 잡힌다.
+
+### D2. 모듈 파일 컨벤션
+
+`src/kernbench/benches/<slug>.py` 는 다음 중 하나여야 한다:
+
+- **bench 모듈**: 최상위 임포트 경로에서 적어도 한 번 `@bench(...)` 가 실행되어
+  최소 하나의 bench를 등록한다.
+- **helper 모듈**: 파일명이 `_` 로 시작 (예: `_shared_helpers.py`). `iter_modules`
+  순회에서 스킵된다.
+
+audit (`_audit_modules`) 는 helper가 아닌데도 @bench를 호출하지 않은 모듈을
+허용하지 않는다. 의도된 결과: 새 파일을 `benches/` 에 추가하기만 하면 자동
+등록되며, helper와의 구분은 **파일명 접두사** 하나로 명확하게 표시된다.
+
+### D3. bench 함수 시그너처는 `def run(torch)` 다
+
+데코레이터는 함수 이름을 강제하지 않지만, **CLI 디스패치는 `spec_entry.run`
+(즉 데코레이트된 callable) 을 호출**한다. 따라서 컨벤션은:
+
+- 함수 이름: `run`. 다른 이름으로 데코레이트해도 동작은 하지만 readability /
+  grep-ability 측면에서 항상 `run`.
+- 인자: 단일 위치 인자 `torch`. 실제로는 `RuntimeContext` 인스턴스이며 PyTorch
+  스타일의 namespace (zeros/empty/launch/distributed/...)를 노출한다 (ADR-0024 D3).
+- 반환값: 임의 (`Any`). 현재 `run_bench` 는 반환값을 무시하고 `ctx.handles()` /
+  `engine.get_completion()` 로 완료를 추적한다.
+
+`torch` 이름은 PyTorch 호환 idiom을 흉내내기 위함이며, 실제로 PyTorch 모듈이
+들어오는 것은 아니다 (ADR-0024 의 "rank = SIP" launcher 컨벤션과 정렬).
+
+### D4. bench는 최소 한 번의 submit을 수행해야 한다
+
+`run_bench` 는 `ctx.handles()` 가 비어 있는 경우 BenchResult.completion 을
+`ok=False, error_code="NO_REQUESTS"` 로 반환한다. 따라서 의미 있는 bench는
+다음 중 하나 이상을 호출해야 한다:
+
+- 텐서 생성 API: `torch.zeros(...)`, `torch.empty(...)` — 내부적으로
+  `MmuMapMsg` 와 (zeros 의 경우) `MemoryWriteMsg` 가 submit 됨.
+- 커널 실행 API: `torch.launch(name, fn, *args)` — `KernelLaunchMsg` 를 SIP 별로
+  submit.
+- (예외) 빈 placeholder bench: `ipcq_allreduce.py` 처럼 `print(...)` 만 하는
+  스텁은 NO_REQUESTS 결과를 받게 됨. CI 측에서 placeholder임을 인지하고 별도
+  처리하는 것을 가정한다.
+
+### D5. 단일-디바이스 컨벤션 + 멀티-SIP 예외 (ADR-0024/0027)
+
+CLAUDE.md Part 2 CLI Semantics 가 명시하는 **"benchmarks MUST remain
+single-device"** 컨벤션은 다음과 같이 해석된다:
+
+- **일반 bench (single-SIP 사용)**: `dp = DPPolicy(...)` 로 텐서 placement를
+  정의하고 `torch.launch(...)` 로 커널 발사. SIP 인덱스는 `--device` 가
+  결정한다 (CLI 측 책임).
+- **CCL bench (멀티-SIP 사용)**: 예외적으로 `torch.distributed.init_process_group
+  (backend="ahbm")` + `torch.multiprocessing.spawn(_worker, ..., nprocs=ws)` 로
+  rank = SIP 패턴 (ADR-0024 D3) 을 따른다. `--device` 는 무시되며 (또는
+  `all` 로 가정), 각 spawned worker가 `torch.ahbm.set_device(rank)` 로 자신의
+  SIP를 바인딩한다.
+
+이 두 패턴 외의 멀티-디바이스 호출 (예: 한 bench 함수가 동일 process에서 여러
+SIP을 직접 launch) 은 본 ADR이 금지한다. CLI 가 `--device all` 로 호출되어도
+bench는 한 번만 실행되며, 그 안에서 멀티-SIP을 다루려면 D5의 두 번째 패턴을
+사용한다.
+
+### D6. 이름·인덱스 해석 (`resolve`)
+
+`resolve(identifier: str)` 는 다음 순서로 BenchSpec을 반환한다:
+
+1. `identifier.isdigit()` → 정수 변환 후 `_REGISTRY` 의 entries에서 `index ==`
+   인 spec 반환. 없으면 `ValueError("No bench with index ..."`)`.
+2. `identifier in _REGISTRY` → 직접 lookup.
+3. 그 외 → `ValueError("Unknown bench ...")`.
+
+빈/공백 identifier 는 `ValueError("bench identifier must be a non-empty string.")`.
+
+CLI 는 `--bench` 의 인자를 그대로 `resolve` 에 넘긴다. 따라서 사용자는
+`kernbench run --bench gemm-single-pe` 또는 `kernbench run --bench 2` 형식 모두
+사용 가능.
+
+### D7. 인덱스는 안정 API가 아니다
+
+`_finalize()` 가 `_PENDING` 을 **이름 알파벳 정렬** 후 1-based index를 부여하므로,
+새 bench 가 추가되면 기존 bench의 index가 밀릴 수 있다. 따라서:
+
+- 사람-친화적 인터랙티브 사용: 인덱스 OK.
+- 스크립트 / CI 자동화: 반드시 이름을 사용한다.
+
+이 사실은 `registry.py` 모듈 docstring 에 명시되어 있다.
+
+### D8. RuntimeContext 가 bench에 노출하는 표면
+
+bench 함수가 `torch` 파라미터를 통해 정상적으로 사용할 수 있는 표면:
+
+- **텐서 생성**: `torch.empty(shape, dtype=..., dp=DPPolicy(...), name=...)`,
+  `torch.zeros(...)`, `torch.from_numpy(arr)`. 모두 host-side 메타 + 디바이스
+  배포 (MmuMap + MemoryWrite) 를 submit 한다.
+- **커널 발사**: `torch.launch(kernel_name, kernel_fn, *args)` —
+  `(Tensor, int, float)` 위치 인자를 `TensorArg` / `ScalarArg` 로 변환하여
+  SIP 별 `KernelLaunchMsg` 발행 후 drain.
+- **동기화**: `torch.wait(handle)`, `torch.wait_all()` (run_bench 가 자동 호출).
+- **분산**: `torch.distributed.init_process_group(backend="ahbm")`,
+  `torch.distributed.get_world_size()`, `torch.distributed.all_reduce(t, op=...)`
+  (ADR-0024/0027).
+- **멀티-프로세스 (rank=SIP)**: `torch.multiprocessing.spawn(_worker, ..., nprocs=ws)`
+  (ADR-0024 D3 / ADR-0027).
+- **디바이스 바인딩**: `torch.ahbm.set_device(rank)` 또는
+  `torch.accelerator.set_device_index(rank)` (둘 다 같은 namespace를 가리킴).
+- **IPCQ 설치**: `torch.install_ipcq(algorithm=..., ccl_yaml=...)` (ADR-0023 D10).
+- **스펙 조회**: `torch.spec` — 토폴로지 빌더가 만든 dict (시스템·cube_mesh·HBM
+  파라미터 등). bench가 toplogy.yaml 파라미터에 의존하지 않게 짜기 위함.
+
+bench는 위에 열거되지 않은 RuntimeContext 의 private 멤버 (`_handles`, `_traces`,
+`_allocators` 등) 에 직접 접근해선 안 된다. ADR-0007 의 layer boundary 정신과
+정렬: bench → runtime API → sim_engine 한 방향만 허용.
+
+### D9. 환경 변수로 파라미터화는 허용된다
+
+`matmul_composite.py` 처럼 `os.environ.get("MATMUL_M", ...)` 등으로 bench
+파라미터를 외부에서 주입하는 패턴은 허용한다. 이유:
+
+- bench 함수 시그너처는 D3 에 의해 `def run(torch)` 로 고정되어 있어 위치/키워드
+  인자로 파라미터를 받기 곤란.
+- 환경 변수 패턴은 `MATMUL_VARIANT` 같은 운영-시 스윕을 위한 자연스러운 hook.
+- `scripts/gemm_sweep.py` 같은 외부 드라이버 (ADR-0044) 가 이 hook을 사용한다.
+
+단, 환경 변수가 bench의 동작을 바꾼다면 모듈 docstring 에 모든 변수를 명시할 것
+(matmul_composite.py 가 그 예시).
+
+## Alternatives Considered
+
+### A1. 명시적 manifest 파일 (YAML)에 bench 목록 두기
+
+기각. @bench 데코레이터 + audit 패턴은 "파일 추가 = 자동 등록" 을 보장하여 신규
+bench 작성자의 인지 비용을 한 곳 (파일 작성)으로 집중시킨다. 별도 manifest는
+유지보수 측에서 drift 위험이 크고, helper 분리는 이미 `_` 접두로 명확하다.
+
+### A2. bench 함수 이름을 데코레이터 인자로 받기 (`@bench(name=..., entry="run_xxx")`)
+
+기각. 디스패치 측에서 `spec.run` 하나만 호출하면 되는 단순함을 깬다. `run` 컨벤션
+하나로 충분하며, 변종이 필요하면 같은 모듈에 여러 함수를 등록하면 된다 (각각
+@bench 데코레이트).
+
+### A3. CCL bench를 위한 별도 `@multi_device_bench` 데코레이터
+
+기각. D5에서 명시한 두 패턴 (single + ADR-0024 멀티-SIP) 만으로 현재 8개 bench가
+모두 표현 가능. 별도 데코레이터는 디스패치 측에서 분기를 강제하여 복잡도를 늘리며,
+멀티-SIP 사용 의도는 bench 함수 본문의 `init_process_group(...)` 호출로 충분히
+드러난다.
+
+### A4. 인덱스를 안정 API로 만들기 (등록 순서 / explicit index= 인자)
+
+기각. D7에서 명시한 trade-off — 사용자 친화성 (알파벳 정렬된 인덱스가 list 출력
+에서 자연스럽게 1, 2, 3...) 우선. 스크립트는 이름으로 지정하면 충분.
+
+## Consequences
+
+- "bench 추가 방법" 이 한 ADR로 정리됨 → 신규 작성자가 코드 grep 없이 D1-D3,
+  D8 만 따르면 됨.
+- helper 모듈을 `_` 접두로 분리하는 패턴이 ADR-level에서 정당화되어, 향후
+  `benches/_*.py` 식의 공유 helper 작성이 자유로워짐.
+- CLAUDE.md Part 2 CLI Semantics 의 single-device 컨벤션이 멀티-SIP CCL bench
+  와 모순되지 않음을 D5 가 명시 — 둘은 직교한다.
+- ADR-0044 (GEMM eval harness) 의 `scripts/gemm_sweep.py` 가 환경 변수 hook을
+  사용하는 근거 (D9) 가 본 ADR에 굳어짐.
+- 인덱스가 불안정함 (D7) 이 명시되어, CI 측 `kernbench run --bench 3` 같은
+  코드는 본 ADR 수락 직후 점검 대상.

docs/adr-ko/ADR-0046-prog-tl-context-contract.md

@@ -0,0 +1,307 @@
+# ADR-0046: TLContext — Kernel-side `tl.*` API Contract
+
+## Status
+
+Accepted (2026-05-22).
+
+`src/kernbench/triton_emu/` 의 `TLContext` 가 노출하는 `tl.*` primitive
+집합과 그 의미, 그리고 두 실행 모드 (command-list / greenlet runner) 의
+계약을 명시한다. ADR-0014/0020 가 PE 파이프라인과 2-pass 실행 모델을
+정의하나, **bench 의 kernel 함수가 호출하는 `tl.*` 표면 자체**는 ADR-level
+에 정리되어 있지 않았다.
+
+## First action (제일 처음에 하는 일)
+
+`TLContext(pe_id, num_programs, dispatch_cycles, runner, cube_id, num_cubes,
+scratch_base, scratch_size)` 생성 시 가장 먼저 다음 6개 필드를 초기화한다:
+
+- `self._pe_id`, `self._num_programs`, `self._cube_id`, `self._num_cubes` —
+  `tl.program_id` / `tl.num_programs` 가 반환할 값.
+- `self._dispatch_cycles` — 모든 `tl.*` API 호출 시작에서 자동으로 발행될
+  `PeCpuOverheadCmd(cycles)` 의 cycle 수.
+- `self._runner` — `KernelRunner` 인스턴스 (있으면 greenlet 모드, 없으면
+  command-list 모드).
+- `self._commands: list[PeCommand] = []` — command-list 모드에서 누적할
+  command 시퀀스.
+- `self._handle_counter = 0`, `self._completion_counter = 0` — 새 TensorHandle /
+  CompletionHandle id 생성용.
+- `self._scratch_base`, `self._scratch_size`, `self._scratch_cursor = 0` —
+  PE-로컬 scratch 영역 (math/dot/composite 의 output handle 주소 할당용).
+
+즉, **TLContext 의 첫 일은 "이 kernel 인스턴스가 어디서 (sip/cube/pe) 어떤
+규모 (num_programs/num_cubes) 로 실행되며, 어느 모드 (runner 유무) 로
+명령을 발사할지 메타데이터를 채우는 것"** 이다. 이 시점에 SimPy event 는
+없으며 command 도 발사되지 않는다.
+
+런타임 첫 동작은 kernel 함수가 `tl.<api>()` 를 처음 호출할 때 발생한다.
+모든 `tl.*` API 의 표준 entry 동작은:
+
+1. `self._emit_dispatch_overhead()` 호출 — `dispatch_cycles > 0` 인 경우
+   `PeCpuOverheadCmd(dispatch_cycles)` 를 즉시 `_emit`.
+2. API 별 처리 (TensorHandle 생성, command 구성).
+3. `self._emit(cmd)` — runner 모드면 greenlet.switch 로 SimPy 측에 cmd 전달,
+   아니면 `self._commands` 에 append.
+
+## Context
+
+`tl.*` 표면은 `TLContext` 가 노출하는 메소드들로 구성되며, kernel 함수가
+받는 `tl` 매개변수가 이 객체다. 사용자(bench 작성자) 입장에서 보이는
+contract:
+
+- 어떤 primitive 가 있는가
+- 각 primitive 가 어떤 데이터 흐름을 발생시키는가 (DMA / compute / IPCQ /
+  metadata-only)
+- TensorHandle 의 `space` 와 `addr` 가 어떻게 결정되는가
+- command-list 모드와 greenlet 모드의 차이
+
+ADR-0014 (PE pipeline) 가 PE_SCHEDULER 가 받는 PeCommand 들을 정의하나,
+`tl.*` 가 이들을 어떻게 emit 하는지는 코드 컨벤션에만 존재한다. 또한
+ADR-0020 (2-pass data execution) 가 greenlet 모드의 존재를 D3 에서
+언급하나, runner / non-runner 두 경로의 시그너처 차이 (return value 처리)
+는 ADR-level 에 명시되어 있지 않다. 이 ADR 이 그 빈자리를 채운다.
+
+## Decision
+
+### D1. `tl` 매개변수는 `TLContext` 인스턴스다
+
+bench 의 kernel 함수는 다음 시그너처를 따른다:
+
+```python
+def _kernel(arg1, arg2, ..., tl, **kwargs):
+    ...
+```
+
+`tl` 의 정체는 `kernbench.triton_emu.tl_context.TLContext` 인스턴스이다.
+real Triton 의 `triton.language` 모듈을 흉내내기 위한 이름이며, real
+Triton 모듈이 들어오는 것은 아니다.
+
+kernel 함수는 일반 Python 함수이며 `yield` / `async` 가 없다. `tl.*`
+호출이 SimPy event 를 발생시키지만, 호출자(kernel) 쪽에서는 동기 호출처럼
+보인다 — greenlet 모드에서 KernelRunner 가 SimPy ↔ kernel 사이를 중계
+하기 때문 (ADR-0020 D3).
+
+### D2. 두 실행 모드 — command-list / greenlet runner
+
+- **command-list 모드 (`runner is None`)**: `tl.*` 호출이 `self._commands`
+  리스트에 PeCommand 를 누적. DMA / GEMM / Math 가 실제 SimPy 시간을
+  소비하지 않으며, return value 가 metadata-only TensorHandle (data=None) 다.
+  이후 PE_SCHEDULER / sim_engine 가 command 시퀀스를 시간상 재생.
+
+- **greenlet runner 모드 (`runner is not None`)**: `tl.*` 호출이
+  `self._emit(cmd)` 를 통해 `runner.switch_to_simpy(cmd)` 로 부모 greenlet
+  (SimPy) 으로 컨트롤을 넘김. 부모는 cmd 를 컴포넌트에 분배하여 SimPy 시간을
+  소비한 뒤, DMA read 의 경우 실제 numpy 데이터를 반환. kernel 은 그
+  결과를 받아 다음 line 으로 진행 (ADR-0020 D3 의 데이터 인지 실행 모델).
+
+mode 선택은 KernelRunner 인스턴스를 TLContext 에 주입하는지 여부로 결정
+되며, `tl.*` 메소드들은 이 차이를 인지하지 않고 `_emit()` 헬퍼를 통해
+일관되게 동작한다.
+
+### D3. Primitive 카테고리
+
+#### D3.1. Reference (no DMA, metadata only)
+
+- `tl.ref(ptr, shape, dtype="f16") -> TensorHandle`: HBM 데이터를 참조하는
+  핸들만 만들고 DMA 는 발행하지 않음. composite scheduler 가 per-tile 로
+  스트리밍할 때 사용 (예: GEMM 의 b 피연산자).
+
+#### D3.2. Data movement (blocking, DMA engine)
+
+- `tl.load(ptr, shape, dtype="f16") -> TensorHandle`: HBM → 결과 핸들.
+  `DmaReadCmd` 발행. greenlet 모드에서는 결과 핸들의 `.data` 에 실제
+  numpy 배열 첨부; command-list 모드에서는 placeholder. 반환 핸들의
+  `space="hbm"`, `pinned=True`.
+- `tl.store(ptr, handle) -> None`: TCM → HBM. `DmaWriteCmd` 발행. greenlet
+  모드에서는 `handle.data` 가 있을 때만 `_store.write("hbm", ptr, data)` 를
+  먼저 호출 (visibility = issue time, ADR-0020 D3).
+
+#### D3.3. GEMM / compute (blocking)
+
+- `tl.dot(a, b) -> TensorHandle`: `a @ b`. 두 피연산자는 TCM 이어야 하며,
+  shape (M,K) × (K,N) → (M,N). `GemmCmd` 발행, output handle 은
+  `_make_compute_out(shape, dtype)` 로 PE-로컬 scratch 에 할당.
+- `tl.composite(op, a, b=None, out_ptr=0, math_op=None, epilogue=None,
+  acc_dtype=None, tile_shape=None) -> CompletionHandle`: 비차단(non-blocking)
+  tiled pipeline. `CompositeCmd` 발행. `epilogue` 는 dict list, 각 dict 는
+  `"op"` 키 + op-specific 필드 + 옵션 `"scope"` (k_tile / output_tile);
+  unknown op 나 missing field 는 즉시 ValueError. 반환된 CompletionHandle 은
+  `tl.wait(h)` 로 동기화.
+
+#### D3.4. Math: unary (blocking)
+
+- `tl.exp(x)`, `tl.log(x)`, `tl.sqrt(x)`, `tl.abs(x)`, `tl.sigmoid(x)`,
+  `tl.cos(x)`, `tl.sin(x)` — 모두 `MathCmd(op=<name>, inputs=(x,), out=)`
+  발행. `out` 은 동일 shape/dtype 의 scratch 할당.
+
+#### D3.5. Math: binary (blocking)
+
+- `tl.maximum(a, b)`, `tl.minimum(a, b)` — `_binary_math`.
+- `tl.fma(a, b, c)` — `a*b + c`. inputs 3개.
+- `tl.clamp(x, min, max)` — `MathCmd(op="clamp", inputs=(x, min, max))`.
+- `tl.where(cond, a, b)` — `MathCmd(op="where", inputs=(cond, a, b))`.
+- `tl.softmax(x, axis=-1)` — 단일 MathCmd(op="softmax") 로 시간 회계는
+  한 번에. Phase 2 DataExecutor 가 canonical (x-max → exp → sum → div) 로
+  expand 한다.
+
+#### D3.6. Reduction (blocking)
+
+- `tl.sum(x, axis)`, `tl.max(x, axis)`, `tl.min(x, axis)` — 해당 axis 의
+  크기를 1 로 줄인 output handle 을 반환. `MathCmd(op=<name>, inputs=(x,),
+  out=, axis=axis)` 발행.
+
+#### D3.7. Index / scalar (PE_CPU, no engine)
+
+- `tl.program_id(axis=0) -> int`: `axis==0` → pe_id (cube-local PE 인덱스),
+  `axis==1` → cube_id (ADR-0022).
+- `tl.num_programs(axis=0) -> int`: `axis==0` → num_programs (cube 당
+  PE 수), `axis==1` → num_cubes.
+- `tl.arange(start, end, dtype="i32") -> TensorHandle`: TCM 의 인덱스
+  range. command 발사 없이 metadata 만.
+- `tl.zeros(shape, dtype="f16") -> TensorHandle`, `tl.full(shape, value,
+  dtype="f16") -> TensorHandle`: TCM 에 placeholder. command 발사 없음.
+
+#### D3.8. Scalar helpers (no command, no engine)
+
+- `TLContext.cdiv(a, b) -> int` (static): ceiling division
+  `-(-a // b)`. real Triton 의 `tl.cdiv` 모방.
+
+#### D3.9. Metadata-only (no compute, no DMA)
+
+- `tl.trans(x) -> TensorHandle`: shape 의 마지막 두 dim 을 swap 한 새
+  핸들. 같은 addr/data 를 공유, command 발사 없음.
+
+#### D3.10. IPCQ (CCL) primitives (ADR-0023 D4)
+
+- `tl.send(dir, src=None, *, src_addr=None, nbytes=None, shape=None,
+  dtype="f16", space="tcm") -> None`: blocking send. handle 형태 또는
+  raw 주소 형태 둘 다 허용. `IpcqSendCmd` 발행. handle 의 `.data` 스냅샷이
+  명령에 실리는 경우, recv 측에서 받은 데이터의 race 회피.
+- `tl.recv(dir=None, shape=(), dtype="f16", space="tcm", dst_addr=None,
+  dst_space=None) -> TensorHandle`: blocking recv. `dst_addr/dst_space`
+  둘 다 주면 "copy_to_dst" 모드, 아니면 "return_slot" 모드. greenlet
+  모드에서 핸들의 `.data` 에 실제 데이터 첨부.
+- `tl.recv_no_consume(dir=None, shape=(), dtype="f16") -> TensorHandle`:
+  **DIAGNOSTIC ONLY**. recv blocking 동기화는 그대로 적용되나 slot-read
+  latency (slot-IO + PE↔bank fabric drain) 는 건너뛴다. pe2pe overview
+  플롯에서 `tl.store` 와의 apples-to-apples 비교용. production kernel 은
+  사용 금지 — `consume=False` 라는 별도 명령 분기로 격리되어 있어 실수
+  flag 가 작동하지 않는다.
+- `tl.recv_async(dir, shape=(), dtype="f16") -> RecvFuture`: non-blocking
+  recv. `RecvFuture` 를 반환; 이후 `tl.wait(future)` 로 결과 수령.
+
+#### D3.11. Composite + control
+
+- `tl.composite(...)`: D3.3 에서 설명.
+- `tl.wait(handle=None)`: `CompletionHandle` (composite) 또는 `RecvFuture`
+  (async recv) 또는 `None` (모든 pending composite) 대기.
+- `tl.cycles(n)`: PE_CPU scalar 실행 overhead 를 명시적으로 선언.
+  `PeCpuOverheadCmd(cycles=n)` 발행.
+
+### D4. TensorHandle 산술 연산자 — thread-local TLContext
+
+`tl_context.py` 모듈 로드 시점에 `_enable_tensor_ops()` 가 호출되어
+`TensorHandle.__add__`, `__sub__`, `__mul__`, `__truediv__` 를 patch한다.
+각 연산자는 thread-local `_ctx` (모듈 변수) 에 저장된 active TLContext 의
+`_binary_math` 를 호출한다.
+
+따라서 kernel 안에서 `c = a + b` 는 `MathCmd(op="add", inputs=(a,b),
+out=)` 발행 + new TensorHandle 반환 패턴과 동일하다.
+
+active TLContext 관리:
+
+- `TLContext._set_active(ctx)`: 현재 thread/greenlet 의 active ctx 설정.
+- `TLContext._get_active()`: 조회 (없으면 RuntimeError).
+- `run_kernel(kernel_fn, tl_ctx, *args, **kwargs)`: helper. 진입 시
+  active 설정, kernel 실행, 종료 시 None 으로 복원.
+
+`KernelRunner` 는 매 cmd 분배 시 `_switch_kernel` 가 직접 `_set_active(tl)`
+를 호출하여, 같은 thread 안의 다른 PE runner 가 active 를 덮어쓴 경우에도
+복원되도록 한다.
+
+### D5. Scratch allocator — compute output handles
+
+`tl.dot`, `tl.exp`, `tl.add` (TensorHandle `__add__`) 등 결과를 만드는 op 는
+`_make_compute_out(shape, dtype)` 를 호출하여 16-byte aligned scratch
+주소를 할당한다. 이 주소는 `space="tcm"` 로 발행되며, 이후 `tl.send` /
+`tl.store` 가 이 handle 을 source 로 사용할 수 있다.
+
+`_scratch_base == 0` (command-list 모드 등) 이면 할당 주소가 0으로
+반환되어 handle 은 send/store 의 source 로 사용 불가 (이 경우 `tl.load`
+로 받은 핸들만 source 가 될 수 있다).
+
+cursor 가 `_scratch_size` (default 1 MiB) 를 초과하면 RuntimeError.
+cursor 는 매 kernel invocation 시작 시 0 으로 리셋되어야 하나 (현재 코드는
+KernelRunner 가 새 TLContext 를 매번 생성하여 자연스럽게 리셋됨).
+
+### D6. Dispatch overhead — `PeCpuOverheadCmd(dispatch_cycles)`
+
+모든 non-metadata `tl.*` 호출의 entry 에서 `_emit_dispatch_overhead()` 가
+호출되며 `dispatch_cycles > 0` 일 때 `PeCpuOverheadCmd(dispatch_cycles)`
+를 발행한다. PE_CPU 가 명령 dispatch 자체에 소비하는 cycle 비용을
+모델링하기 위함이다.
+
+기본값:
+
+- `TLContext.__init__` 의 `dispatch_cycles` 매개변수 기본값: 1 cycle.
+- `KernelRunner` 가 만드는 TLContext: 0 cycles (greenlet 모드는 cycle
+  회계가 별도, ADR-0020 D3 정신).
+
+### D7. Kernel registry (`triton_emu/registry.py`)
+
+별도의 `_kernels: dict[str, Callable]` 가 kernel 이름 → 함수 매핑을 보유:
+
+- `register_kernel(name, fn)`: duplicate 등록 시 ValueError.
+- `get_kernel(name)`: 미등록 시 KeyError.
+- `clear_registry()`: 테스트 전용.
+
+`RuntimeContext.launch(kernel_name, kernel_fn, *args)` 가 매 호출마다
+`_kernels[kernel_name] = kernel_fn` 으로 idempotent 덮어쓴다 (last call
+wins). 이는 ADR-0045 D8 의 launch 동작과 정합된다.
+
+PE_CPU 는 `KernelRef.name` 으로 registry 에서 kernel 함수를 lookup 한 뒤
+KernelRunner 로 실행한다.
+
+## Alternatives Considered
+
+### A1. tl.* 를 ADR-0014 / ADR-0020 안으로 통합
+
+기각. ADR-0014 는 PE pipeline (PeCommand 의 sim_engine 측 소비) 를, ADR-0020
+은 2-pass 실행 (Phase 1 timing / Phase 2 data) 을 다룬다. `tl.*` 는 kernel
+작성자가 만나는 API 표면이라 독립 분리하는 것이 검색성·온보딩 측면에서
+낫다.
+
+### A2. command-list 모드 deprecation
+
+기각 (현재). 단순한 unit test 와 kernel verification 에서 command-list
+모드가 가볍게 동작한다. greenlet 의존성 없이 PeCommand 시퀀스를 검사할 수
+있는 출입구로 유지한다. greenlet 모드만의 의미 (실데이터, Phase 2) 가
+필요하면 D2 의 mode 선택으로 명시적으로 들어간다.
+
+### A3. TensorHandle 산술 연산자 제거
+
+기각. real Triton 의 kernel 코드 가독성을 흉내내기 위함이며 (예: `c = a +
+b`), thread-local active ctx 패턴이 깔끔하게 작동 중. 명시적 `tl.add(a, b)`
+도 D3.5 에 노출되어 있어, 연산자가 헷갈리면 함수형 호출로 대체 가능.
+
+### A4. softmax 를 명시적 시퀀스 (max → exp → sum → div) 로 expand
+
+부분 채택. `tl.softmax` 는 단일 `MathCmd(op="softmax")` 로 timing 회계는
+한 번에 처리한다 (D3.5). 실 데이터 expansion 은 Phase 2 DataExecutor 가
+canonical 시퀀스로 풀어준다. 즉, 시간 모델은 atomic, 데이터 모델은
+expansion — 두 마리 토끼를 의도적으로 분리.
+
+## Consequences
+
+- bench 작성자가 만나는 모든 `tl.*` primitive 가 한 ADR 에 분류·정의됨.
+  ADR-0045 D8 의 host-side surface (torch.empty 등) 와 짝을 이루어 "kernel
+  안 / 밖" 양쪽 작성 가이드가 완성.
+- command-list / greenlet 두 모드의 차이가 D2 에 명시되어, 새로운 `tl.*`
+  primitive 추가 시 `_emit()` 패턴만 따르면 양쪽 자동 호환됨.
+- thread-local active ctx 패턴 (D4) 이 ADR-level 에서 정당화되어, 향후
+  multi-PE 동일-thread 실행 시 reset 책임이 어디인지 명확해짐
+  (`_switch_kernel` 가 cmd 분배 시 active 복원 — KernelRunner.run 의
+  contract).
+- `tl.recv_no_consume` 의 진단 전용 격리(D3.10) 가 ADR 에 굳어져, 실수로
+  production kernel 에서 사용되는 것을 막는 layer 가 명확.
+- registry (D7) 가 별도 D 항목으로 분리되어, kernel 이름 충돌 / 동적
+  재등록 동작의 사양이 명시.

docs/adr-ko/ADR-0047-par-ahbm-ccl-backend.md

@@ -0,0 +1,243 @@
+# ADR-0047: AHBM CCL Backend — `torch.distributed`-compat shim
+
+## Status
+
+Accepted (2026-05-22).
+
+`runtime_api/distributed.py` 의 `AhbmCCLBackend` + `DistributedContext` —
+즉 `torch.distributed.init_process_group(backend="ahbm")` 진입점이 실제로
+무엇을 설치하고 어떤 의미로 `all_reduce`/`barrier`/`get_rank` 등을
+구현하는지를 명시한다. ADR-0023 D11 이 "torch.distributed compatibility"
+의도를 언급하나, **backend 자체의 동작 모델**은 ADR-level 에 없었다.
+
+## First action (제일 처음에 하는 일)
+
+`RuntimeContext.__post_init__` 가 자동으로 `DistributedContext()` 인스턴스를
+만들어 `self.distributed` 에 attach 한다. 그 시점의 첫 일은:
+
+1. `self._backend: AhbmCCLBackend | None = None` 으로 초기화 (아직 init
+   되지 않은 상태).
+2. `self._rank_by_greenlet: dict = {}` 로 greenlet-local rank 레지스트리
+   초기화 (ADR-0024 D2).
+3. 호출자(RuntimeContext) 측에서 `dc._ctx_ref = self` 로 back-reference 를
+   심어, 이후 `init_process_group` 가 `ctx.engine` / `ctx.spec` / `ctx.launch`
+   에 도달할 수 있게 한다.
+
+즉, **DistributedContext 의 첫 일은 "RuntimeContext 에 자기 자신을
+back-reference 와 함께 부착하고 backend 슬롯을 비워두는 것"**. 실제 backend
+설치(IPCQ install, world_size 산출, 알고리즘 모듈 로드)는 사용자 코드의
+`torch.distributed.init_process_group(backend="ahbm")` 호출 시점에 비로소
+일어난다.
+
+해당 시점의 `init_process_group` 의 첫 일은:
+
+1. `backend != "ahbm"` 이면 즉시 `ValueError("Unsupported backend ...")`.
+2. `getattr(self, "_ctx_ref", None)` 가 None 이면
+   `RuntimeError("DistributedContext not bound to a RuntimeContext")`.
+3. `self._backend = AhbmCCLBackend(torch_ctx=ctx)` — 이 생성자 안에서
+   ccl.yaml load + 알고리즘 모듈 import + world_size 산출 + SFR 설정 +
+   IPCQ install 이 모두 일어난다.
+4. `self._backend._dist_ctx = self` — backend 가 거꾸로
+   `_rank_by_greenlet` 에 접근할 수 있게 함.
+
+## Context
+
+PyTorch DDP 의 collective 호출 (`init_process_group`, `all_reduce` 등) 을
+그대로 사용할 수 있게 만들어, bench 코드가 "진짜 DDP training script" 와
+동일한 모습이 되도록 하는 것이 `AhbmCCLBackend` 의 목적이다 (ADR-0024 +
+ADR-0027 의 launcher 모델과 정렬).
+
+이 backend 가 책임지는 것:
+
+- `init_process_group` 시점에 **IPCQ neighbor table 을 한 번 설치** (real
+  NCCL communicator creation 과 유사).
+- `all_reduce(tensor, op="sum")` 호출 시 **설정된 algorithm 의 kernel 함수
+  를 `ctx.launch(...)` 로 발사**.
+- `get_world_size` / `get_rank` 를 greenlet-local rank 레지스트리와
+  ccl.yaml/topology 로부터 일관되게 답함.
+
+ADR-0023 D10 (IPCQ install plan), ADR-0024 (SIP launcher) 가 부분적으로
+이를 다루나, **`AhbmCCLBackend` 자체의 책임 범위와 의사결정 순서**는
+어디에도 명시되어 있지 않다. 본 ADR 이 채운다.
+
+## Decision
+
+### D1. backend 는 `init_process_group(backend="ahbm")` 시점에만 생성된다
+
+`DistributedContext` 는 `__init__` 시점에 `_backend = None` 으로 시작한다.
+backend 객체는 사용자가 `dist.init_process_group(backend="ahbm")` 를
+호출하기 전까지 존재하지 않으며, 그 외 API (`is_initialized`,
+`get_world_size`, `all_reduce`, `barrier`) 가 backend 가 None 인 채로
+호출되면 `RuntimeError("Default process group has not been initialized...")`
+를 던진다 (`_ensure_initialized` 헬퍼).
+
+`backend != "ahbm"` 은 즉시 `ValueError`. 다른 backend 명 (nccl, gloo
+등) 은 인식하지 않는다.
+
+### D2. world_size 산출 우선순위 — algorithm > defaults > topology
+
+`AhbmCCLBackend._resolve_world_size` (ADR-0024 D1) 의 결정 순서:
+
+1. `ccl.yaml` 의 algorithm entry 에 `world_size` 가 있으면 그 값.
+2. `defaults.world_size` 가 있으면 그 값.
+3. 둘 다 없으면 `spec.system.sips.count` (=topology 의 SIP 개수).
+
+기본 의미는 **rank = SIP** (ADR-0024). cube/PE-level parallelism 은 각
+rank 안에서 DPPolicy 로 표현되며 world_size 에 영향을 주지 않는다. 명시적
+`ccl.yaml` 의 world_size override 가 있으면 legacy "rank = flat PE 인덱스"
+테스트 경로를 위해 그대로 존중된다.
+
+`init_process_group(world_size=..., rank=...)` 의 사용자 인자는 **수신하나
+무시**된다 (real PyTorch 의 `RANK` / `WORLD_SIZE` env var 와 같은 의미).
+
+### D3. `init_process_group` 가 즉시 하는 4가지 설치 작업
+
+`AhbmCCLBackend.__init__` 안에서 다음이 순차 실행된다:
+
+1. **ccl.yaml 로딩**: `kernbench.ccl.install.load_ccl_config()` →
+   `resolve_algorithm_config(_cfg_all)` 로 `defaults.algorithm` (또는
+   사용자가 지정한 알고리즘) 의 merged config 산출.
+2. **알고리즘 모듈 import**: `importlib.import_module(self._merged["module"])`.
+   이 모듈은 `kernel` 함수, `kernel_args(world_size, n_elem, cube_w, cube_h)`
+   helper, optional `TOPO_NAME_TO_KIND` 매핑을 노출해야 한다.
+3. **world_size 산출** (D2).
+4. **topology 메타 수집**: `spec` 으로부터 `n_sips`, `sip_topo` (`ring_1d`
+   기본), `cube_w`/`cube_h`, `sips.w`/`sips.h`. SIP topology 가 ring_1d 가
+   아니면 explicit `w`/`h` 또는 square root 로 (`w*h == n_sips` 보장)
+   `_sip_topo_w/h` 산출. 불일치 시 `ValueError`.
+5. **SFR + IPCQ 설치**: `kernbench.ccl.sfr_config.configure_sfr_intercube_multisip
+   (engine, spec, self._merged)` 를 호출. 이 함수가 모든 SIP/cube 의 pe0 에
+   IPCQ neighbor table 을 푸시 (real NCCL communicator 의 일회성 설정에
+   해당).
+
+이 순서가 변하면 (예: SFR 전에 algorithm 모듈 load 가 실패하면) 부분 초기화
+상태가 발생할 수 있다. 따라서 D3 는 atomic 한 4-단계로 본다 — 실패 시
+backend 는 미설치 상태로 남는다.
+
+### D4. greenlet-local rank 등록 (ADR-0024 D2)
+
+`DistributedContext._rank_by_greenlet: dict[greenlet, int]` 은 spawn 된
+worker greenlet 각각에 rank 를 매핑한다. bench launcher (예:
+`torch.multiprocessing.spawn`) 가 worker 를 띄울 때
+`dc._bind_rank(g, rank)` 를 호출하여 등록한다.
+
+`get_rank()` 는 `getcurrent()` 의 greenlet 을 lookup. 미등록 greenlet은
+fallback 으로 0 을 반환 — single-driver / 테스트 호환성 유지.
+
+backend 는 `_dist_ctx._rank_by_greenlet` 를 통해 `all_reduce` 시 현재
+greenlet 의 rank 를 가져온다 (D5).
+
+### D5. `all_reduce(tensor, op="sum")` 동작
+
+검증 단계:
+
+- `op != "sum"` → `NotImplementedError`. 현재 kernel 들은 add reduction만 구현.
+- `tensor._handle is None` → `RuntimeError("not deployed")`.
+- `tensor._handle.shards` 가 비면 `RuntimeError("no shards")`.
+
+준비 단계:
+
+- `n_elem = shards[0].nbytes // tensor.itemsize` — 단일 shard 의 element 수.
+- `kernel_fn = self._algo_module.kernel` — D3 에서 import 된 알고리즘 모듈의
+  진입 함수.
+- effective cube dims 결정: 첫 번째 SIP 의 cube 갯수가 1 이면 (1,1) 으로
+  scalar 처리, 아니면 토폴로지의 `cube_w`/`cube_h` 사용. TP 가 일부 cube
+  만 쓰는 경우를 자연스럽게 흡수.
+- `kernel_args = self._algo_module.kernel_args(world_size, n_elem, cube_w,
+  cube_h)` — 알고리즘이 자기 kernel 에 넘길 인자 셋을 결정.
+
+dispatch:
+
+- 현재 greenlet 의 rank 를 `_rank_by_greenlet.get(g, 0)` 로 lookup.
+- `extra_args = (sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)` 를 append.
+- `pending = self.ctx.launch(algorithm_name, kernel_fn, tensor, *kernel_args,
+  *extra_args, _defer_wait=True)` — `_defer_wait=True` 로 collective drain
+  을 메인 scheduler 에 위임 (ADR-0027 D0.4).
+
+drain:
+
+- 부모 greenlet 이 살아있으면 (multi-greenlet 모드) `_pending_collective_handles`
+  에 enqueue 한 뒤 부모로 switch. 메인 scheduler 가 모든 rank 의 launch 후
+  일괄 drain.
+- 단일-driver 모드면 inline 으로 `for h, _sip_id, meta in pending:
+  self.ctx.wait(h, _meta=meta)` 즉시 drain.
+
+### D6. `barrier()` 는 no-op 이다 (single-driver 모델)
+
+kernbench 는 하나의 Python process 안에서 모든 rank 를 greenlet 으로 다룬다.
+process 간 동기화가 필요한 상황이 없으므로 `barrier()` 는 호출 가능하지만
+실제 어떤 동기화도 수행하지 않는다. real PyTorch DDP 와의 API 호환성을
+위해 유지 (호출자가 NotImplementedError 를 받지 않도록).
+
+장래에 multi-process kernbench (예: SimPy event loop 가 process 별로
+독립) 가 도입되면 D6 를 supersede 하는 새 ADR 이 필요.
+
+### D7. `get_rank` / `get_world_size` / `get_backend` 의 의미
+
+- `get_rank()` (D4): 현재 greenlet 의 bound rank. 미등록은 0.
+- `get_world_size()` (D2): backend 가 D3 에서 산출한 world_size.
+- `get_backend()`: 항상 `"ahbm"` 문자열. backend 객체가 존재하지 않으면
+  `_ensure_initialized` 에서 RuntimeError.
+
+real PyTorch 와의 차이:
+
+- real PyTorch `get_rank()` 는 process global 값이지만, kernbench 는
+  greenlet-local. spawn 된 worker 안에서 호출하면 rank, main thread 에서
+  호출하면 0. bench 작성자는 worker 함수 안에서만 의미 있는 rank 를 기대해야
+  한다.
+
+### D8. 지원하는 API 표면 (final)
+
+`DistributedContext` 가 노출하는 API:
+
+- `init_process_group(backend="ahbm", world_size=None, rank=None, **kwargs)`
+- `is_initialized() -> bool`
+- `get_world_size() -> int`
+- `get_rank() -> int`
+- `get_backend() -> str`
+- `all_reduce(tensor, op="sum") -> None`
+- `barrier() -> None`
+- (internal) `_bind_rank(g, rank)`
+
+이외의 PyTorch distributed API (broadcast, reduce, all_gather, gather,
+scatter, send/recv 등) 는 **아직 구현되어 있지 않다**. kernel 레벨에서는
+`tl.send`/`tl.recv` (ADR-0046 D3.10) 로 직접 표현 가능하나, dist.* surface
+로는 노출되지 않는다. 추가 collective 가 필요해질 시 별도 알고리즘 모듈
++ `DistributedContext` 메소드 한 쌍을 추가하여 D8 를 확장한다.
+
+## Alternatives Considered
+
+### A1. backend 를 `RuntimeContext.__init__` 에서 즉시 생성
+
+기각. ccl.yaml 이 없거나 알고리즘 모듈을 import 할 수 없는 경우, bench 가
+distributed 기능을 안 쓰는데도 RuntimeContext 생성 자체가 실패하게 된다.
+"호출 시점에 비로소 설치" (D1) 가 lazy 의미상 옳다.
+
+### A2. world_size 를 항상 topology 로부터 자동 산출 (override 금지)
+
+기각. ADR-0024 D1 의 "explicit override" 경로가 legacy 테스트에서 사용 중.
+한 SIP 안에서 PE-level rank 를 따로 정의해야 하는 진단 시나리오를 위해
+유지.
+
+### A3. `op != "sum"` 을 silent fallback 으로 처리
+
+기각. 사용자가 `op="prod"` / `"max"` / `"avg"` 를 의도했는데 silently sum
+이 실행되면 결과 검증이 매우 어렵다. 명시적 `NotImplementedError` 가 안전.
+
+### A4. `barrier` 를 SimPy event 로 구현
+
+기각 (현재). single-driver 모델에서 cross-process 동기화 의미가 없으므로
+no-op 가 의미적으로 정확. SimPy fake-barrier 는 의미 없이 코드 복잡도만
+높임. multi-process kernbench 도입 시 재평가.
+
+## Consequences
+
+- `torch.distributed.init_process_group(backend="ahbm")` 의 4-단계 설치
+  (D3) 가 ADR-level 에서 굳어져, 향후 새 collective 알고리즘이 어디에
+  훅을 걸어야 하는지 명확.
+- D2 의 우선순위 (algorithm > defaults > topology) 가 명시되어, ccl.yaml
+  변경 시 영향 범위를 빠르게 가늠 가능.
+- D6 의 barrier no-op 결정이 ADR-level 에 굳어져, multi-process kernbench
+  도입 시 별도 ADR 로 supersede 해야 함이 분명.
+- D8 의 미지원 API 목록이 명시되어, 사용자가 `dist.broadcast(...)` 를
+  호출하려 할 때의 명확한 거절 근거 제공.

docs/adr-ko/ADR-0048-mem-allocator-algorithms.md

@@ -0,0 +1,262 @@
+# ADR-0048: Memory Allocator Algorithms — VirtualAllocator + PEMemAllocator
+
+## Status
+
+Accepted (2026-05-22).
+
+`policy/address/allocator.py` 의 `_FreeList` / `PEMemAllocator` 와
+`va_allocator.py` 의 `VirtualAllocator` 가 사용하는 free-list 알고리즘,
+페이지 정렬, coalescing 규칙을 명시한다. ADR-0001 (PhysAddr 레이아웃) 과
+ADR-0011 (PA/VA/LA 모델) 이 주소 스킴을 정의하나, **할당 알고리즘**은 별도
+ADR 이 없었다.
+
+## First action (제일 처음에 하는 일)
+
+### `_FreeList(capacity)`
+
+생성 즉시 `self._capacity = capacity`, `self._used = 0`, `self._free =
+[(0, capacity)]` 로 초기화. 첫 일은 **전 영역을 single free block 으로
+세우는 것** — 즉 `(offset=0, size=capacity)` 한 튜플이 free list 의 유일한
+원소다.
+
+### `PEMemAllocator(sip_id, die_id, pe_id, cfg)`
+
+생성 즉시 두 개의 `_FreeList` 를 만든다:
+
+- `self._hbm = _FreeList(cfg.hbm_slice_bytes)` — 이 PE 가 소유한 HBM
+  slice 의 바이트 크기 (`hbm_bytes_per_cube // hbm_slices_per_cube`) 만큼.
+- `self._tcm = _FreeList(cfg.tcm_allocatable_bytes)` — `tcm_bytes_per_pe -
+  tcm_scheduler_reserved_bytes` 만큼 (scheduler 예약분은 사전 분리).
+
+따라서 PEMemAllocator 의 첫 일은 **이 PE 의 HBM slice 와 사용자
+TCM 영역을 각각 단일 free block 으로 세우는 것**.
+
+### `VirtualAllocator(va_base, va_size, page_size=2*1024*1024)`
+
+생성 즉시 `self._va_base = va_base`, `self._va_size = va_size`,
+`self._page_size = page_size`, `self._used = 0`, `self._free = [(va_base,
+va_size)]`. 첫 일은 **VA base 부터 size 까지 single block 으로 세우고
+page_size 를 회수**.
+
+## Context
+
+`runtime_api/context.py::_ensure_allocators` 는 다음 단계로 allocator 세트를
+구성한다:
+
+1. spec 으로부터 `hbm_total_gb_per_cube`, `hbm_slices_per_cube`,
+   `tcm_size_mb`, target_device 별 SIP 범위 등을 읽음.
+2. `AddressConfig` 로 모든 파라미터를 frozen 하게 패킹.
+3. target SIP 범위 × cube × PE 의 모든 조합에 대해
+   `PEMemAllocator(sip, cube, pe, cfg)` 인스턴스를 1개씩 생성.
+4. `VirtualAllocator(va_base=0x1_0000_0000, va_size=64 GiB,
+   page_size=pe_mmu.page_size)` 를 1개 생성.
+
+allocator 들의 책임:
+
+- **PEMemAllocator**: PE-로컬 HBM slice / TCM 의 PA-공간 할당 (PhysAddr
+  encoding 까지 포함).
+- **VirtualAllocator**: device-wide VA 공간을 페이지 정렬로 할당. 이후
+  `RuntimeContext._create_tensor` 가 VA → PA 매핑을 `MmuMapMsg` 로 fabric
+  에 push.
+
+이 알고리즘들은:
+
+- **first-fit** 으로 단순.
+- 자유 블록 리스트는 **offset 정렬 (sorted by start)** 유지.
+- `free()` 시 **양쪽 인접 블록과 coalesce**.
+
+이런 결정의 근거가 어디에도 없으므로, 향후 누군가 "왜 best-fit 이 아닌가",
+"왜 buddy allocator 가 아닌가", "왜 partial overlap free 가 silently
+허용되는가" 라는 질문에 답할 기준이 필요. 본 ADR 이 그 기준을 마련한다.
+
+## Decision
+
+### D1. `_FreeList` — offset-기반 first-fit + coalescing
+
+`policy/address/allocator.py::_FreeList`:
+
+- 내부 표현: `list[tuple[int, int]]` = `[(start_offset, size), ...]` —
+  start offset 으로 정렬된 자유 블록의 sorted list.
+- `alloc(nbytes)`:
+  1. free list 를 앞에서부터 순회 (first-fit).
+  2. 처음 만나는 `size >= nbytes` 인 블록에서 앞부분을 잘라 사용.
+  3. 정확히 일치하면 블록 통째로 제거; 아니면 `(start+nbytes, size-nbytes)`
+     로 축소.
+  4. `_used += nbytes`, 잘라낸 `start` 반환.
+  5. 맞는 블록이 없으면 `AllocationError("overflow ... largest free block
+     ...")`.
+- `free(offset, nbytes)`:
+  1. `_used -= nbytes`.
+  2. `bisect_left(self._free, (offset,))` 로 삽입 위치 결정.
+  3. 직전 블록과 인접 (`prev_start + prev_size == offset`) 하면 흡수.
+  4. 직후 블록과 인접 (`offset+nbytes == next_start`) 하면 흡수.
+  5. coalesced range 를 정렬 위치에 insert.
+
+이 알고리즘은 fragmentation 에 약점이 있으나 (best-fit / buddy 대비), 본
+시뮬레이터의 워크로드 특성상 (deploy/free 패턴이 거의 stack-like) 충분
+하다는 것이 디자인 가정이다. 워크로드가 변하면 D1 supersede 후보.
+
+### D2. partial overlap free 는 **검사하지 않는다**
+
+`_FreeList.free(offset, nbytes)` 는 호출자가 정확한 (offset, nbytes) 를
+넘긴다고 신뢰한다. 다음을 검증하지 않는다:
+
+- 그 range 가 실제로 alloc 된 것인지.
+- 그 range 가 다른 alloc 된 영역과 겹치지 않는지.
+
+이유: 시뮬레이터 컨텍스트에서 호출자는 항상 `alloc()` 의 반환값을 그대로
+저장했다가 `free()` 에 넘기는 패턴이며, 외부 사용자 입력이 아니다. 안전성
+검사를 추가하면 매 free 마다 O(N) 비용이 들어 시뮬 wall-clock 에 영향.
+
+이 신뢰 모델이 깨지면 (예: 두 텐서가 같은 PA 를 가리키는 코드 경로 도입)
+즉시 ADR-level 으로 재검토.
+
+### D3. `PEMemAllocator` — HBM/TCM 두 채널 분리
+
+`PEMemAllocator(sip_id, die_id, pe_id, cfg)` 는 두 `_FreeList` 를 보유:
+
+- `_hbm`: `cfg.hbm_slice_bytes` 크기.
+- `_tcm`: `cfg.tcm_allocatable_bytes` (= `tcm_bytes_per_pe -
+  tcm_scheduler_reserved_bytes`) 크기.
+
+`alloc_hbm(nbytes) -> PhysAddr`:
+
+- `_hbm.alloc(nbytes)` 로 offset 획득.
+- `PhysAddr.pe_hbm_addr(sip_id, die_id, pe_id, pe_local_hbm_offset=offset,
+  slice_size_bytes=cfg.hbm_slice_bytes)` 로 PA 인코딩.
+- 실패 시 `AllocationError("HBM overflow ...")`.
+
+`free_hbm(pa, nbytes)`:
+
+- `pa.hbm_offset - pe_id * cfg.hbm_slice_bytes` 로 PE-local offset 복원.
+- `_hbm.free(offset, nbytes)`.
+
+`alloc_tcm(nbytes) -> PhysAddr`: 유사하게 `PhysAddr.pe_tcm_addr` 로 인코딩.
+
+`free_tcm(pa, nbytes)`: `pa.sub_offset` 을 그대로 사용 (TCM 은 PE-local
+offset 이 곧 sub_offset).
+
+scheduler-reserved TCM 영역 (`cfg.tcm_scheduler_reserved_bytes`) 은
+allocator 가 인지하지 않는다 (`_tcm` 의 capacity 에서 사전 차감되어 있음).
+이는 ADR-0014 의 PE_SCHEDULER 내부 buffer 예약과 정합된다.
+
+### D4. `VirtualAllocator` — 페이지 정렬 first-fit + coalescing
+
+`policy/address/va_allocator.py::VirtualAllocator`:
+
+- 내부 표현: `_FreeList` 와 동일한 sorted `list[tuple[int, int]]`.
+  최초: `[(va_base, va_size)]`.
+- `_align_up(nbytes) = ceil(nbytes / page_size) * page_size`.
+- `alloc(nbytes) -> int`:
+  1. `aligned = _align_up(nbytes)`.
+  2. first-fit 으로 `size >= aligned` 인 블록 탐색.
+  3. 블록 앞부분 `aligned` 만큼 잘라 사용. 정확히 일치하면 제거.
+  4. `_used += aligned`. 블록 `start` (= aligned 된 VA) 반환.
+  5. 실패 시 `VaAllocationError`.
+- `free(va, nbytes)`: `_align_up(nbytes)` 단위로 free. _FreeList 와 동일한
+  coalesce 알고리즘.
+
+`page_size` 의 실제 값은 두 곳에서 다른 기본을 갖는다:
+
+- `VirtualAllocator.__init__` 의 매개변수 기본값: `2 MiB`. 직접 호출하는
+  테스트가 그대로 받는다.
+- `RuntimeContext._ensure_allocators` 가 인스턴스화할 때:
+  `pe_mmu.attrs.get("page_size", 4096)` — `topology.yaml` 의
+  `pe_mmu.attrs.page_size` 가 있으면 그 값, 없으면 fallback 4 KiB.
+
+두 기본이 다른 이유: VirtualAllocator 의 standalone 기본은 ADR-0039 의
+PE_MMU stopgap 기본 (2 MiB) 과 정합되어 직접 테스트가 자연스럽고, context
+fallback 의 4 KiB 는 topology 미설정 시 안전한 minimum page 다. 실제 사용
+경로는 항상 후자이며 (`_ensure_allocators` 가 인스턴스화하므로),
+`topology.yaml` 에서 `page_size` 가 명시되면 그 값이 양쪽 (MMU + VA
+allocator) 으로 일관되게 흐른다.
+
+만약 이 일치가 깨지면 (예: VirtualAllocator 의 page_size 를 PE_MMU 와
+다르게 인스턴스화) MMU `map()` 가 서브-페이지 region 모드 (ADR-0039 D3) 로
+흐른다.
+
+VA 기본 범위: `va_base = 0x1_0000_0000` (= 4 GiB), `va_size = 64 GiB`. 이
+값은 `_ensure_allocators` 에 하드코딩되어 있으며 ADR-0011 의 VA 모델에서
+직접적인 의미를 갖지는 않는다 — 단지 host 코드와 충돌하지 않을 만큼 큰
+주소 공간을 device-wide 로 잡아둔 것.
+
+### D5. allocator 인스턴스의 lifecycle
+
+- `RuntimeContext._ensure_allocators` 가 lazy 하게 호출됨 (`_create_tensor`
+  의 첫 호출 시점).
+- 한 번 생성된 allocator dict (`self._allocators`) 는 RuntimeContext 의
+  lifetime 동안 재사용. 같은 process 안의 두 번째 deploy 는 새 객체를
+  만들지 않는다.
+- `RuntimeContext.cleanup()` 이 모든 living tensor 의 `_free_tensor()` 를
+  호출 → MMU unmap + `va_allocator.free` + `pemem_allocator.free_hbm` 으로
+  free list 가 원상복구. 다음 RuntimeContext 가 다시 만들면 초기 상태부터.
+
+allocator 상태가 RuntimeContext 간에 공유되지 않는 점이 단일 process 안의
+연속 실행에서 deploy → cleanup → deploy 의 결정성을 보장한다.
+
+### D6. Allocator 실패는 raise 한다 (silent OOM 금지)
+
+`_FreeList.alloc` / `VirtualAllocator.alloc` 모두 충분한 free block 이
+없으면 `AllocationError` / `VaAllocationError` 를 던진다. 메시지에는
+"required size + largest available block" 가 포함되어, fragmentation
+인지 진짜 OOM 인지 진단 가능.
+
+silent fallback (예: 가장 큰 블록만큼만 alloc) 는 절대 금지 — 부분 할당된
+텐서가 SimPy 단계에 들어가면 라우팅·DMA 가 잘못된 PA 를 인지하여 시뮬
+정확도가 깨진다.
+
+### D7. address space 와 allocator 의 1:1 대응
+
+물리 주소 공간 분리는 PhysAddr 의 sub-unit (ADR-0001 D2.3) 으로 표현되며,
+각 sub-unit 마다 별도 allocator 인스턴스를 둔다:
+
+- HBM slice → `PEMemAllocator._hbm`.
+- PE TCM → `PEMemAllocator._tcm`.
+- (현재 미사용) M_CPU local memory, CUBE SRAM → 별도 allocator 필요. 현재
+  구현은 아직 IPCQ-only slot 으로 처리 (ADR-0023 D9.7) 하며 PA 공간을
+  share 하지 않으므로 별도 free-list 가 없음.
+
+cube-level SRAM allocator 가 필요해지면 `_FreeList(cfg.sram_bytes_per_cube)`
+인스턴스를 cube 단위로 추가한다 (`cfg.sram_bytes_per_cube` 는 이미
+`AddressConfig` 에 정의되어 있어 데이터 모델은 준비됨).
+
+## Alternatives Considered
+
+### A1. best-fit / buddy allocator
+
+기각 (현재). 워크로드의 alloc/free 패턴이 stack-like (deploy 순서 = free
+순서) 라 first-fit + coalescing 으로 fragmentation 이 충분히 통제된다.
+LLM kernel sweep 에서 long-running fragmentation 이 관찰되면 buddy 로
+교체하는 ADR 을 별도로 만든다.
+
+### A2. partial overlap free 검증 추가
+
+기각. D2 의 신뢰 모델 + O(N) 검사 비용. 단, 디버그 모드 (`KERNBENCH_DEBUG`
+env var 등) 에서 활성화하는 옵션은 후속 작업으로 가능.
+
+### A3. VA 와 PA 의 통합 allocator
+
+기각. VA 공간 (64 GiB device-wide) 과 PA 공간 (slice 별 ~6 GiB) 는 의미
+차원이 다르다. VA 는 host kernel 의 view, PA 는 device sub-unit 의 view.
+ADR-0011 의 VA 모델 정신 (MMU 가 둘 사이를 매핑) 과 정합하기 위해
+allocator 도 분리.
+
+### A4. page_size 의 multi-tier 지원 (large page + small page)
+
+기각 (현재). 단일 page_size (현재 2 MiB) 가 LLM kernel 의 텐서 단위 (수
+MiB~수 GiB) 에 맞고, ADR-0039 D3 의 서브-페이지 region 으로 작은 매핑이
+필요할 때 흡수된다. multi-tier page 는 MMU 자체 모델을 확장해야 하므로
+별도 ADR 후보.
+
+## Consequences
+
+- allocator 알고리즘이 ADR-level 에서 굳어져 (D1·D3·D4), 새로운 시뮬
+  시나리오에서 fragmentation 이슈가 발생할 때 "여기서 first-fit + coalesce
+  를 쓰고 있다" 가 명확.
+- D2 의 신뢰 모델이 명시되어, 향후 사용자 입력으로부터 직접 alloc/free 를
+  받는 경로가 도입되면 본 ADR supersede 가 필요함을 일찍 인지 가능.
+- D7 의 sub-unit별 allocator 1:1 대응이 명시되어, M_CPU/SRAM 별도 영역이
+  필요해질 때 어디에 free-list 를 추가해야 하는지 명확.
+- `VirtualAllocator` 의 page_size 가 PE_MMU 설정과 일치해야 함이 D4 에
+  적혀 있어, 향후 topology.yaml 의 page_size 변경 시 ADR-0039 stopgap 동작
+  과의 상호작용을 빠르게 가늠 가능.

docs/adr-ko/ADR-0049-ver-probe-subcommand.md

@@ -0,0 +1,231 @@
+# ADR-0049: `kernbench probe` Subcommand — Traffic-Pattern Verification Harness
+
+## Status
+
+Accepted (2026-05-22).
+
+`probes/probe.py` 의 `run_probe(...)` 가 노출하는 traffic-pattern catalog,
+formula vs actual 비교, 그리고 monotonicity / D2H≥H2D 같은 invariant
+체크의 의미를 명시한다. ADR-0010 (CLI surface) 가 `kernbench probe`
+subcommand 를 enumerate 하나, **probe 가 실제로 측정하는 것**과 **어떤
+invariant 를 PASS/FAIL 로 판정하는가**는 ADR-level 에 없었다.
+
+## First action (제일 처음에 하는 일)
+
+`run_probe(topology_path, case_filter=None)` 의 첫 4가지 작업:
+
+1. `Path(topology_path).expanduser().resolve()` 로 절대 경로 산출.
+2. `load_topology(path)` → `TopologyGraph` 인스턴스 (그래프 + spec).
+3. `_build_edge_map(graph)` → `{(src, dst): Edge}` 빠른 lookup 테이블.
+4. `AddressResolver(graph)` + `PathRouter(graph)` 인스턴스화.
+
+그 다음 `nbytes = 32768` (= 32 KiB, summary table 의 기준 데이터 크기) 와
+`show_all = (case_filter is None or case_filter == "all")` 를 설정.
+
+즉, **probe 의 첫 일은 "토폴로지를 한 번 로드하여 edge map / resolver /
+router 를 준비하고, 32 KiB 라는 표준 측정 크기를 픽스하는 것"**. 그 이후
+H2D → D2H → PE DMA 세 카테고리의 case 들이 각각 별도의 `GraphEngine`
+인스턴스에서 실행된다 (case 간 cross-talk 차단).
+
+## Context
+
+`kernbench probe` 는 다음 의도로 도입된 verification 도구다:
+
+- **수동 분석 ground truth**: 실 시뮬레이션 (`kernbench run --bench ...`)
+  결과의 latency 가 비정상으로 보일 때, 단순 traffic pattern 의 정답을 별도
+  로 얻어 비교.
+- **formula vs actual 비교**: 분석 모델 (wire latency + overhead + drain)
+  과 시뮬레이션 결과 (`total_ns`) 가 일치하는지 확인. 일치하지 않으면 모델
+  단순화 가정 (ADR-0033) 어디가 빠진 것인지 단서.
+- **monotonicity check**: hop 수가 늘면 latency 가 단조 증가해야 한다는
+  invariant 의 자동 확인.
+- **utilization sweep**: 데이터 크기 (4 KiB ~ 1 MiB) 별 BW 활용률 표.
+
+이 도구의 동작 사양이 ADR-level 에 없으면:
+
+- 다른 형식의 traffic pattern (예: MCpuDma, IPCQ) 을 추가하려는 사람이 기존
+  카테고리의 표 포맷 / 측정 단위를 일관되게 따르기 어렵다.
+- monotonicity 가 무엇을 기준으로 검사되는지 (hop 수? cube 거리? wire
+  길이?) 모호.
+- 32 KiB 라는 기준 크기와 `[4 KiB, 16 KiB, 64 KiB, 256 KiB, 1 MiB]` sweep
+  의 의미가 코드 grep 으로만 확인 가능.
+
+## Decision
+
+### D1. 세 가지 case category — H2D / D2H / PE DMA
+
+각 category 는 토폴로지 상 별개의 데이터 경로를 가지며, 별도의 summary
+table + sweep table + route detail block 으로 출력된다.
+
+- **H2D (Host→Device Write)**: `MemoryWriteMsg(dst_sip=0, dst_cube,
+  dst_pe=0, pattern="zero")` 가 `pcie_ep → io_cpu → m_cpu → hbm_ctrl` 경로
+  를 흐른다. cube 인덱스로 hop 수가 증가:
+  - h2d-1hop: cube=0, hops=1
+  - h2d-2hop: cube=4, hops=2
+  - h2d-3hop: cube=8, hops=3
+  - h2d-4hop: cube=12, hops=4
+- **D2H (Device→Host Read)**: `MemoryReadMsg(src_sip=0, src_cube, src_pe=0)`.
+  forward command path + reverse data path 의 합 latency. 같은 4 hops
+  카테고리.
+- **PE DMA (PE-initiated)**: `PeDmaMsg(src_sip, src_cube, src_pe, dst_pa)`.
+  5 가지 케이스로 cube/PE 위치 변화:
+  - pe-local-hbm: same cube, same PE
+  - pe-same-half-hbm: same cube, different PE (PE 1)
+  - pe-cross-half-hbm: same cube, far PE (PE 4)
+  - pe-cross-cube-hbm-best: adjacent cube (cube 1)
+  - pe-cross-cube-hbm-worst: diagonal far cube (cube 15)
+
+cube 인덱스가 4/8/12 (H2D), 1/4/15 (PE DMA) 같이 의미 있는 이유는
+4x4 cube mesh (sip.cube_mesh.w=4, h=4) 에서의 거리 정의 — 추후 cube_mesh
+크기 변경 시 이 값들이 같이 갱신되어야 한다.
+
+### D2. 표준 측정 크기 — `nbytes = 32768` (32 KiB)
+
+모든 case 의 summary table 은 `nbytes=32768` 로 한 번 실행한 결과를
+보여준다. 32 KiB 가 선택된 이유:
+
+- DMA overhead 와 BW drain 이 한쪽으로 치우치지 않는 적당한 크기.
+- 다수 sub-unit (TCM, register file) 의 1회 transfer 단위와 비교 가능.
+
+크기별 utilization 변화는 별도 sweep table 이 보여준다 (D3).
+
+### D3. Utilization sweep — `[4 KiB, 16 KiB, 64 KiB, 256 KiB, 1 MiB]`
+
+`SWEEP_SIZES = [4096, 16384, 65536, 262144, 1048576]`, `SWEEP_LABELS =
+["4KB", "16KB", "64KB", "256KB", "1MB"]`. 매 size 마다 다음 공식:
+
+```
+drain   = nbytes / bottleneck_bw
+total   = overhead + wire + drain
+eff_bw  = nbytes / total
+util%   = eff_bw / bottleneck_bw × 100
+```
+
+`bn_bw is None or <= 0` 이면 그 컬럼은 0.0 % 로 출력. 의미: hop 수가 늘
+수록 작은 transfer 는 overhead-bound, 큰 transfer 는 drain-bound 가 되는
+패턴을 한 표에서 확인.
+
+### D4. 측정 항목 — actual / formula / breakdown
+
+각 case 행에 표시되는 컬럼:
+
+- `Actual` (total_ns): SimPy 실행 결과의 `trace["total_ns"]`.
+- `Ovhd`: 경로상 모든 node 의 `node.attrs["overhead_ns"]` 합 (formula
+  breakdown).
+- `Drain`: `nbytes / min(edge.bw_gbs over path)` (formula).
+- `Wire`: `Σ edge.distance_mm * (ns_per_mm from spec)`.
+- `Ovhd%` / `Drain%`: Ovhd/Drain 이 Actual 에서 차지하는 비율 (formula 의
+  Wire 는 통상 매우 작아 표시하지 않음).
+- `Eff.BW`: `nbytes / total_ns` (실 측정 BW).
+- `BN.BW`: bottleneck bandwidth (formula). path 상 모든 edge 의 BW 중 최소.
+  edge BW 가 없으면 "-".
+- `Util%`: `Eff.BW / BN.BW × 100`. 100% 면 single-stream BW upper bound 에
+  도달.
+
+formula 의 합 (`wire + ovhd + drain`) 과 actual 의 차이가 크면 모델
+단순화가 잡지 못하는 요소가 있다는 신호 (ADR-0033 의 가정 점검).
+
+### D5. Invariant 자동 체크 — PASS/FAIL
+
+다음 invariant 들이 자동으로 확인되어 `[v] PASS` / `[x] FAIL` 로 출력:
+
+- **H2D / D2H monotonic increase**: hop 수가 늘면 actual latency 가
+  단조 증가해야 함. `all(lats[i] < lats[i+1] for ...)`.
+- **D2H ≥ H2D**: 같은 hop 인덱스에서 D2H ≥ H2D (D2H 는 forward command
+  + reverse data 두 leg 이므로). `all(d2h[i].total >= h2d[i].total)`.
+- **PE DMA best < worst**: cross-cube best (adjacent) latency < cross-cube
+  worst (diagonal) latency.
+- **PE DMA local vs remote**: local BN BW vs remote BN BW 의 비교 출력
+  (PASS/FAIL 이 아닌 정보성).
+
+체크가 FAIL 이면 사람이 즉시 모델/토폴로지 회귀를 인지할 수 있도록 한
+줄로 분명하게 출력.
+
+### D6. Route detail — per-hop timestamp trace
+
+summary 와 sweep 표 이후 각 case 의 path 와 per-hop 누적 시간 (
+`_hop_timestamps`) 가 별도 섹션에서 출력된다:
+
+- H2D: leg1 (`pcie_ep → io_cpu`) + leg2 (`io_cpu → m_cpu`) + leg3
+  (`m_cpu → hbm_ctrl`) + per-hop trace.
+- D2H: forward (cmd, no data) + reverse (data) trace 분리 표시.
+- PE DMA: `pe_dma → router → hbm_ctrl` path + per-hop trace.
+
+각 hop 의 timestamp 는 cumulative `wire_ns + overhead_ns` 누적. terminal
+hop 의 annotation 에 `drain:Xns` 가 붙는다. bottleneck edge 는
+`<BN:XXGB/s>` 로 표시되어 시각적으로 식별 가능.
+
+### D7. case_filter 인자의 의미
+
+- `None` 또는 `"all"`: 모든 case 실행 (default).
+- 다른 문자열: 그 이름과 정확히 일치하는 case 만 실행. 예: `kernbench
+  probe --case h2d-2hop`.
+
+각 카테고리 안에서 `name != case_filter` 면 skip 되며, 그 카테고리의
+monotonicity / D2H≥H2D 비교는 데이터가 1개일 때 자연히 skip 된다.
+
+CLI parser 의 `--case` 기본값은 `"all"`이라 인자 생략 시 전체 실행.
+
+### D8. 매 case 별 fresh GraphEngine
+
+H2D 4개, D2H 4개, PE DMA 5개의 case 가 각각 **새로운 GraphEngine**
+인스턴스에서 실행된다 (`engine = GraphEngine(graph)`). 이유:
+
+- case 간 누적 상태 (op_log, completion 추적, allocator 등) 가 cross-talk
+  하지 않도록 격리.
+- 한 case 의 traffic 이 다른 case 의 BW 측정에 영향을 주지 않도록 보장.
+
+이 격리는 probe 의 측정 결과를 **각 case 단독 single-flow** 의 latency 로
+해석할 수 있게 한다. multi-flow contention 측정은 별도 도구 (예:
+`pe2pe_overview` 플롯, ADR-0033 의 multi-flow merging 모델) 책임.
+
+### D9. 출력 포맷의 안정성
+
+probe 의 stdout 출력은 사람이 읽기 위함이며, 정확한 컬럼 폭/구분자/공백 은
+machine-readable contract 가 아니다. 자동화된 도구가 probe 결과를 파싱
+하려면 별도 JSON 출력 모드를 추가해야 한다 (현재 미구현).
+
+PASS/FAIL 줄의 `[v]` / `[x]` 접두사는 CI grep 용 anchor 로 안정 보장.
+
+## Alternatives Considered
+
+### A1. Probe 를 별도 bench 로 등록 (`@bench(name="probe")`)
+
+기각. probe 는 bench 가 아니라 verification 도구로 의도된다 — sweep / 분석
+용 multi-engine 실행과 invariant PASS/FAIL 출력이 본질이며, ADR-0045 의
+"단일 디바이스 + 단일 RuntimeContext" bench 모델과 맞지 않는다.
+
+### A2. monotonicity 위반 시 exit code 1
+
+기각 (현재). 인간 검사 도구 위주로 의도되어 있어 PASS/FAIL 줄을 출력하고
+exit 0 로 종료. CI 가 violation 으로 fail 하길 원하면 별도 wrapper 가
+`grep "\[x\]"` 결과로 판단하면 됨. 후속으로 strict-mode flag (`--strict`)
+도입 가능.
+
+### A3. probe 의 case 정의를 외부 YAML 로
+
+기각 (현재). 8개 case (4 H2D + 4 D2H + 5 PE DMA — 합 13개) 는 코드에
+하드코딩되어 있고 의미가 토폴로지 mesh 구조에 단단히 묶여 있다. 외부
+YAML 로 옮기면 cube 인덱스의 의미 (4, 8, 12 / 1, 4, 15) 를 별도로 문서화
+해야 하므로 응집도 손실. 케이스 추가가 잦아지면 그때 별도 ADR 로 도입.
+
+### A4. multi-flow contention 측정 추가
+
+기각 (probe 범위 밖). D8 에서 명시한 single-flow 격리 모델이 probe 의 핵심
+의도. multi-flow contention 은 ADR-0033 latency model 의 다른 영역으로,
+별도 도구 또는 별도 case category 로 처리.
+
+## Consequences
+
+- probe 의 case catalog (D1) 와 측정 단위 (D2/D3) 가 ADR-level 에서 명시
+  되어, 새 traffic 카테고리 추가 시 어떤 표 포맷을 따라야 하는지 분명.
+- formula vs actual 의 컬럼 의미 (D4) 가 굳어져, probe 결과를 보고 "왜
+  Drain% 가 5% 인가 / 70% 인가" 같은 질문을 빠르게 ADR-0033 가정 점검으로
+  연결 가능.
+- invariant 자동 체크 (D5) 가 ADR 에 굳어져, 향후 latency 모델 변경 시
+  monotonicity / D2H≥H2D 회귀를 probe 가 즉시 잡아낸다는 안전망 정착.
+- D8 의 case 간 격리가 명시되어, probe 결과를 single-flow 측정으로 안전
+  하게 해석 가능. multi-flow 측정이 필요해지면 별도 도구 트랙이 필요함이
+  분명.
+- A2 의 strict-mode flag 가 후속 작업 후보로 기록되어, CI 통합 요구 시
+  최소 추가 작업으로 도입 가능.

docs/adr-ko/ADR-0050-par-ccl-algorithm-module-contract.md

@@ -0,0 +1,308 @@
+# ADR-0050: CCL Algorithm Module Contract — `ccl/algorithms/*.py`
+
+## Status
+
+Accepted (2026-05-22).
+
+`src/kernbench/ccl/algorithms/` 디렉터리 안의 모듈이 AHBM CCL backend
+(ADR-0047) 에서 collective algorithm 으로 사용되려면 갖춰야 할 인터페이스,
+kernel 시그너처, 그리고 새 알고리즘 추가 절차를 명시한다. ADR-0047 D3 가
+"algorithm 모듈은 `kernel`, `kernel_args`, optional `TOPO_NAME_TO_KIND` 를
+expose 해야 한다" 라고만 한 줄로 언급하나, **algorithm 모듈 작성자가 따라야
+할 contract** 는 ADR-level 에서 정리된 적이 없다. ADR-0045 가 bench 모듈
+contract 를 다루는 것과 짝을 이룬다.
+
+## First action (제일 처음에 하는 일)
+
+알고리즘 모듈이 import 되는 시점은 두 가지다:
+
+1. **AHBM backend 진입**: 사용자 코드가 `dist.init_process_group(backend="ahbm")`
+   를 호출하면, `AhbmCCLBackend.__init__` 안에서 `self._algo_module =
+   importlib.import_module(self._merged["module"])` 가 실행된다. 이때 모듈
+   레벨에서 가장 먼저 일어나는 일:
+   - `SIP_TOPO_RING/TORUS/MESH` 같은 정수 상수가 모듈 namespace 에 노출.
+   - `TOPO_NAME_TO_KIND` 사전이 모듈 namespace 에 노출 — backend 가
+     `topo_map = getattr(self._algo_module, "TOPO_NAME_TO_KIND", None)` 로
+     조회.
+   - `kernel_args` 함수 정의 — 호출 시 호출자가 사용.
+   - `allreduce_intercube_multidevice` 같은 알고리즘 함수 정의.
+   - 모듈 마지막 줄에서 `kernel = allreduce_intercube_multidevice` 로
+     alias 가 노출.
+
+2. **ccl.yaml install 단계**: `kernbench.ccl.install.install_ipcq` 가 호출
+   되어 IPCQ neighbor table 을 푸시할 때 같은 알고리즘 모듈이 import 됨.
+
+즉, **algorithm 모듈의 첫 일은 "topology-kind 상수, `TOPO_NAME_TO_KIND`
+사전, `kernel_args` 함수, 그리고 `kernel` alias 를 모듈 namespace 에 노출
+하는 것"** 이다. 모든 노출은 import-time 부수효과로 충분하며 별도 초기화
+함수 호출이 필요하지 않다.
+
+## Context
+
+`AhbmCCLBackend` (ADR-0047) 는 process group 초기화 시점에 `ccl.yaml` 의
+`defaults.algorithm` (또는 사용자가 지정한 알고리즘 이름) 으로부터 모듈
+경로를 얻어 dynamic import 한다. backend 는 그 모듈로부터 다음 4 가지를
+기대한다:
+
+- `kernel`: collective 의 진입 함수.
+- `kernel_args(world_size, n_elem, cube_w=, cube_h=) -> tuple`: kernel 에
+  넘길 위치 인자 묶음.
+- `TOPO_NAME_TO_KIND` (optional): `topology.yaml` 의 `sips.topology`
+  문자열 (예: `"ring_1d"`, `"torus_2d"`, `"mesh_2d_no_wrap"`) 을 정수
+  상수로 매핑하는 dict.
+- (간접) IPCQ neighbor table 설치: `configure_sfr_intercube_multisip` 가
+  알고리즘 모듈의 `TOPO_NAME_TO_KIND` 와 `cube_w/h` 를 보고 SFR 을 결정.
+
+현재 코퍼스의 유일한 algorithm 모듈은 `lrab_hierarchical_allreduce.py`
+(248 줄) 이다. 이름은 "**l**eft-**r**ight **a**lternating **b**roadcast
+**hierarchical allreduce**". 향후 `ring_allreduce`, `tree_allreduce`,
+`broadcast` 같은 모듈이 추가될 때마다 이 contract 를 따라야 일관된
+디스패치가 가능하다.
+
+이 contract 가 ADR-level 에 없으면:
+
+- 새 algorithm 작성자가 ADR-0047 D3 의 한 줄 만으로 시그너처를 추론해야.
+- kernel 함수 인자 순서 (특히 `t_ptr, n_elem, cube_w, cube_h, n_sips,
+  sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h, tl`) 의 의미가 코드
+  grep 없이는 명확하지 않다.
+- `kernel_args` 가 어떤 인자를 받고 어떤 tuple 을 돌려줘야 하는지 관례
+  로만 굳어진다.
+
+## Decision
+
+### D1. algorithm 모듈은 4 가지 public symbol 을 노출한다
+
+```python
+# src/kernbench/ccl/algorithms/<name>.py
+
+from __future__ import annotations
+
+# (필수) topology-kind 상수 — 알고리즘 내부에서 사용
+SIP_TOPO_RING = 0
+SIP_TOPO_TORUS = 1
+SIP_TOPO_MESH = 2
+
+# (선택) topology 이름 → kind 매핑. backend 가 ccl.yaml/topology 의
+# 문자열 SIP topology 를 정수로 변환하는 데 사용.
+TOPO_NAME_TO_KIND = {
+    "ring_1d": SIP_TOPO_RING,
+    "torus_2d": SIP_TOPO_TORUS,
+    "mesh_2d_no_wrap": SIP_TOPO_MESH,
+}
+
+# (필수) kernel 인자 빌더
+def kernel_args(world_size: int, n_elem: int, *, cube_w: int = 4, cube_h: int = 4) -> tuple:
+    return (n_elem, cube_w, cube_h, world_size)
+
+# (필수) kernel 함수 (`tl=...` 키워드를 통해 TLContext 가 주입됨)
+def my_allreduce_kernel(t_ptr, n_elem, cube_w, cube_h, n_sips,
+                         sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h, *, tl):
+    ...
+
+# (필수) kernel alias — backend 가 `module.kernel` 로 접근
+kernel = my_allreduce_kernel
+```
+
+- `kernel` alias 는 backend 가 직접 호출하는 entry point 다. 함수 이름이
+  무엇이든 (`allreduce_intercube_multidevice` 처럼) `module.kernel = fn`
+  으로 노출해야 한다.
+- `kernel_args` 가 없으면 backend 가 알고리즘 인자를 만들 방법이 없다.
+  signature 는 D2 참고.
+- `TOPO_NAME_TO_KIND` 가 없으면 backend 는 `sip_topo_kind = 0` 으로
+  fallback 한다. 단일 topology 만 지원하는 알고리즘이라면 생략 가능.
+
+### D2. `kernel_args` 시그너처 — `(world_size, n_elem, *, cube_w, cube_h)`
+
+```python
+def kernel_args(world_size: int, n_elem: int, *,
+                cube_w: int = 4, cube_h: int = 4) -> tuple:
+    return (n_elem, cube_w, cube_h, world_size)
+```
+
+- **위치 인자**: `world_size` (= rank 수), `n_elem` (= 단일 shard 의
+  element 수, f16 기준).
+- **키워드 인자**: `cube_w`, `cube_h` (= cube mesh 크기). default 는
+  4×4 — `topology.yaml` 의 `sip.cube_mesh` 기본값과 정합.
+- **반환**: kernel 의 위치 인자 순서대로 묶은 tuple.
+
+backend 의 `all_reduce` 가 호출 시:
+
+```python
+kernel_args_tuple = self._algo_module.kernel_args(
+    self._world_size, n_elem, cube_w=eff_cube_w, cube_h=eff_cube_h,
+)
+extra_args = (sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)
+pending = self.ctx.launch(
+    self._merged["algorithm"], kernel_fn, tensor,
+    *kernel_args_tuple, *extra_args, _defer_wait=True,
+)
+```
+
+즉 kernel 의 최종 위치 인자는: `(tensor_ptr, *kernel_args_tuple,
+sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)` 이며, 거기에 `tl=...` 가
+키워드로 자동 주입된다. `kernel_args` 가 돌려주는 tuple 의 길이/순서는
+**kernel signature 와 1:1 일치** 해야 한다.
+
+### D3. `kernel` 함수 시그너처 — 정형화된 9 + tl 인자
+
+권장 시그너처:
+
+```python
+def my_kernel(
+    t_ptr: int,             # VA base of the row-wise-sharded tensor on this SIP
+    n_elem: int,            # element count per cube tile (or per shard)
+    cube_w: int,            # cube mesh width  (kernel_args 에서 옴)
+    cube_h: int,            # cube mesh height (kernel_args 에서 옴)
+    n_sips: int,            # world_size 와 동일 (rank = SIP, ADR-0024)
+    sip_rank: int,          # 이 SIP 의 rank
+    sip_topo_kind: int,     # TOPO_NAME_TO_KIND lookup 결과
+    sip_topo_w: int,        # SIP mesh width (ring_1d 면 0)
+    sip_topo_h: int,        # SIP mesh height (ring_1d 면 0)
+    *, tl,                  # TLContext (auto-injected)
+) -> None:
+```
+
+`kernel_args` 가 다른 위치 인자 순서를 채택하더라도, kernel 의 **마지막
+4 개 위치 인자는 항상 `(sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)`**
+이며 backend 가 `extra_args` 로 append 한다 (ADR-0047 D5). 이 4 개 인자는
+사용자 정의 algorithm 도 받아야 하지만, 알고리즘이 single-SIP 이라면
+그냥 무시하면 된다.
+
+`tl` 은 위치 인자가 아닌 키워드로 주입된다 — `RuntimeContext.launch` 가
+kernel 호출 직전에 `tl=tl_ctx` 를 추가한다. 따라서 kernel signature 의
+`tl` 은 keyword-only (`*, tl`) 또는 마지막 키워드 매개변수 형태여야
+한다.
+
+### D4. kernel body 의 자유도와 제약
+
+kernel body 안에서 사용 가능한 표면: ADR-0046 D3 의 모든 `tl.*` primitive.
+
+특히 자주 쓰이는 패턴:
+
+- `cube_id = tl.program_id(axis=1)` — 이 PE 가 속한 cube 인덱스.
+- `pe_addr = t_ptr + cube_id * nbytes` — cube-별 tile 의 VA 계산.
+- `acc = tl.load(pe_addr, shape=(n_elem,), dtype="f16")` — local 데이터
+  로드.
+- `tl.send(dir=...)` / `tl.recv(dir=..., shape=, dtype=)` — IPCQ
+  collective.
+- `acc = acc + recv` — TensorHandle 산술 연산자 (ADR-0046 D4).
+- `tl.store(pe_addr, acc)` — 결과 저장.
+
+kernel body 는 일반 Python 함수이며, branching/looping 자유. 단:
+
+- SimPy `yield` 또는 `async` 금지 (ADR-0046 D1).
+- TensorHandle 의 `.data` 직접 접근 금지 — phase 1 timing 모델은
+  데이터 의존을 모른다 (ADR-0020 의 2-pass 분리).
+- kernel 실행은 deterministic 해야 한다 — 같은 입력으로 두 번 실행하면
+  같은 op 시퀀스 발사. random / external IO 금지.
+
+### D5. SIP topology semantics — `sip_topo_kind` 의 의미
+
+backend 가 `topology.yaml` 의 `system.sips.topology` 문자열을 algorithm
+모듈의 `TOPO_NAME_TO_KIND` 로 lookup 하여 `sip_topo_kind` 정수로 변환.
+algorithm 은 이 정수를 보고 분기:
+
+```python
+if sip_topo_kind == SIP_TOPO_RING:
+    acc = _inter_sip_ring(...)
+elif sip_topo_kind == SIP_TOPO_TORUS:
+    acc = _inter_sip_torus_2d(...)
+elif sip_topo_kind == SIP_TOPO_MESH:
+    acc = _inter_sip_mesh_2d(...)
+```
+
+각 topology branch 는 IPCQ direction 이름 (예: `"global_E"`, `"W"`, `"S"`,
+`"N"`) 을 통해 peer 와 통신. direction 의 의미는 ADR-0023/0025 가 정의
+하며, `configure_sfr_intercube_multisip` 가 IPCQ neighbor table 을 그에
+맞춰 설치한다.
+
+algorithm 모듈은 자기가 지원하지 않는 topology kind 가 들어오면 silent
+no-op 으로 두기보다 명시적으로 `raise ValueError(f"unsupported topology
+kind {sip_topo_kind}")` 하는 것을 권장 — 실수로 backend 에 잘못 dispatch
+된 경우 빠르게 fail.
+
+### D6. ccl.yaml 의 algorithm entry 구조
+
+algorithm 모듈은 `ccl.yaml` 의 entry 와 짝을 이룬다 (ADR-0023 D10 +
+ADR-0047 D3):
+
+```yaml
+defaults:
+  algorithm: lrab_hierarchical_allreduce
+  n_elem: 8
+
+algorithms:
+  lrab_hierarchical_allreduce:
+    module: kernbench.ccl.algorithms.lrab_hierarchical_allreduce
+    # optional: world_size override
+    # optional: per-algorithm parameters consumed by configure_sfr_intercube_multisip
+```
+
+- `module`: full Python module path. backend 의 `importlib.import_module`
+  가 이 문자열을 그대로 사용.
+- `world_size` (optional): 명시되면 topology fallback 을 override
+  (ADR-0047 D2).
+- algorithm-specific parameters 는 `configure_sfr_intercube_multisip` 가
+  소비.
+
+새 algorithm 추가 시:
+
+1. `src/kernbench/ccl/algorithms/<name>.py` 작성 (D1 컨벤션).
+2. `ccl.yaml` 의 `algorithms` 섹션에 entry 추가.
+3. (필요 시) `kernbench.ccl.sfr_config` 에 SFR 설치 분기 추가.
+4. test 추가 (예: `tests/sccl/test_<name>.py`, ADR-0043 의 eval harness
+   확장).
+
+### D7. legacy "rank = flat PE index" 모드
+
+ADR-0047 D2 가 명시한 `ccl.yaml` 의 `world_size` override 경로는 legacy
+"rank = flat PE index" 테스트가 사용한다. algorithm 모듈은 이 모드 에서도
+`n_sips=world_size` 만큼의 rank 가 들어옴을 가정하면 된다 — backend 가
+rank↔(SIP, cube, PE) 매핑을 사전에 분리해 두므로 algorithm 본체에서는
+modal 분기가 필요 없다.
+
+단, single-cube workload 에서는 `cube_w=cube_h=1` 이 들어와 mesh-기반
+phase 들이 skip 되도록 작성해야 한다 (`lrab_hierarchical_allreduce.py`
+의 `single_cube = (cube_w == 1 and cube_h == 1)` 패턴 참고).
+
+## Alternatives Considered
+
+### A1. algorithm 모듈을 class 로 구조화 (`class Allreduce: kernel(...)` 등)
+
+기각. Python 모듈 namespace 자체가 algorithm 의 identity 로 사용 중이며
+(ADR-0047 D3 의 `importlib.import_module`), class 한 겹은 추가 indirection
+만 늘리고 dispatch 측 코드를 두텁게 만든다. 모듈-레벨 free function
++ `kernel` alias 패턴이 충분히 명확.
+
+### A2. `kernel_args` 를 명시적 dataclass 로 typing
+
+기각 (현재). algorithm 마다 인자 갯수가 다른 것이 정상이며, dataclass 한
+종류를 강제하면 다양한 algorithm 간 호환이 어려워진다. tuple 반환은 simple
+하고 backend 측 `*kernel_args_tuple` unpacking 과 깨끗이 맞물린다.
+algorithm 별 자체 타입 강도가 필요해지면 그 algorithm 모듈 안에서 NamedTuple
+사용은 자유.
+
+### A3. SFR 설치를 algorithm 모듈 안으로
+
+기각. SFR 설치 (`configure_sfr_intercube_multisip`) 는 topology + algorithm
+모두를 보고 IPCQ neighbor table 을 설치하는 cross-module 결정이라, algorithm
+모듈 내부보다 `kernbench.ccl.sfr_config` 같은 전용 위치가 자연스럽다. D6 의
+"필요 시 sfr_config 분기 추가" 워크플로우가 책임 분리 측면에서 더 명확.
+
+### A4. algorithm name 을 모듈 namespace 에 자동 등록 (decorator)
+
+기각. ADR-0045 (bench) 와 달리 algorithm 은 ccl.yaml entry 와 직접 묶여
+있어 추가 등록 레지스트리가 중복이다. `module` 문자열 매핑 하나면 충분.
+
+## Consequences
+
+- ADR-0047 D3 의 한 줄 contract 가 D1–D7 의 작성자-친화적 가이드로 확장
+  되어, 새 algorithm 추가 시 시그너처를 grep 으로 추론할 필요 없음.
+- D3 의 9 + tl 인자 시그너처가 표준화되어, backend 의 `extra_args` append
+  (ADR-0047 D5) 와 자연스럽게 맞물림. 향후 single-SIP-only algorithm 도
+  4 개의 sip_* 인자를 받아야 함이 명시.
+- D5 의 fail-loud 권장으로, ccl.yaml 의 topology 가 algorithm 미지원
+  topology 로 잘못 설정되면 backend 가 silent wrong-result 가 아닌
+  ValueError 로 fail.
+- D6 의 단계별 추가 절차가 명시되어, 새 algorithm 추가가 sfr_config /
+  test / ccl.yaml 어디까지 손대야 하는지 분명.

docs/adr-ko/ADR-0051-lat-routing-helper-api.md

@@ -0,0 +1,267 @@
+# ADR-0051: Routing Helper API — `AddressResolver` + `PathRouter`
+
+## Status
+
+Accepted (2026-05-22).
+
+`policy/routing/router.py` 가 노출하는 두 helper 클래스
+(`AddressResolver`, `PathRouter`) 의 모든 public API, 인자, 반환 값,
+그리고 네 가지 다른 adjacency graph 의 사용처를 명시한다. ADR-0002 가
+routing distance 와 ordering, bypass 규칙을 정의하나, **helper API 표면
+자체** 는 ADR-level 에 정리된 적이 없다.
+
+## First action (제일 처음에 하는 일)
+
+### `AddressResolver(graph)`
+
+생성 즉시 다음 두 가지를 캐시한다:
+
+1. `self._node_ids = set(graph.nodes)` — 모든 node id 의 set (lookup 용).
+2. `self._hbm_slice_bytes = hbm_total_gb * (1 << 30) // slices_per_cube` —
+   `graph.spec.cube.memory_map` 으로부터 산출 (기본 `48 GB / 8 slices = 6
+   GB`). 이 값이 `resolve()` 가 HBM PA 의 `hbm_offset` 에서 `pe_id` 를
+   복원하는 데 쓰인다.
+
+즉, **AddressResolver 의 첫 일은 "전체 node id 집합과 HBM slice 크기를
+미리 계산해 두는 것"** 이다. graph 자체는 보유하지 않는다.
+
+### `PathRouter(graph)`
+
+생성 즉시 **네 개의 별도 adjacency graph 를 동시 구축**한다:
+
+1. `self._adj_all`: 모든 edge 포함 (component-to-component routing 용).
+2. `self._adj`: `kind != "command"` 인 edge 만 (PE DMA / 일반 data path).
+3. `self._adj_mcpu_dma`: `_MCPU_DMA_EXCLUDE = {"pe_internal",
+   "pe_to_router"}` 를 제외 (M_CPU DMA 가 PE pipeline 노드로 잘못 라우팅
+   되지 않게).
+4. `self._adj_local`: `_UCIE_KINDS` 8 종을 제외 (cube-local routing 용 —
+   UCIe 가 zero-distance bus 처럼 보여 Dijkstra 가 mesh 보다 선호하는
+   것을 막음).
+
+각 그래프는 `defaultdict(list)` of `(neighbor, weight)` 형태이며,
+`edge.routing_weight_mm or edge.distance_mm` 이 weight 로 쓰인다.
+
+즉, **PathRouter 의 첫 일은 "topology edge 들을 4개의 다른 정책으로 동시
+분류하여 4 개의 인접 리스트로 구축하는 것"**. 매 `find_*()` 호출 시 적절
+한 그래프를 골라 Dijkstra 를 돌린다.
+
+## Context
+
+`policy/routing/router.py` 는 다음 두 책임을 함께 수행한다:
+
+- **이름 매핑**: 토폴로지 명명 규칙 (`sip{S}.cube{C}.<comp>`,
+  `sip{S}.io{I}.pcie_ep` 등) 의 단일 소유자. 컴포넌트 / probe / IPCQ
+  install / runtime API 가 이름 문자열을 직접 만들지 않고 helper 를 호출.
+- **경로 결정**: edge 의 `kind` 에 따른 정책 분리. 같은 src→dst 라도
+  routing 의도 (PE DMA vs M_CPU DMA vs general component routing) 에 따라
+  다른 adjacency 를 사용해야 결과가 달라진다.
+
+이 helper API 가 코드 전반에서 광범위하게 소비되는데도 (probe.py /
+distributed.py / install.py / 각종 component / tests), ADR-level 에서
+**정확한 시그너처 / 반환 의미 / 어떤 adjacency 를 쓰는지** 가 한 곳에
+정리되어 있지 않다. 본 ADR 이 그 빈자리를 채운다.
+
+## Decision
+
+### D1. `AddressResolver` 의 5 개 public API
+
+#### D1.1. `resolve(addr: PhysAddr) -> str`
+
+`PhysAddr` 인스턴스를 토폴로지의 destination node id 로 변환.
+
+```
+addr.kind == "hbm"             → f"sip{s}.cube{d}.hbm_ctrl.pe{pe_id}"
+  where pe_id = addr.hbm_offset // self._hbm_slice_bytes  (ADR-0017 D4/D9)
+
+addr.kind == "pe_resource":
+  addr.unit_type == PE         → f"sip{s}.cube{d}.pe{addr.pe_id}.pe_tcm"
+  addr.unit_type == SRAM       → f"sip{s}.cube{d}.sram"
+  addr.unit_type == MCPU       → f"sip{s}.cube{d}.m_cpu"
+  그 외                          → RoutingError("unsupported unit_type")
+
+다른 kind                       → RoutingError("unsupported address kind")
+```
+
+산출된 node id 가 `self._node_ids` 에 없으면 `RoutingError(f"node {node_id}
+not found in topology")`. 즉, address 의 syntax 가 valid 해도 topology 에
+실제로 매핑되는 노드가 없으면 fail-loud.
+
+#### D1.2. `find_m_cpu(sip, cube) -> str`
+
+`f"sip{sip}.cube{cube}.m_cpu"`. 없으면 `RoutingError`.
+
+#### D1.3. `find_pcie_ep(sip, io_id="io0") -> str`
+
+`f"sip{sip}.{io_id}.pcie_ep"`. 없으면 `RoutingError`.
+
+#### D1.4. `find_io_cpu(sip, io_id="io0") -> str`
+
+`f"sip{sip}.{io_id}.io_cpu"`. 없으면 `RoutingError`.
+
+#### D1.5. `find_all_pcie_eps() -> list[str]`
+
+전 SIP 의 PCIE_EP node id 를 정렬된 리스트로 반환. `endswith(".pcie_ep")`
+필터링. cross-SIP IPCQ 가 모든 PCIE_EP 를 enumerate 할 때 사용.
+
+명명 규칙 (`sip{S}.cube{C}.<comp>`, `sip{S}.{io_id}.<comp>`) 의 단일
+소유자가 이 클래스다 (ADR-0015 D4). 토폴로지 빌더가 같은 명명 규칙으로
+노드를 만들고, 컴포넌트는 이름 문자열을 절대 직접 구성하지 않는다 —
+모두 helper 를 거친다.
+
+### D2. `PathRouter` 의 4 개 adjacency graph
+
+생성자가 한 번에 구축. edge `kind` 가 정책을 결정:
+
+| graph             | 제외 edge kinds                               | 용도                                       |
+|-------------------|-----------------------------------------------|--------------------------------------------|
+| `_adj_all`        | (none)                                        | M_CPU↔NOC command 포함, IO_CPU/M_CPU routing |
+| `_adj`            | `"command"`                                   | PE DMA / 일반 data path                    |
+| `_adj_mcpu_dma`   | `"pe_internal"`, `"pe_to_router"`            | M_CPU DMA (PE pipeline 우회)               |
+| `_adj_local`      | `_UCIE_KINDS` (`ucie_internal`, `ucie_conn_to_router`, `router_to_ucie_conn`, `ucie_conn_to_noc`, `noc_to_ucie_conn`, `ucie_mesh`, `io_to_cube`, `cube_to_io`) | same-cube routing (UCIe bus 우회) |
+
+각 그래프는 `dict[node_id, list[(neighbor, weight)]]` 이며, weight 는
+`edge.routing_weight_mm or edge.distance_mm`. command edge 의 routing
+영향력을 명시적으로 가르고, UCIe 의 "0-distance bus" 가 mesh 보다 선호
+되는 것을 막기 위한 `_adj_local` 분리가 ADR-0017 D7 의 cross-PE-slice
+mesh-distance 요구와 정합.
+
+### D3. `PathRouter` 의 6 개 public API (+ 2 backward-compat)
+
+#### D3.1. `find_path(src_pe: str, dst_node: str) -> list[str]`
+
+**PE DMA routing**. `src_pe` 는 PE prefix (예: `"sip0.cube0.pe0"`) 이며,
+함수가 `.pe_dma` 를 자동으로 prepend 하여 실제 시작 노드를
+`"sip0.cube0.pe0.pe_dma"` 로 설정.
+
+cube-local 여부 (`_same_cube`) 에 따라 adjacency 선택:
+
+- **same-cube** (src 와 dst 가 `sip{S}.cube{C}.` prefix 공유):
+  `_adj_local` 사용. UCIe 우회를 막아 cross-PE-slice 가 mesh 거리를 정확
+  히 지불 (ADR-0017 D7).
+- **cross-cube**: `_adj` 사용. UCIe 가 자연스럽게 cross-cube path 의
+  최적 선택지로 포함됨.
+
+#### D3.2. `find_path_with_distance(src_pe, dst_node) -> tuple[list[str], float]`
+
+D3.1 과 동일한 adjacency 정책을 사용하나, 결과로 `(path, total_distance)`
+를 함께 반환. probe / 분석 도구에서 distance 메트릭이 필요할 때 사용.
+
+#### D3.3. `find_mcpu_dma_path(m_cpu_id: str, dst_hbm_id: str) -> list[str]`
+
+**M_CPU DMA path**. cube 가 같으면 `_adj_local` (mesh 안에서 마무리), 다르
+면 `_adj_all` (UCIe 경유). `_MCPU_DMA_EXCLUDE` 가 PE pipeline 노드를 자동
+배제하므로, M_CPU 가 PE 의 내부 stage 를 거쳐 routing 되는 잘못된 경로가
+나오지 않는다.
+
+#### D3.4. `find_memory_path(src: str, dst: str) -> list[str]`
+
+`pcie_ep → io_noc → cube → router mesh → hbm_ctrl` 같은 직접 메모리
+경로. `_adj_mcpu_dma` 를 사용하여 `pe_internal` 및 `pe_to_router` edge
+를 제외 — host-issued read/write 가 PE pipeline 으로 새지 않게 보장.
+probe (ADR-0049 D1 의 H2D/D2H case) 에서 직접 호출.
+
+#### D3.5. `find_node_path(src: str, dst: str) -> list[str]`
+
+임의의 두 node 사이의 path. **command edge 포함** (`_adj_all` 사용). M_CPU
+↔ NOC 같은 command-kind link 를 거쳐야 하는 IoCpuComponent /
+MCpuComponent 등이 호출.
+
+#### D3.6. backward-compat shims
+
+- `_dijkstra(start, goal) -> list[str]` — `_run_dijkstra(self._adj, …)`
+  의 thin wrapper.
+- `_dijkstra_with_dist(start, goal) -> tuple[list[str], float]` — distance
+  포함 버전.
+
+언더스코어 prefix 에서 보듯이 내부 API 인 척이지만 기존 테스트가 직접
+호출. 새 코드는 D3.1–D3.5 를 사용하고, 이 두 shim 은 deprecation 후보.
+
+### D4. Dijkstra 알고리즘 — single-source shortest path
+
+`_run_dijkstra_with_dist(adj, start, goal)`:
+
+- `heapq` priority queue.
+- `best: dict[node, distance]` — 노드별 최단 거리 캐시.
+- `prev: dict[node, predecessor]` — path reconstruction.
+- weight 는 `routing_weight_mm or distance_mm`. UCIe 처럼 routing_weight 가
+  명시되어 distance 와 다른 edge 가 있으므로 weight 분리가 의도된 것.
+
+`start == goal` 은 빠른 path `([start], 0.0)` 반환. 도달 불가는
+`RoutingError(f"no path from {start} to {goal}")`.
+
+이 알고리즘은 **deterministic** 하다 — 같은 graph + start/goal 이면 같은
+경로. 이는 SPEC R1 의 "Routing MUST be deterministic" 요구와 정합. tie-
+break 는 `heapq` 의 push 순서를 따른다 (Python list 순서가 deterministic).
+
+### D5. helper API 의 단일 소유자 원칙
+
+다음 정보는 오직 router.py 안에서만 결정된다:
+
+- 명명 규칙: `sip{S}.cube{C}.<comp>`, `sip{S}.{io_id}.<comp>`,
+  `sip{S}.cube{C}.hbm_ctrl.pe{pe_id}`.
+- adjacency 정책: 어떤 edge kind 가 어떤 그래프에 포함되는가.
+- HBM slice 크기로부터 PE id 복원 방법.
+- Dijkstra의 weight 결정 (`routing_weight_mm or distance_mm`).
+
+이 단일 소유자 원칙이 깨지면 (예: 컴포넌트가 자체적으로 `f"sip{s}..."` 를
+구성하기 시작하면) 명명 규칙 변경 시 영향 범위가 폭발한다. ADR-0015 D4 의
+정신과 정렬.
+
+### D6. helper API consumer 의 목록
+
+본 helper 가 노출하는 메소드를 호출하는 곳을 명시 (현재 코퍼스 기준):
+
+- `probes/probe.py` (ADR-0049): `find_pcie_ep`, `find_io_cpu`,
+  `find_m_cpu`, `find_node_path`, `find_mcpu_dma_path`,
+  `find_memory_path`, `find_path`, `resolve`.
+- `runtime_api/distributed.py` (ADR-0047): 간접 (engine 내부 routing).
+- `ccl/install.py` (ADR-0023): `find_all_pcie_eps`, `resolve`.
+- `sim_engine/event_log.py`: probe 와 유사하게 `find_pcie_ep`,
+  `find_memory_path`.
+- `components/builtin/m_cpu.py`, `components/builtin/io_cpu.py`:
+  `find_node_path`, `find_mcpu_dma_path`.
+- 각종 tests (test_routing.py, test_cross_sip_routing.py 등): D3.1–D3.5
+  대부분.
+
+새 consumer 가 추가될 때 본 ADR 의 D1/D3 가 그 의도에 맞는 메소드가
+이미 있는지 / 새 메소드를 추가해야 하는지 1차 판단의 기준이 된다.
+
+## Alternatives Considered
+
+### A1. 단일 adjacency graph + edge-kind filter 동적 적용
+
+기각. 매 `find_*()` 마다 graph filtering 을 다시 하면 Dijkstra 의 cache
+locality 와 성능이 떨어진다. 4 개 그래프 동시 구축 (D2) 은 메모리 비용
+이 작고 (edge ≤ 수만 건 규모), 호출 시점에 정책 선택이 O(1) 로 결정.
+
+### A2. adjacency 분리를 edge 의 `kind` 가 아닌 별도 metadata 로
+
+기각. edge `kind` 는 이미 topology builder 가 부여하며 (ADR-0015 D4 +
+ADR-0017), 별도 metadata 를 도입하면 두 시스템이 동기화되어야 하는
+중복이 생긴다.
+
+### A3. Dijkstra 대신 BFS + uniform weight
+
+기각. routing_weight_mm 이 edge 별로 다른 (mesh link / UCIe / IO-internal)
+현실에서 BFS 는 hop 수 최소화일 뿐 latency / distance 최단을 보장하지
+않는다. SPEC R1 + R2 의 결정적·정확한 routing 요구에 어긋남.
+
+### A4. helper API 를 클래스 메서드가 아닌 모듈 함수로
+
+기각. 두 클래스 (`AddressResolver`, `PathRouter`) 가 각각 cache 상태
+(`_node_ids`, `_hbm_slice_bytes`, 4 adjacency graphs) 를 보유해야 하며,
+같은 graph 인스턴스에 여러 routing 질의가 발생한다. 모듈 함수는 매 호출
+시 state 를 다시 만들거나 global 로 두어야 해서 안전성/성능 저하.
+
+## Consequences
+
+- 컴포넌트 / probe / IPCQ install / runtime API 가 모두 router.py 의
+  helper 만 호출하면 명명 규칙 변경 (예: `.io0.` → `.iochiplet0.`) 이
+  단 한 파일 수정으로 끝남 (D5).
+- D2 의 4 그래프 분리가 ADR 에 굳어져, 새 edge kind 가 추가될 때 (예:
+  Inter-die UCIe link 의 새 kind) 어느 그래프에 포함시킬지 결정의 명확
+  한 기준 제공.
+- D3.1 의 cube-local vs cross-cube 분기 (ADR-0017 D7) 가 명시되어, 향후
+  routing 동작을 변경하려는 사람이 어느 adjacency 를 건드려야 할지 안다.
+- D6 의 consumer 목록이 명시되어, helper API 변경 시 PR review 범위가
+  분명. backward-compat shim (D3.6) 의 deprecation 후보가 식별됨.

docs/adr-ko/ADR-0052-dev-oplog-memory-store-schemas.md

@@ -0,0 +1,352 @@
+# ADR-0052: OpLog + MemoryStore Schemas — sim_engine internals
+
+## Status
+
+Accepted (2026-05-22).
+
+`sim_engine/op_log.py` 의 `OpRecord` 스키마와 `OpLogger` 의 record_start /
+record_end / record_copy 동작, 그리고 `sim_engine/memory_store.py` 의
+`MemoryStore` 가 사용하는 (space, addr) 주소공간 namespace 와 read/write
+의미를 명시한다. ADR-0020 (2-pass data execution) 가 두 인프라의 존재를
+선언하나, **레코드의 정확한 필드와 의미** 는 ADR-level 에서 정리되지
+않았고 ADR-0046 D3.2 (`tl.store` visibility), ADR-0023 D9 (IPCQ copy
+record) 등 여러 ADR 이 이들의 동작에 의존하고 있다.
+
+## First action (제일 처음에 하는 일)
+
+### `OpLogger(memory_store=None)`
+
+생성 즉시 다음 3 가지 필드 초기화:
+
+1. `self._records: list[OpRecord] = []` — 누적된 op record.
+2. `self._pending: dict[int, dict] = {}` — `id(msg)` 키로 partial record
+   (record_start 시점에 만들어졌고 record_end 가 아직 안 온 것).
+3. `self._memory_store = memory_store` — 옵션 MemoryStore reference.
+   math op 의 input 스냅샷 + dma_write 의 HBM source 스냅샷 캡처에 사용.
+
+생성 시점에는 records / pending 모두 비어 있으며, `record_*` 호출이
+순차적으로 데이터를 누적한다.
+
+### `MemoryStore()`
+
+생성 즉시 `self._storage: dict[str, dict[int, np.ndarray]] = {}` 단 하나
+의 필드 초기화. 두 단계 dict (`space → addr → ndarray`) 이며 lazy 하게
+필요한 space 가 생길 때마다 inner dict 가 채워진다.
+
+즉, **두 인프라의 첫 일은 "비어 있는 누적 buffer + space-별 sparse dict
+를 만들어 두는 것"** 이다. 첫 record / write 가 실제로 도착하면 그때
+필드가 채워지기 시작한다.
+
+## Context
+
+ADR-0020 (2-pass data execution) 의 D2/D5/D7 가 다음을 선언:
+
+- Phase 1 (timing) 동안 `ComponentBase._on_process_start/end` hook 이
+  `OpLogger.record_start/end` 를 호출하여 모든 data op 의 시간 + 메타
+  데이터를 기록.
+- Phase 2 (data) 가 op_log 를 t_start 순으로 재생하여 실 데이터 결과를
+  계산.
+- 데이터 페이로드 자체는 `MemoryStore` 에 (space, addr) 키로 보관.
+
+ADR-0023 D9 (IPCQ atomic write), ADR-0027 (Megatron TP scratch
+overwrite 회피), ADR-0046 D3.2 (`tl.store` visibility) 등 후속 ADR 들이
+op_log 와 MemoryStore 의 동작에 의존하지만, **정확한 record 필드 / space
+이름 / 스냅샷 시점** 은 코드 grep 으로만 확인 가능하다. 본 ADR 이 이를
+정리한다.
+
+## Decision
+
+### D1. `OpRecord` 스키마 — 7 개 필드
+
+```python
+@dataclass
+class OpRecord:
+    t_start: float
+    t_end: float
+    component_id: str
+    op_kind: str               # "memory" | "gemm" | "math" | "unknown"
+    op_name: str               # e.g. "dma_read", "gemm_f16", "exp",
+                               #     "TileToken/DMA_READ", "composite_gemm",
+                               #     "ipcq_copy"
+    params: dict[str, Any]
+    dependency_ids: list[int] = field(default_factory=list)
+```
+
+- **`t_start` / `t_end`**: SimPy 시간 (float ns). `t_start` 는 component
+  가 op 를 시작한 시점, `t_end` 는 완료 시점. duration = `t_end - t_start`.
+- **`component_id`**: op 가 발생한 node id (예:
+  `"sip0.cube0.pe0.pe_dma"`).
+- **`op_kind`**: 4 가지 중 하나. Phase 2 DataExecutor 가 이 값으로 분기.
+- **`op_name`**: 디버깅 / 분석용 사람-친화 이름. TileToken 일 경우
+  `"TileToken/{stage_type}"` (예: `"TileToken/DMA_READ"`) 로 stage 를
+  구분.
+- **`params`**: op-종속 메타데이터 dict (D3 참고).
+- **`dependency_ids`**: 현재 사용되지 않음 (default `[]`). 향후 cross-op
+  dependency 추적이 필요해질 때를 위한 자리.
+
+### D2. `OpLogger.records` — t_start 정렬 보장
+
+```python
+@property
+def records(self) -> list[OpRecord]:
+    self._records.sort(key=lambda r: r.t_start)
+    return self._records
+```
+
+매 접근 시 `t_start` 로 stable sort. 즉 같은 t_start 인 record 들은 insertion
+순서를 유지. ADR-0020 D5 의 "t_start stable ordering" 요구와 정합.
+
+Phase 2 DataExecutor 는 항상 `records` property 를 통해 접근하므로,
+record_end 호출이 t_start 와 다른 순서로 도착해도 (예: 짧은 op 가 긴
+op 보다 늦게 시작했으나 먼저 끝남) 재정렬되어 일관된 시퀀스를 받는다.
+
+### D3. op_name 별 `params` 스키마 (`_extract_op_info` 매핑)
+
+#### D3.1. `op_kind="memory", op_name="dma_read"` (DmaReadCmd)
+
+```python
+{"src_addr": int, "nbytes": int, "handle_id": str}
+```
+
+#### D3.2. `op_kind="memory", op_name="dma_write"` (DmaWriteCmd)
+
+```python
+{
+    "src_space": str,   # handle.space ("tcm"|"hbm"|"sram"), default "tcm"
+    "src_addr": int,    # handle.addr
+    "shape": tuple, "dtype": str,
+    "dst_space": "hbm", # DmaWrite 는 항상 HBM 으로
+    "dst_addr": int,
+    "nbytes": int,
+    "handle_id": str,
+    # record_end 시점에 src_space == "hbm" 이면 snapshot 추가 (D4)
+    "snapshot": np.ndarray | None,
+}
+```
+
+#### D3.3. `op_kind="gemm", op_name=f"gemm_{dtype_a}"` (GemmCmd)
+
+```python
+{
+    "src_a_addr": int, "src_b_addr": int, "dst_addr": int,
+    "shape_a": tuple, "shape_b": tuple, "shape_out": tuple,
+    "dtype_in": str, "dtype_out": str,
+    "m": int, "k": int, "n": int,
+    # ADR-0027: per-operand + output spaces 보존
+    "src_a_space": str, "src_b_space": str, "dst_space": str,
+}
+```
+
+#### D3.4. `op_kind="math", op_name=msg.op` (MathCmd; op = "exp", "sum", "add", "where" 등)
+
+```python
+{
+    "input_addrs": list[int],   # 입력 핸들들의 addr
+    "input_shapes": list[tuple],
+    "input_spaces": list[str],
+    "input_dtypes": list[str],
+    "dst_addr": int, "dst_space": str,
+    "shape_out": tuple, "dtype": str,
+    "axis": int | None,         # reduction 인 경우만 의미 있음
+    # record_end 시점에 모든 input 의 스냅샷이 채워짐 (D4)
+    "input_snapshots": list[np.ndarray | None],
+}
+```
+
+#### D3.5. `op_kind="gemm" or "math", op_name=f"composite_{op}"` (CompositeCmd)
+
+```python
+{
+    "op": str,              # "gemm" | "math"
+    "out_addr": int, "out_nbytes": int,
+    # op == "gemm" 인 경우 GemmCmd 와 같은 필드 추가:
+    "src_a_addr": int, "src_b_addr": int,
+    "shape_a": tuple, "shape_b": tuple,
+    "dtype_in": str, "dtype_out": str,
+    "src_a_space": str, "src_b_space": str,
+    "dst_space": "hbm", "dst_addr": int,  # = out_addr
+}
+```
+
+`op == "gemm"` 이면 `op_kind = "gemm"`, 아니면 `"math"`. Phase 2 측에서
+GemmCmd 와 동일 path 로 재생되도록 alias.
+
+#### D3.6. `op_kind="memory", op_name="ipcq_copy"` (record_copy 전용 경로)
+
+```python
+{
+    "src_space": str, "src_addr": int,
+    "dst_space": str, "dst_addr": int,
+    "shape": tuple, "dtype": str, "nbytes": int,
+    "snapshot": np.ndarray | None,   # 호출자가 전달, 없으면 record_copy 가 fresh read
+}
+```
+
+`PE_DMA._handle_ipcq_inbound` (ADR-0023 D9) 가 이 record 를 발사하여 IPCQ
+slot 의 inbound copy 를 Phase 2 가 재생 가능하게 한다. 이 record 는
+`record_start` / `record_end` 를 거치지 않고 직접 `record_copy()` 로 push.
+
+#### D3.7. `op_kind="unknown", op_name=type(msg).__name__`
+
+`_extract_op_info` 가 인식 못 한 message 의 fallback. params = `{}`.
+DataExecutor 가 이 op_kind 를 만나면 skip — Phase 2 replay 에 영향 없음.
+
+### D4. snapshot 캡처 시점
+
+`OpLogger._memory_store` 가 set 되어 있을 때 record_end 가 다음을 수행:
+
+- **math op**: 모든 input addr/shape/space/dtype 으로
+  `self._memory_store.read(...)` 를 호출하여 `params["input_snapshots"]` 에
+  ndarray copy 첨부. read 실패 시 None.
+- **dma_write op**: `src_space == "hbm"` 인 경우에만 source HBM 의
+  스냅샷을 `params["snapshot"]` 에 첨부. TCM source 는 **명시적으로
+  스킵** — TCM (PE scratch) 은 Phase 2 math/gemm 재생이 다시 채우므로,
+  Phase-1-time snapshot 을 잡으면 이전 kernel 의 stale 데이터를 잡을 위험
+  (ADR-0027 postmortem: TP gemm → all_reduce race).
+- **ipcq_copy**: `record_copy` 호출자가 `snapshot=token.data` 같이 in-flight
+  스냅샷을 전달. 없으면 record_copy 가 fresh read 로 대체 시도.
+
+스냅샷은 `.copy()` 가 호출되어 (`ndarray.copy()` 가 fresh allocation) 이후
+storage mutation 으로부터 안전. ADR-0027 의 "cross-PE Phase 2 ordering"
+race 회피의 근간.
+
+`memory_store` 가 None 인 경우 (Phase 1 timing-only 모드) 스냅샷 단계는
+전부 skip. record 의 timing 정보만 보존되며 데이터 replay 는 불가능.
+
+### D5. TileToken 처리 — record_start 가 stage 정보를 캡처
+
+ADR-0014 D6 의 self-routing tile token (pipeline 모드) 은 stage_idx 가
+record_end 시점에 이미 advance 되어 있을 수 있다 (TileToken 이 다음
+component 로 이동하면서 next stage 의 params 를 캐시). 따라서:
+
+`record_start` 가 다음을 `pending[id(msg)]["snap"]` 에 미리 저장:
+
+```python
+snap["stage_type"] = stage.stage_type.name        # "DMA_READ", "GEMM", 등
+snap["stage_params"] = dict(stage.params)         # 시점의 params 복사본
+```
+
+`record_end` 에서 이 snap 을 꺼내 params 에 merge:
+
+- `params["stage_type"]` 가 final params 에 추가.
+- `stage_params` 의 key 들이 (이미 있으면 보존) merge.
+- `op_name == "TileToken"` 이면 `op_name = f"TileToken/{stage_type}"` 로
+  rewrite (예: `"TileToken/DMA_READ"`) — 같은 component 에서 발생한 서로
+  다른 stage 의 record 를 disambiguate.
+
+이 메커니즘 덕분에 DMA_READ vs DMA_WRITE, FETCH vs STORE 가 같은 component
+(예: pe_dma) 에서 발생하더라도 reporting 측에서 구분 가능.
+
+### D6. `MemoryStore` — (space, addr) 두 단계 dict
+
+```python
+class MemoryStore:
+    def __init__(self) -> None:
+        self._storage: dict[str, dict[int, np.ndarray]] = {}
+
+    def write(self, space, addr, data): self._storage[space][addr] = data
+    def read(self, space, addr, shape=None, dtype=None) -> np.ndarray: ...
+    def has(self, space, addr) -> bool: ...
+    def snapshot(self) -> MemoryStore: ...
+```
+
+#### D6.1. space namespace
+
+문자열 키. 표준 값:
+
+- `"hbm"`: HBM 데이터 (deploy_tensor + Phase 2 dma_write 결과).
+- `"tcm"`: PE-로컬 TCM (Phase 2 math/gemm 결과).
+- `"sram"`: cube-level SRAM (ADR-0023 D9.7 IPCQ slot tier).
+
+다른 space (예: `"reg"`) 도 자유롭게 허용 — `_storage` 가 lazy dict 라
+새 space 가 write 호출과 함께 자동 생성.
+
+#### D6.2. address keying
+
+`addr` 는 정수. **physical address (PA) 또는 virtual address (VA)** 일 수
+있다 — MemoryStore 자체는 address space 의 의미를 모르고 그저 키로 쓴다.
+Phase 1 의 `MemoryWriteMsg` 는 PA + VA 둘 다 write (`_create_tensor` 에서
+PA 로 zero-init, VA base 로도 zero-init), Phase 2 는 op_log 가 captured
+한 address 로 read/write.
+
+`addr` 의 의미는 호출자가 결정한다 — `MemoryStore` 는 lookup 만 제공.
+
+#### D6.3. read/write 의미 — reference store (no copy)
+
+`write(space, addr, data)`: `data` ndarray 의 reference 를 저장. **copy
+하지 않음**. 호출자가 같은 ndarray 를 이후 mutate 하면 stored value 도
+변경된다.
+
+`read(space, addr, shape=None, dtype=None)`: 저장된 ndarray 의 reference
+반환. `shape` 또는 `dtype` 이 제공되면:
+
+- `dtype != stored.dtype`: `arr.view(np_dtype)` 로 reinterpret cast (no
+  copy).
+- `shape != stored.shape`: `nbytes` 가 일치하면 `arr.reshape(shape)` (view).
+- `nbytes` 불일치: `ValueError`.
+
+데이터를 안전하게 분리하려면 호출자가 `arr.copy()` 호출. ADR-0027 의
+race 회피가 op_log snapshot 단계에서 명시적 copy 를 강제하는 이유.
+
+#### D6.4. `has(space, addr) -> bool`
+
+해당 키의 존재 여부만 확인. 데이터 인스턴스화는 안 함.
+
+#### D6.5. `snapshot() -> MemoryStore`
+
+shallow copy. inner dict 의 새 인스턴스를 만들되 ndarray reference 는
+공유. Phase 2 초기화 시점에 Phase 1 의 store 를 fork 하여 Phase 2 의
+mutation 이 Phase 1 의 다른 사용처에 영향을 주지 않게 분리하는 데 사용.
+
+### D7. op_log 가 SimPy 단일-스레드를 가정한다
+
+`OpLogger` 의 `_records`, `_pending` 은 lock 없이 사용. SimPy 가 single-
+threaded 라 `record_start` → `record_end` 사이에 다른 thread 가 끼어들
+수 없다는 가정.
+
+향후 multi-process kernbench (ADR-0047 D6) 가 도입되면 OpLogger 도 process
+별로 분리되어야 함이 명시. 단일 OpLogger 인스턴스가 multiple process 의
+record 를 받지 못한다.
+
+## Alternatives Considered
+
+### A1. op_log 를 SQLite / parquet 같은 외부 store 로
+
+기각 (현재). in-memory list 가 Phase 1 → Phase 2 의 핸드오프 latency 를
+최소화한다. 외부화는 long-running batch run 에서 의미가 있겠으나, 현재
+single-run 워크로드 에서는 overhead 만 추가.
+
+### A2. snapshot 을 record_start 시점에 캡처
+
+기각. record_start 시점은 input 이 아직 채워지지 않은 상황 (예: math
+op 의 input 이 직전 op 의 output 일 때) 이 흔하다. record_end 가 정확한
+시점.
+
+### A3. MemoryStore 를 component-별 store 로 분리
+
+기각. (space, addr) 키가 이미 충분히 disambiguation 을 제공하며, component
+별 분리는 cross-PE IPCQ copy (ADR-0023 D9) 가 source/destination 양쪽
+store 를 접근해야 하는 케이스를 복잡하게 만든다.
+
+### A4. op_log 에 cross-op dependency edge 명시
+
+부분 채택. `dependency_ids` 필드가 OpRecord 에 자리 잡고 있지만 현재
+사용되지 않음 (D1). Phase 2 DataExecutor 가 t_start 정렬 + secondary sort
+(memory ops before math at same t_start) 로 ordering 을 결정하며, 명시적
+dependency graph 가 필요해지면 이 필드가 채워질 자리. 현재는 ordering rule
+이 충분하므로 미사용.
+
+## Consequences
+
+- ADR-0020 의 op_log / MemoryStore 선언이 D1–D6 의 구체 schema 로 확장
+  되어, Phase 2 DataExecutor 작성/수정 시 정확한 필드 의미를 grep 없이
+  ADR 에서 확인 가능.
+- D3 의 op_name 별 params 스키마가 명시되어, 새 op (예: 새 reduction
+  type) 추가 시 `_extract_op_info` 분기 어디에 끼울지 명확.
+- D4 의 snapshot 시점 차이 (math = input snapshot, dma_write = HBM-only
+  snapshot) 가 ADR 에 굳어져, ADR-0027 의 cross-PE race 회피 결정이 향후
+  refactor 에서 silently 깨지지 않음.
+- D6.3 의 reference-store 의미가 명시되어, 호출자가 mutation safety 책임
+  을 인지. ADR-0027 의 explicit `.copy()` 패턴이 정당화됨.
+- D7 의 single-thread 가정이 명시되어, multi-process kernbench (ADR-0047
+  D6 supersession 후보) 도입 시 OpLogger 분리가 필요함이 분명.

docs/adr-ko/ADR-0053-dev-topology-builder-algorithms.md

@@ -0,0 +1,307 @@
+# ADR-0053: Topology Builder + Visualizer Algorithms
+
+## Status
+
+Accepted (2026-05-22).
+
+`topology/builder.py`, `topology/mesh_gen.py`, `topology/visualizer.py` 가
+함께 수행하는 토폴로지 컴파일·시각화 파이프라인의 핵심 알고리즘 선택
+(placement-driven router attachment, mesh auto-layout, source_hash 캐시,
+view projection, SVG rendering) 을 명시한다. ADR-0006 가 topology
+compilation 의 high-level intent (compiled topology, distance extraction,
+automatic diagram generation) 를 정의하나, **builder 가 실제로 어떤
+알고리즘을 사용하는지** 는 코드 grep 으로만 확인 가능했다.
+
+## First action (제일 처음에 하는 일)
+
+`resolve_topology(path_str)` 가 호출되면 다음 4 단계가 순서대로 일어난다:
+
+1. **경로 검증** (`builder.py::resolve_topology`):
+   `Path(path_str).expanduser().resolve()`, 존재 확인, file 여부 확인.
+   실패 시 `FileNotFoundError` 또는 `ValueError`.
+2. **YAML 파싱** (`_read_spec`): `yaml.safe_load`. parse error 면 line/
+   column 정보 포함한 `ValueError`. dict 가 아니면 reject.
+3. **mesh 자동 생성** (`mesh_gen.ensure_mesh_file`): topology yaml 과
+   같은 디렉터리에 `cube_mesh.yaml` 을 만들거나 (캐시 invalid 시) 재사용
+   (캐시 hit 시). 이 단계가 cube NoC 의 라우터 grid 와 부착 정보를 결정.
+4. **graph 컴파일** (`_compile_graph`): system → IO chiplets → cubes →
+   inter-cube edges → IO↔cube edges → system↔IO edges 순으로 nodes/edges
+   를 누적, 그 다음 4 개의 view projection (system, sip, cube, pe) 을
+   생성하여 `TopologyGraph` 로 묶음.
+
+즉, **topology compile 의 첫 일은 "topology.yaml 을 dict 로 읽고, 동일
+디렉터리에 cube_mesh.yaml 을 생성/검증한 뒤, system→sip→cube→pe 순으로
+flat graph + 4-view projection 을 만드는 것"** 이다.
+
+## Context
+
+`topology/` 패키지의 책임:
+
+- **builder.py** (1207 줄): topology.yaml 을 받아 `TopologyGraph` (nodes
+  + edges + 4 view projections) 를 컴파일.
+- **mesh_gen.py** (305 줄): cube NoC 의 라우터 grid 와 PE/UCIe/M_CPU/SRAM
+  부착 위치를 자동 결정하여 `cube_mesh.yaml` 로 캐시.
+- **visualizer.py** (887 줄): `TopologyGraph` 로부터 SVG 다이어그램 4종
+  (system / sip / cube / pe) 을 생성.
+
+ADR-0006 가 "topology compilation 의 결과는 distance metadata 와 diagram
+generation 의 single source" 라는 high-level 결정을 정의하나, 구체 알고리즘
+(예: placement-driven nearest-router attachment, HBM 제외 zone 산출,
+source_hash 의 어떤 필드가 invalidation 을 트리거하는가) 은 ADR 에 없다.
+
+특히 다음 결정들이 ADR-level 에 부재:
+
+- 왜 mesh_gen 이 별도 파일 (`cube_mesh.yaml`) 로 캐시되는가?
+- source_hash 가 어떤 필드를 포함하며, 어떤 변경이 재생성을 강제하는가?
+- placement coordinate 가 cube 좌표가 아닌 mm 단위인 이유?
+- HBM zone 제외와 UCIe N/S/E/W 분배가 mesh 안에서 어떻게 결정되는가?
+- view projection 4 개 (system/sip/cube/pe) 의 추상화 레벨 차이?
+
+이 ADR 이 이 결정들을 한 곳에 정리한다.
+
+## Decision
+
+### D1. compile 파이프라인 — 6 단계
+
+`_compile_graph(spec)`:
+
+1. **시스템 노드 생성** (`_instantiate_system`): `fabric.switch0`, host CPU
+   등 system-level 노드 추가.
+2. **per-SIP loop** (`for sip_id in range(system.sips.count)`):
+   - **IO chiplets** (`_instantiate_io_chiplets`): pcie_ep / io_cpu /
+     io_noc / io_ucie PHY / conn 노드 + 내부 양방향 edge 생성.
+   - **cube instantiation** (`_instantiate_cube`): cube_mesh.yaml 의 router
+     grid 를 토대로 cube-별 라우터, PE sub-components (pe_cpu, pe_dma,
+     pe_fetch_store, pe_gemm, pe_math, pe_mmu, pe_tcm, pe_scheduler,
+     pe_ipcq), m_cpu, sram, hbm_ctrl 인스턴스화 + 내부 edge 깔기.
+   - **inter-cube edges** (`_add_inter_cube_edges`): UCIe N/S/E/W mesh
+     edge.
+   - **IO ↔ cube edges** (`_add_io_to_cube_edges`): io_noc 와 cube 의
+     edge UCIe phy 사이 연결.
+3. **switch ↔ IO edges** (`_add_system_to_io_edges`): `fabric.switch0`
+   와 각 SIP 의 `pcie_ep` 사이 양방향 edge (ADR-0038 D3 + ADR-0010 의
+   cross-SIP IPCQ 경로).
+4. **view projections** 4 종 build:
+   - `_build_system_view(spec)` — Tray 레벨, SIP 들과 system switch.
+   - `_build_sip_view(spec)` — SIP 안의 cube mesh + IO chiplet.
+   - `_build_cube_view(spec)` — 단일 cube 안의 router grid + PE/M_CPU/SRAM/
+     HBM_CTRL 부착.
+   - `_build_pe_view(spec)` — 단일 PE 안의 9 sub-components + 내부 edge.
+5. **TopologyGraph 리턴**: `TopologyGraph(spec, nodes, edges, system_view,
+   sip_view, cube_view, pe_view)`.
+
+이 6 단계는 **순서가 의미를 가진다**: cubes 가 만들어진 후에야 inter-cube
+edges 가 valid 한 src/dst 를 갖고, IO chiplet 이 먼저 만들어져야 IO ↔ cube
+edge 가 그를 참조할 수 있다. 새 노드 종류를 끼울 때는 의존 관계를 보고
+적절한 위치에 삽입해야 한다.
+
+### D2. `cube_mesh.yaml` — 별도 파일 + source_hash 캐시
+
+`mesh_gen.ensure_mesh_file(cube_spec, mesh_path)`:
+
+1. `source_hash = _compute_source_hash(cube_spec)` 산출. 입력 필드:
+   - `geometry` (cube_mm.w/h 등).
+   - `pe_layout` (corners, pe_per_corner).
+   - `ucie.n_connections`.
+   - `memory_map.hbm_mapping_mode`.
+   - `placement` (m_cpu/sram pos_mm).
+2. `mesh_path` (= `topology.yaml` 와 같은 디렉터리의 `cube_mesh.yaml`) 이
+   존재하고 `existing.source_hash == source_hash` 면 재사용 (캐시 hit).
+3. 아니면 `_generate_mesh(cube_spec, source_hash)` 로 새 mesh 생성 후
+   yaml 로 저장.
+
+별도 파일로 캐시하는 이유:
+
+- mesh 생성은 PE/UCIe/router 부착 계산이 들어가 매번 다시 하기 무거움.
+- 같은 cube spec 으로 여러 번 실행 시 동일 mesh 가 보장되어야 함.
+- 사람이 직접 mesh 를 inspect / debug 할 수 있는 artifact 가 됨.
+
+`source_hash` 가 list 한 5 개 필드가 mesh 형상을 결정하는 핵심이며, 그
+외 (예: bandwidth, overhead_ns) 변경은 mesh 재생성을 트리거하지 않는다.
+
+### D3. cube NoC mesh auto-layout 알고리즘
+
+`_generate_mesh(cube_spec)`:
+
+#### D3.1. 행/열 결정
+
+- `pe_positions = _corner_pe_positions(cube_w, cube_h)`: 4 corner (NW/NE/
+  SW/SE) 마다 PE center 좌표 (mm). hardcoded `(1.5, 1.5)` / `(cube_w-1.5,
+  cube_h-1.5)` 패턴 + `pe_per_corner=2` 면 각 corner 에 2 PE 위치.
+- `col_xs = _compute_col_positions(...)`: PE 들의 x 좌표 union + `max_spacing
+  = 3.0 mm` 보다 큰 gap 에 relay 컬럼 삽입.
+- `row_ys, rows_per_half = _compute_row_positions(cube_h, n_connections,
+  pe_positions)`:
+  - `n_conn = max(n_connections, 2)` (hot path minimum).
+  - `rows_per_half = ceil(n_conn / 2)`.
+  - top 절반 + HBM 두 row + bottom 절반. HBM 은 `(cube_h/2 - 1.5, cube_h/2
+    + 1.5)` 에 위치. PE rows 와 HBM rows 사이 `hbm_gap = 1.5 mm`.
+
+#### D3.2. HBM 제외 zone
+
+`hbm_row_start = rows_per_half`, `hbm_row_end = rows_per_half + 1`.
+`hbm_col_start = n_cols // 2 - 1`, `hbm_col_end = n_cols // 2`.
+
+이 (row, col) 사각형 안의 router 슬롯은 `None` 으로 마킹 (라우터 없음).
+실제 HBM 컨트롤러는 별도 `hbm_ctrl.pe{X}` 노드로 ADR-0017 D9 의 per-PE
+파티션 패턴을 따라 부착.
+
+#### D3.3. PE 부착
+
+각 corner 의 PE 들은 다음 row 에 매핑:
+
+- Top half: NW → row 0, NE → row 1 (top_corners 안의 index).
+- Bottom half: SW → row `hbm_row_end + 1`, SE → row `hbm_row_end + 2`.
+
+각 PE 의 x 좌표가 가장 가까운 col 의 router 에 부착 (`min(range(n_cols),
+key=lambda c: abs(col_xs[c] - pe_x))`). 부착 항목은 `pe{pe_idx}.dma`,
+`pe{pe_idx}.cpu`, `pe{pe_idx}.hbm` 세 가지 (router 별 attach list 에 push).
+
+#### D3.4. M_CPU / SRAM 부착 — nearest router by Euclidean distance
+
+`placement.m_cpu.pos_mm` (default `[1.5, 5.5]`) 와 `placement.sram.pos_mm`
+(default `[1.5, 8.5]`) 의 좌표에서 가장 가까운 router 를 Euclidean
+distance 로 찾아 attach list 에 `"m_cpu"` / `"sram"` 추가.
+
+#### D3.5. UCIe N/S/E/W 분배
+
+`ucie_pe_rows = top_pe_rows + bot_pe_rows` (총 `2 * rows_per_half` 개).
+
+- UCIe-E: 매 PE row 마다 rightmost col 의 router 에 `ucie_e.c{i}`.
+- UCIe-W: leftmost col 의 router 에 `ucie_w.c{i}` (E 의 mirror).
+- UCIe-N/S: PE column 들 중 절반을 좌측, 절반을 우측으로 나눠 top row /
+  bottom row 의 해당 col 에 부착.
+
+각 UCIe connection 은 `c{i}` index 가 붙어 ucie_n_connections 만큼의 PHY
+가 분산된다 (ADR-0017 D5+).
+
+### D4. node 명명 규칙 — 단일 소유자
+
+builder.py 는 다음 명명 규칙으로 노드를 만든다 (ADR-0051 D5 의 단일
+소유자 원칙):
+
+- `fabric.switch0` — system-level switch.
+- `sip{S}.{io_id}.{pcie_ep|io_cpu|io_noc|io_ucie.{dir}|conn.{id}}` — IO
+  chiplet.
+- `sip{S}.cube{C}.{m_cpu|sram|hbm_ctrl.pe{X}|noc.r{R}c{C}|...}` — cube 내부.
+- `sip{S}.cube{C}.pe{P}.{pe_cpu|pe_dma|pe_fetch_store|pe_gemm|pe_math|pe_mmu|pe_tcm|pe_scheduler|pe_ipcq}` — PE sub-components.
+
+이 명명 규칙을 변경하려면 builder.py 와 router.py (ADR-0051) 의 helper
+양쪽이 함께 갱신되어야 한다. 컴포넌트는 명명 규칙을 직접 알지 못하고
+helper 만 호출한다.
+
+### D5. edge `kind` 분류
+
+각 edge 가 부여받는 `kind` 가 라우팅 정책 (ADR-0051 D2) 의 입력. 주요
+kind 값:
+
+- `"pe_internal"` — PE 내부 sub-component 간.
+- `"pe_to_router"` — PE_DMA ↔ cube NoC router.
+- `"router_mesh"` — cube NoC router 간.
+- `"router_to_hbm"`, `"router_to_mcpu"`, `"router_to_sram"`,
+  `"sram_to_router"` 등 — cube-attached component 간.
+- `"ucie_internal"`, `"ucie_conn_to_router"`, `"router_to_ucie_conn"`,
+  `"ucie_conn_to_noc"`, `"noc_to_ucie_conn"`, `"ucie_mesh"` — UCIe 관련.
+- `"io_internal"` — IO chiplet 내부.
+- `"io_to_cube"`, `"cube_to_io"` — IO ↔ cube 경계.
+- `"pcie"` — switch ↔ pcie_ep.
+- `"command"` — control-plane only edges (M_CPU ↔ NOC 등; PE DMA path 에서
+  제외).
+
+새 edge kind 를 추가하면 router.py 의 4 adjacency graph (ADR-0051 D2) 의
+어느 카테고리에 속할지 결정해야 한다 — 그렇지 않으면 default 로 `_adj_all`
+에만 포함되어 의도와 다른 routing 발생 가능.
+
+### D6. view projection — 4 추상화 레벨
+
+`TopologyGraph` 는 flat (nodes + edges) 외에 4 개의 view projection 을
+보유:
+
+- **system_view** (`_build_system_view`): Tray 레벨. SIP 박스들 + `fabric.
+  switch0`. PCIE 링크 표시. 외부 발표용 high-level overview.
+- **sip_view** (`_build_sip_view`): 한 SIP 안. cube mesh + IO chiplet
+  (pcie_ep + io_cpu + io_noc). UCIe N/S/E/W 가 cube 간 연결로 보임.
+- **cube_view** (`_build_cube_view`): 한 cube 안. router grid + PE/M_CPU/
+  SRAM/HBM_CTRL 부착 + UCIe PHY edge 부분. cube 내부 라우팅 / placement
+  진단용.
+- **pe_view** (`_build_pe_view`): 한 PE 안. 9 sub-components + 내부 edge
+  (pe_internal kind). 자세한 PE 내부 dataflow 검토용.
+
+view 는 spec 에서 `visualization.emit_views: [system, sip, cube]` 같이
+선택적으로 출력 (ADR-0006). pe view 는 기본 출력에서 빠져 있으나 코드는
+유지 (자세한 디버그용).
+
+### D7. visualizer.py — SVG 다이어그램 출력
+
+`emit_diagrams(graph, out_dir)` 가 모든 view 를 SVG 로 렌더. 핵심 함수:
+
+- `_render_view_svg(view)` — 일반적인 view 렌더 (router grid 가 없는
+  경우).
+- `_render_cube_view_svg(view, spec)` — cube view 전용 (HBM block 그리기,
+  router grid layout, PE/M_CPU/SRAM/HBM positioning).
+- `_draw_node`, `_draw_edge` — 노드 / edge 의 시각적 표현.
+- `_pick_scale`, `_compute_node_sizes` — 자동 스케일링.
+
+visualizer 는 **derived artifact** (ADR-0006) 로 분류되며, 코드 변경 시
+production check 대상이 아니다. CLAUDE.md 의 "Derived Artifacts" 항목과
+정합.
+
+### D8. spec 변경의 영향 범위
+
+| spec 필드                             | 영향              | mesh 재생성 |
+|---------------------------------------|-------------------|-------------|
+| `system.sips.count`                   | SIP 갯수, node 수 | No          |
+| `sip.cube_mesh.w/h`                   | cube mesh 형상    | No          |
+| `cube.geometry.cube_mm.w/h`           | cube 크기 (mm)    | **Yes**     |
+| `cube.pe_layout.corners/pe_per_corner`| PE 부착 위치       | **Yes**     |
+| `cube.ucie.n_connections`             | UCIe PHY 분배     | **Yes**     |
+| `cube.memory_map.hbm_mapping_mode`    | HBM 분배 모드     | **Yes**     |
+| `cube.placement`                      | M_CPU/SRAM 위치   | **Yes**     |
+| `cube.memory_map.*` (위 제외)         | HBM 용량 / BW     | No          |
+| `*.links.*.bw_gbs`                    | edge bandwidth    | No          |
+| `*.attrs.overhead_ns`                 | 컴포넌트 latency  | No          |
+
+위 표가 D2 의 `_compute_source_hash` 입력과 일치. mesh 재생성이 필요한
+변경은 `cube_mesh.yaml` 의 source_hash 가 자동 invalidate.
+
+## Alternatives Considered
+
+### A1. mesh 를 별도 캐시 파일 없이 매 compile 시 재생성
+
+기각. 같은 spec 으로 여러 번 호출되는 케이스 (CLI run, probe, test) 마다
+mesh 생성 비용을 다시 지불. 또한 사람이 mesh 를 inspect 할 수 있는 artifact
+가 사라짐.
+
+### A2. mesh 생성을 builder.py 에 합치기
+
+기각 (현재). 305 줄 짜리 자체 알고리즘이며, mesh layout 의 결정 (placement-
+driven router attachment, HBM exclusion zone) 이 builder 의 일반적인
+node/edge 생성 책임과 다르다. 분리 유지가 단일 책임 원칙에 더 부합.
+
+### A3. placement coordinate 를 cube 좌표 (col/row) 로 표현
+
+기각. mm 단위 좌표가 시각화 측 (visualizer) 과 mesh layout 측 (nearest-
+router 산출) 양쪽에서 일관되게 쓰인다. cube 좌표는 router grid 가 결정
+되기 전까지는 정의되지 않으므로 placement 입력에 부적절.
+
+### A4. view projection 을 lazy 하게 생성
+
+기각 (현재). 4 개 view 의 생성 비용이 작고 (보통 < 100 ms), eager 생성이
+`TopologyGraph` 를 통한 single source of truth 를 보장.
+
+### A5. visualizer 출력 형식을 SVG 외 (PNG/PDF) 도
+
+기각. SVG 가 vector + 텍스트 검색 가능 + 브라우저 직접 렌더가 가능한 가장
+유연한 형식. PNG 변환이 필요하면 별도 도구 (rsvg-convert 등) 로 후처리.
+
+## Consequences
+
+- ADR-0006 의 high-level intent 가 D1–D7 로 구체화되어, topology 변경
+  영향을 D8 표로 빠르게 가늠 가능.
+- D3 의 mesh auto-layout 알고리즘이 ADR-level 에서 굳어져, 추후 새 PE
+  부착 패턴 (예: HBM 의 6-zone 분할) 도입 시 어느 단계가 영향받는지 명확.
+- D5 의 edge kind 목록과 D7 의 view 구조가 명시되어, 새 component 종류
+  추가 시 (builder + router + visualizer) 어디까지 손대야 하는지 PR
+  reviewer 가 한눈에 파악 가능.
+- D2 의 source_hash invalidation 규칙이 명시되어, cube_mesh.yaml 이 stale
+  하게 남는 경우 (예: bw 값만 바꿨을 때) 가 정상 동작임이 분명.

docs/adr-ko/ADR-0054-eval-milestone-benches.md

@@ -0,0 +1,138 @@
+# ADR-0054: 마일스톤 평가 bench — 자기완결적 sweep + figure bench
+
+## Status
+
+Accepted (2026-05-22).
+
+ADR-0044(D1/D2)와 ADR-0045(D5)를 개정하고, ADR-0043/0044의 "로직이
+`scripts/` + `tests/`에 산다" 배치를 대체한다: GEMM/allreduce 평가
+하니스가 이제 사용자가 실행하여 모든 결과 + figure를 재생성하는
+자기완결적 **bench**가 된다.
+
+## Context
+
+ADR-0043(allreduce 평가)과 ADR-0044(GEMM 평가)는 각 하니스를 **sweep**
+(수동 `scripts/` 드라이버, 또는 allreduce의 경우 parametrized 테스트
+자체) + committed 데이터를 렌더링하는 **figure 테스트**로 분리했다.
+따라서 sweep/render 로직은 `scripts/gemm_sweep.py`,
+`tests/gemm/_gemm_plot_helpers.py`, `tests/sccl/_allreduce_helpers.py`에
+존재했다.
+
+마일스톤 요구사항("사용자가 *하나의 bench*를 실행해 모든 결과와 플롯을
+생성하도록 allreduce + GEMM 평가를 리팩터")은 그 배치로는 충족 불가다:
+bench는 production 코드이며 **`tests/`를 import할 수 없다**(ADR-0007 레이어
+방향). 평가 로직은 bench에서 닿을 수 있도록 production으로 이동해야 했다.
+
+선택한 home은 별도 `kernbench.eval` 패키지가 아니라 bench 모듈 자체다.
+bench 파일은 임의의 모듈 레벨 코드를 가질 수 있으며, 하니스를 bench로
+합치면 도메인당 파일 하나가 유지되고 패키지 레이어가 하나 줄어든다.
+
+## Decision
+
+### D1. 두 마일스톤 bench가 평가 로직을 보유
+
+- `src/kernbench/benches/milestone_1h_gemm.py` — GEMM shape×variant sweep
+  + 세 figure renderer(`scripts/gemm_sweep.py` +
+  `tests/gemm/_gemm_plot_helpers.py`에서 이동).
+- `src/kernbench/benches/milestone_1h_ccl.py` — distributed allreduce
+  드라이버, latency + buffer-kind sweep, topology diagram, FSIM 비교, 그리고
+  direct-launch 패리티 레퍼런스(`tests/sccl/_allreduce_helpers.py`에서 이동).
+
+각 파일은 해당 도메인 평가 로직의 **단일 home**이다.
+
+### D2. "평가 bench" 패턴 (ADR-0045 D5 확장)
+
+ADR-0045 D5는 bench를 단일 구성(single-SIP, 또는 ADR-0024 multi-SIP CCL
+예외)으로 고정했다. 본 ADR은 세 번째 패턴을 추가한다:
+
+- **평가 bench**는 *여러* 구성을 구동하고 figure를 렌더링할 수 있다. 외부
+  `run_bench` 엔진 대신 sweep 지점마다 자체 `GraphEngine` /
+  `RuntimeContext`를 빌드한다.
+- 그러면 외부 ctx에 제출된 handle이 없으므로, bench는 마지막에
+  **sentinel 텐서**(`torch.zeros((1, 1), …)`)를 제출하여 `run_bench`의
+  "최소 한 번 제출" 계약(ADR-0045 D4)을 만족시키고 CLI가 0으로 종료되게
+  한다.
+
+### D3. 출력 위치
+
+두 bench 모두 `src/kernbench/benches/1H_milestone_output/{gemm,ccl}/`에
+쓴다(사용자 요청 — bench 옆 아티팩트). 디렉터리는 생성된 PNG/CSV/JSON만
+보유하며(`.py`/`__init__.py` 없음), 따라서 eager-import audit(ADR-0045
+첫 동작)이 무시한다 — `pkgutil.iter_modules`는 비-패키지 하위 디렉터리를
+yield하지 않는다. `docs/diagrams/` 아티팩트처럼 **커밋된다**(원격에서
+figure를 볼 수 있도록); bench 재실행 시 제자리에서 재생성된다.
+
+### D4. GEMM 무거운 sweep — 기본은 fresh, `MILESTONE_FAST`로 재사용
+
+`milestone-1h-gemm`은 기본적으로 전체 24-sim sweep을 실행한다(분 단위;
+한 shape는 2048 tile). `MILESTONE_FAST=1`은 committed
+`docs/diagrams/gemm_sweep.json`을 재사용하고 렌더링만 한다(초 단위). 이는
+ADR-0044 D1/D2의 "무거운 sweep은 수동/`slow` 단계로 유지"를 뒤집는다:
+bench 실행이 곧 재생성이다. slow 경로는 `@pytest.mark.slow` bench
+테스트로 행사되고, fast 경로는 기본 실행된다.
+
+### D5. 테스트 + 스크립트는 thin re-export shim으로 재사용 (단일 home 유지)
+
+기존 figure 테스트와 `scripts/gemm_sweep.py` 진입점은 유지되며 이제 bench
+모듈을 재사용한다:
+
+- `tests/gemm/_gemm_plot_helpers.py` → renderer +
+  `GEMM_SWEEP_JSON`/`GEMM_PLOTS_DIR`/`ROOT`를
+  `kernbench.benches.milestone_1h_gemm`에서 re-export.
+- `tests/sccl/_allreduce_helpers.py` → 드라이버 코어, config writer, sweep
+  상수, renderer, disk aggregator를 `kernbench.benches.milestone_1h_ccl`에서
+  re-export하고, **pytest 전용** 조각은 로컬 유지: `pytest.param` 행렬
+  (`CONFIGS` / `_sweep_params` / `_bk_params`)과 fixture 결합
+  `_run_distributed`(`monkeypatch.chdir` + `_drive_distributed`) wrapper.
+- `scripts/gemm_sweep.py` → bench의 `run_sweep` 위 thin wrapper.
+
+테스트가 bench 모듈을 import하는 것은 허용된다(테스트는 production 위에
+위치, ADR-0007); 이는 전체 패키지 eager audit을 유발하며, 그것은 이미 매
+`kernbench` 실행 시 동작한다. matplotlib는 renderer 내부에서 lazy import로
+유지되어 audit의 startup 비용은 불변이다.
+
+### D6. 평면 모듈 네이밍 (`benches/` 하위 폴더 없음)
+
+`1H_milestone…`로 명명된 `benches/` 하위 패키지는 불가능하다 — Python
+패키지 이름은 숫자로 시작할 수 없다. 따라서 bench는 평면 모듈
+`milestone_1h_gemm.py` / `milestone_1h_ccl.py`이며 bench 이름은
+`milestone-1h-gemm` / `milestone-1h-ccl`(kebab-case, ADR-0045 D1에 따라
+글자로 시작)이다.
+
+## Consequences
+
+### Positive
+
+- `kernbench run --bench milestone-1h-gemm`(또는 `…-ccl`)이 도메인의 모든
+  결과 + figure를 한 명령으로 재생성한다 — 마일스톤 요구사항.
+- 평가 로직의 단일 소스(bench), shim을 통해 테스트와 스크립트가 재사용;
+  중복 없음.
+- figure 테스트와 `scripts/gemm_sweep.py`는 변경 없이 계속 동작.
+
+### Negative / limitations
+
+- 두 bench 파일이 크다(CCL 쪽은 distributed 드라이버, sweep, matplotlib
+  드로잉을 섞는다). 대부분 평가 하니스인 "bench"는 이례적이며, 본 ADR이
+  이를 정당화한다.
+- 생성 아티팩트가 명시적 요청에 의해 source tree(`src/kernbench/benches/`)
+  안에 살며 커밋된다(원격에서 figure를 볼 수 있도록); bench 재실행 시
+  재생성된다.
+- `milestone-1h-ccl`(및 기본 `milestone-1h-gemm`)은 분 단위 소요 —
+  on-demand 마일스톤 아티팩트에는 수용 가능, 일상 실행에는 아님.
+
+## Dependencies
+
+- **ADR-0007**: 레이어 방향(테스트는 production을 import할 수 있으나 bench는
+  테스트를 import할 수 없는 이유).
+- **ADR-0043 / ADR-0044**: 본 ADR이 bench로 이전하는 allreduce / GEMM 평가
+  하니스.
+- **ADR-0045**: bench 모듈 계약; 여기 D2가 그 D5(single-device 규칙)를
+  평가-bench 패턴으로 확장하고, sentinel을 위해 D4(NO_REQUESTS)에 의존.
+- **ADR-0024**: allreduce sweep이 구동하는 rank = SIP launcher.
+
+## Open questions
+
+- GEMM theoretical 모델 상수(ADR-0044 D5)를 복사 대신 ADR-0033/0014에서
+  소싱해야 하는가? 본 ADR로는 불변.
+- `build_overview_slides.py`가 GEMM 막대를 네이티브로 그리는 대신 마일스톤
+  출력 PNG를 소비해야 하는가? 여전히 open(ADR-0044 D6 / Negative).

docs/adr-ko/INDEX.md

@@ -0,0 +1,175 @@
+# ADR Index
+
+Auto-generated by `tools/generate_adr_index.py`. Total ADRs: **47**.
+
+Classification mirrors the `/report` skill's section assignment. When adding a new ADR, also add an entry to the `CLASSIFICATION` table in `tools/generate_adr_index.py`.
+
+## Design Principles
+
+- [ADR-0013](./ADR-0013-ver-verification-strategy.md) — 검증 전략 및 Phase 1 테스트 계획
+- [ADR-0033](./ADR-0033-lat-latency-model-assumptions.md) — 레이턴시 모델: 가정 및 알려진 단순화
+
+## High-level Architecture
+
+- [ADR-0003](./ADR-0003-dev-target-system-hierarchy.md) — 타겟 시스템 계층 및 모델링 범위  _(System hierarchy (Tray / SIP / CUBE / PE))_
+- [ADR-0007](./ADR-0007-api-runtime-api-boundaries.md) — 런타임 API 및 시뮬레이션 엔진 경계  _(Runtime API ↔ sim_engine boundaries)_
+- [ADR-0016](./ADR-0016-dev-iochiplet-noc-and-memory-path.md) — IOChiplet NoC와 메모리 데이터 경로  _(IOChiplet NOC and memory data path)_
+- [ADR-0017](./ADR-0017-dev-cube-noc-and-hbm-connectivity.md) — 큐브 NoC와 HBM 연결성  _(Cube NOC and HBM connectivity)_
+
+## Detailed Architecture
+
+One subsection per component file under `src/kernbench/components/builtin/`.
+
+### forwarding
+
+- [ADR-0037](./ADR-0037-dev-forwarding-component.md) — Forwarding 컴포넌트 (forwarding_v1)
+
+### hbm_ctrl
+
+- [ADR-0034](./ADR-0034-dev-hbm-controller-internal-design.md) — HBM 컨트롤러 내부 설계
+
+### io_cpu
+
+- [ADR-0036](./ADR-0036-dev-io-cpu-component-model.md) — IO_CPU 컴포넌트 모델
+
+### m_cpu
+
+- [ADR-0035](./ADR-0035-dev-m-cpu-and-m-cpu-dma-component-model.md) — M_CPU 및 M_CPU.DMA 컴포넌트 모델
+
+### pcie_ep
+
+- [ADR-0038](./ADR-0038-dev-pcie-ep-component-model.md) — PCIE_EP Component Model
+
+### pe_cpu
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+
+### pe_dma
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+- [ADR-0023](./ADR-0023-dev-ipcq-pe-collective.md) — PE-level IPCQ — Inter-PE Collective Communication
+
+### pe_fetch_store
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+
+### pe_gemm
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+
+### pe_ipcq
+
+- [ADR-0023](./ADR-0023-dev-ipcq-pe-collective.md) — PE-level IPCQ — Inter-PE Collective Communication
+
+### pe_math
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+
+### pe_mmu
+
+- [ADR-0039](./ADR-0039-dev-pe-mmu-component-model.md) — PE_MMU Component Model — 컴포넌트 + 유틸리티 이중 역할
+
+### pe_scheduler
+
+- [ADR-0014](./ADR-0014-dev-pe-pipeline-execution-model.md) — PE 파이프라인 실행 모델
+
+### pe_tcm
+
+- [ADR-0040](./ADR-0040-dev-pe-tcm-component-model.md) — PE_TCM Component Model — 듀얼 채널 BW 직렬화
+
+### sram
+
+- [ADR-0041](./ADR-0041-dev-cube-sram-component-model.md) — Cube SRAM Component Model — terminal scratchpad on cube NoC
+
+### tiling
+
+- [ADR-0042](./ADR-0042-prog-tile-plan-generators.md) — Tile Plan Generators — GEMM/Math 파이프라인 plan 빌더
+
+## Implementation Decisions
+
+### Address Scheme
+
+- [ADR-0001](./ADR-0001-mem-physaddr-layout.md) — 51비트 물리 주소 레이아웃 및 디코딩 계약
+- [ADR-0011](./ADR-0011-mem-memory-addressing-simplification.md) — 메모리 주소 지정 — PA / VA / LA 주소 모델
+
+### Routing & Helper API
+
+- [ADR-0002](./ADR-0002-lat-routing-distance.md) — 라우팅 거리, 순서 및 우회 규칙
+- [ADR-0051](./ADR-0051-lat-routing-helper-api.md) — Routing Helper API — `AddressResolver` + `PathRouter`
+
+### Memory Semantics & Local-HBM Bandwidth
+
+- [ADR-0004](./ADR-0004-mem-memory-semantics-local-hbm.md) — 메모리 시맨틱 및 로컬 HBM 대역폭 보장
+
+### Topology Compilation, Diagrams & Builder Algorithms
+
+- [ADR-0005](./ADR-0005-dev-diagram-views-distance-layout.md) — 다이어그램 뷰 및 거리 기반 레이아웃 규칙
+- [ADR-0006](./ADR-0006-dev-topology-compilation-distance-diagram.md) — 토폴로지 컴파일, 거리 추출, 그리고 자동 다이어그램 생성
+- [ADR-0053](./ADR-0053-dev-topology-builder-algorithms.md) — Topology Builder + Visualizer Algorithms
+
+### Tensor Deployment and Allocation
+
+- [ADR-0008](./ADR-0008-api-tensor-deploy-and-allocation.md) — 텐서 배포 및 할당 (호스트 할당기, PA 우선)
+
+### Kernel Execution and Host-Device Messaging
+
+- [ADR-0009](./ADR-0009-api-kernel-execution-messaging.md) — 커널 실행 메시징 및 완료 시맨틱
+- [ADR-0012](./ADR-0012-api-host-io-message-schema.md) — Host ↔ IO_CPU 메시지 스키마 (PA-우선, PE-태깅)
+
+### CLI Surface and Semantics
+
+- [ADR-0010](./ADR-0010-api-cli-surface-and-semantics.md) — 명령줄 인터페이스 및 실행 시맨틱
+
+### Component Port/Wire Fabric Model
+
+- [ADR-0015](./ADR-0015-dev-component-port-wire-model.md) — 컴포넌트 포트/와이어 모델과 패브릭 라우팅
+
+### Two-Pass Data Execution
+
+- [ADR-0020](./ADR-0020-prog-data-execution-two-pass.md) — 2-Pass 데이터 실행 모델 (타이밍 / 데이터 분리)
+
+### 2D Grid Program Identity
+
+- [ADR-0022](./ADR-0022-prog-program-id-2d-grid.md) — 2D 그리드 program_id 시맨틱
+
+### Parallelism (Launcher, DP, TP, AHBM backend, CCL algorithm)
+
+- [ADR-0024](./ADR-0024-par-sip-tp-launcher.md) — SIP-level Launcher — rank = SIP
+- [ADR-0026](./ADR-0026-par-dppolicy-intra-device.md) — DPPolicy = Intra-Device Only — sip/num_sips 필드 제거
+- [ADR-0027](./ADR-0027-par-megatron-tp.md) — Megatron-style Tensor Parallelism API
+- [ADR-0047](./ADR-0047-par-ahbm-ccl-backend.md) — AHBM CCL Backend — `torch.distributed`-compat shim
+- [ADR-0050](./ADR-0050-par-ccl-algorithm-module-contract.md) — CCL Algorithm Module Contract — `ccl/algorithms/*.py`
+
+### IPCQ Direction Addressing
+
+- [ADR-0025](./ADR-0025-algo-ipcq-direction-addressing.md) — IPCQ Direction Addressing — address-based matching
+
+### Intercube All-Reduce
+
+- [ADR-0032](./ADR-0032-algo-intercube-allreduce.md) — 큐브 간 All-Reduce — pe0 큐브-메시 리듀스 + 다중-SIP 교환
+
+### Evaluation Harnesses
+
+- [ADR-0043](./ADR-0043-eval-allreduce-harness.md) — Allreduce 평가 하니스 — `tests/sccl/`
+- [ADR-0044](./ADR-0044-eval-gemm-harness.md) — GEMM 평가 하니스 — `scripts/gemm_sweep.py` + `tests/gemm/`
+- [ADR-0054](./ADR-0054-eval-milestone-benches.md) — 마일스톤 평가 bench — 자기완결적 sweep + figure bench
+
+### Bench Module Contract
+
+- [ADR-0045](./ADR-0045-prog-bench-module-contract.md) — Bench Module Contract — registration, dispatch, and authoring
+
+### Kernel-side tl.* API (TLContext)
+
+- [ADR-0046](./ADR-0046-prog-tl-context-contract.md) — TLContext — Kernel-side `tl.*` API Contract
+
+### Memory Allocator Algorithms
+
+- [ADR-0048](./ADR-0048-mem-allocator-algorithms.md) — Memory Allocator Algorithms — VirtualAllocator + PEMemAllocator
+
+### Probe Subcommand
+
+- [ADR-0049](./ADR-0049-ver-probe-subcommand.md) — `kernbench probe` Subcommand — Traffic-Pattern Verification Harness
+
+### Sim-engine Op Log and Memory Store Schemas
+
+- [ADR-0052](./ADR-0052-dev-oplog-memory-store-schemas.md) — OpLog + MemoryStore Schemas — sim_engine internals

docs/adr-proposed/ADR-0030-mem-ipcq-physaddr.md

@@ -2,7 +2,7 @@
 
 ## Status
 
-Proposed (Blocked on ADR-0031 — PhysAddr PE-resource extension)
+Proposed
 
 ## Context
 
@@ -340,7 +340,7 @@ encoding can be plugged in later" 약속이 이행된 것.
 | `src/kernbench/sim_engine/memory_store.py` | D3: IPCQ buffer가 기존 space와 공유되는지 검증 |
 | `src/kernbench/sim_engine/engine.py` | D4: IPCQ token routing이 PhysAddr-based fabric 경로 사용 |
 | `src/kernbench/ccl/diagnostics.py` | D5: pointer_dump를 PhysAddr 포매팅으로 개선 |
-| `docs/adr/ADR-0023-ipcq-pe-collective.md` | D6: D2.5 amendment note |
+| `docs/adr/ADR-0023-dev-ipcq-pe-collective.md` | D6: D2.5 amendment note |
 | `tests/test_ipcq_physaddr.py` (new) | T1 |
 | `tests/test_ipcq_alloc.py` (new) | T2 |
 | `tests/test_ccl_install_plan.py` | T3 확장 |

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

@@ -0,0 +1,747 @@
+# ADR-0060: AHBM GQA Fused Attention 커널 (Llama3-70B)
+
+## Status
+
+Proposed
+
+**Context model:** Llama3-70B.
+**Decision drivers:** agentic workload → 낮은 batch, 긴 context;
+KV-load-bound decode; long-context prefill용 sequence-parallel (Ring KV).
+
+**Supersedes / extends:** 기존 mesh-native 어텐션 커널
+`_attention_mesh_kv`(prefill)와 `_attention_mesh_mlo`(decode), 그리고
+`milestone-gqa-llama70b` eval bench. *§A. 기존 kernbench 작업과의 관계* 참조 —
+그 코드가 본 ADR이 진짜 GQA·causal·long-context 커널로 업그레이드하는
+baseline이다.
+
+**Supporting ADRs** (efficiency / scale enabler — *GQA blocker 아님*; §8 정정
+참조): **ADR-0063** `tl.scratch_scope`(per-tile scratch 재활용 — 현실적 context
+길이에 필요), **ADR-0062** lazy `tl.load`(첫 사용 시점 auto-wait를 갖는
+non-blocking load → load/compute 오버랩), **ADR-0061** `tl.broadcast`(선택적
+mask/범용 편의). 두 GEMM(Q·Kᵀ, P·V)은 scheduler가 관리하는 `tl.composite`
+커맨드로 발행된다(기존 `CompositeCmd`; 새 커맨드 종류 없음); 진짜 GQA 자체는
+커널 재구조화만 필요(§5.2).
+
+**Algorithm lineage.** 이 커널은 **FlashAttention**(tiling + online/streaming
+softmax, P·V fused — full score matrix 미생성)이다. §4의 KV-parallel
+split-and-combine는 **FlashDecoding**(split-KV + log-sum-exp merge). §5.5의 Ring
+경로는 **Ring Attention**(KV 블록을 mesh 주위로 회전, 같은 online softmax로
+fold). 새 수학은 도입하지 않으며, 본 ADR은 이 알려진 알고리즘을 kernbench의
+**greenlet `tl` 프로그래밍 모델**(ADR-0020, ADR-0046)과 **IPCQ** PE↔PE
+collective(ADR-0023/0025)에 매핑한다.
+
+---
+
+## TL;DR — 최종 GQA 커널 (composite hybrid) pseudocode
+
+결정(§1)을 한 곳에: **GEMM → `tl.composite`; softmax 머지 + tree reduction →
+커널; `tl.load`은 lazy.** KV head별로, `G`를 matmul M 차원에 fold(진짜 GQA,
+broadcast 없음). 이것이 reference 형태이며, 산문 섹션이 각 부분을 부연한다.
+
+```python
+def gqa_attention(q_ptr, k_ptr, v_ptr, o_ptr,
+                  counter, start_pe, N, q_block, scale, *, tl):
+    # ---- geometry (kernel arithmetic, §2) ----
+    pe_id   = tl.program_id(axis=rank_axis)
+    G       = H_q // H_kv                       # query heads per KV head (=8)
+    causal  = q_block.is_prefill
+
+    for kv in range(H_kv):                       # one KV head per iteration (§5.2)
+        # Q group: G rows folded into M → [G·T_q, d]. Lazy load (ADR-0062):
+        # issue now, auto-wait at first use inside the first composite.
+        q_g     = tl.load(q_base(kv), (G * q_block.T_q, d))      # [G·T_q, d]
+        my_len  = valid_len(counter, start_pe, pe_id, N) if not causal else S_kv
+        n_tiles = ceil(my_len / TILE)
+
+        # persistent arena (outside scratch_scope, ADR-0063): -inf, 0, zeros
+        m, l, O = init_running(G * q_block.T_q, d)
+
+        for j in range(n_tiles):
+            if causal and tile_all_future(j, q_block):
+                continue                                         # causal skip (kernel if)
+            with tl.scratch_scope():                             # per-tile temporaries (ADR-0063)
+                # --- GEMM #1 on the scheduler: Q·Kⱼᵀ; K streamed by composite ---
+                Sj = tl.composite("gemm", a=q_g,
+                                  b=tl.ref(k_tile(kv, j), (d, TILE))) * scale   # Kᵀ pre-stored [d,TILE]
+                if causal and tile_partial(j, q_block):
+                    Sj = Sj + causal_mask(j, q_block)            # additive boundary mask
+                # --- online softmax merge in the kernel (MATH ops) ---
+                m_new = tl.maximum(m, tl.max(Sj, axis=-1))
+                P     = tl.exp(Sj - m_new)
+                corr  = tl.exp(m - m_new)
+                l     = l * corr + tl.sum(P, axis=-1)
+                # --- GEMM #2 on the scheduler: P·Vⱼ; V streamed by composite ---
+                Oj    = tl.composite("gemm", a=P,
+                                     b=tl.ref(v_tile(kv, j), (TILE, d)))
+                O     = O * corr + Oj                            # running merge
+                m     = m_new
+
+        # ---- cross-PE combine (§4): log-sum-exp tree to root, or store ----
+        if N == 1:
+            tl.store(o_base(kv), O / l)
+        else:
+            tree_reduce_and_store(m, l, O, pe_id, N, o_base(kv))  # tl.send/tl.recv
+```
+
+> **이 형태인 이유:** 두 `tl.composite` 호출이 tiling + K/V DMA 스트리밍 +
+> cross-tile 파이프라이닝을 PE_SCHEDULER에 offload한다(CPU는 coarse descriptor를
+> 내고 앞서 나가 → 엔진이 saturate 유지, §1); softmax 머지와 IPCQ tree reduction은
+> 커널에 남는다(기존 `CompositeCmd`가 cross-tile 상태를 못 들기 때문). `K`는
+> reshape-not-transpose caveat를 피하려 transpose된 `[d, TILE]`로 미리 저장된다
+> (§3, §B). 전용 "flash-composite" 커맨드 종류는 도입하지 **않는다**(§8 항목 4).
+
+---
+
+## A. 기존 kernbench 작업과의 관계 (먼저 읽을 것)
+
+kernbench는 **오늘날 이미 IPCQ 상에서 online-softmax `(m, ℓ, O)` 머지로
+FlashAttention을 돌린다.** 두 커널이 존재한다:
+
+| File | Role | Mechanism |
+|---|---|---|
+| `src/kernbench/benches/_attention_mesh_kv.py` | prefill (Ring K/V) | per-rank partial attention, bidirectional K/V fan-out, online-softmax fold |
+| `src/kernbench/benches/_attention_mesh_mlo.py` | decode (split-KV) | per-rank one-shot partial attention, bidirectional `(m,ℓ,O)` fan-out, log-sum-exp merge |
+
+둘 다 `milestone-gqa-llama70b`(4 패널:
+`{single,multi}_user × {prefill,decode}`,
+`src/kernbench/benches/milestone_gqa_llama70b.py`)이 구동하며
+`tests/attention/test_milestone_gqa_llama70b.py`에서 테스트된다.
+
+**이들은 greenlet `tl` API로 작성되었다:** `tl.load`, `tl.dot`,
+`tl.softmax`/`tl.max`/`tl.sum`/`tl.exp`, `tl.send`/`tl.recv`, 그리고
+`TensorHandle`에 대한 Python `-`/`*`/`/`(각각 `MathCmd` emit) — GEMM은 composite가
+아니라 **blocking `tl.dot`**로. 실행 중 `(m, ℓ, O)`는 루프를 관통하는 Python
+`TensorHandle`일 뿐이다. 본 ADR은 실행 상태와 softmax 머지는 커널에 유지하되
+**두 GEMM은 scheduler가 관리하는 `tl.composite` 경로로 옮긴다**(§1 참조) — 이것이
+**설계에 중요**하다.
+
+**baseline의 의도된 세 한계** — 효율적 GQA 커널이 정확히 들어내야 하는 것:
+
+1. **GQA 재사용 없음.** `h_q == h_kv == 1`
+   (`test_milestone_gqa_llama70b.py:137-142`). 테스트는 이를 *broadcast view*에
+   대한 MemoryStore byte-conservation 실패로 돌리지만, 그 실패는 baseline의
+   **head-packing 핵**의 속성이다(`_view(K, (h_q·d, S_kv))`는 모든 head를 하나의
+   matmul 차원에 뭉치고 `h_q == h_kv`일 때만 byte를 보존). 올바른 수정은
+   broadcast op이 아니라 **커널 재구조화**다: **한 번에 한 KV head**를 처리하고
+   `G` group 행을 matmul **M** 차원에 fold(§5.2). 이는 byte-보존 reshape만 쓰므로
+   진짜 GQA(`h_q = G·h_kv`)가 **신규 primitive 없이** 돈다 — §8 참조.
+2. **O(N) reduction.** baseline은 all-to-all bidirectional fan-out을 하여 *모든*
+   rank가 full 답을 갖는다(`n_ranks − 1` 단계). 어텐션은 query owner에서만 `O`가
+   필요 → root로의 **tree reduction**은 `⌈log₂ N⌉` 단계(§4).
+3. **검증 스케일만.** `S = 16`인 이유는 1 MiB scratch bump allocator가 per-tile
+   임시값을 누수하고(`test_milestone_gqa_llama70b.py:123-148`) causal masking /
+   tiling이 없기 때문 → **ADR-0063**(재활용) + §5(tiling, causal skip) + composite
+   K/V 스트리밍(§3) + **ADR-0062**(lazy load 오버랩)으로 해결.
+
+**문서 부채(범위 밖이나 기록):** baseline은 ADR-0055/0056/0057/0058/0059를
+인용하나 **파일로 존재하지 않는다** — ghost 참조다. 본 ADR은 이를 소급
+작성하지 않으며; 권고는 Detailed Design Document의 *Open Decisions* 참조.
+
+---
+
+## 0. 참조 차원 (Llama3-70B)
+
+| Symbol | Meaning | Value |
+|---|---|---|
+| `H_q` | query heads | 64 |
+| `H_kv` | KV heads | 8 |
+| `G` | GQA group size = `H_q / H_kv` | 8 |
+| `d` | head dim | 128 |
+| `L` | layers | 80 |
+| `D` | model dim | 8192 |
+
+하드웨어 요약:
+- **AHBM (chip)** = **CUBE**(메모리 cube, 각각 **PE**를 담은 logic die) 집합 +
+  **IO die**(ADR-0003).
+- **IPCQ**: PE↔PE 큐, PE당 4 mesh-방향 queue-pair
+  (`N/S/E/W`, ADR-0023 D3; inter-SIP용 `global_*`, ADR-0032). 커널
+  API: `tl.send(dir, src)` / `tl.recv(dir, shape, dtype)`
+  (`tl_context.py:402-499`).
+- **Composite command** (`CompositeCmd`, `pe_commands.py:144-162`): 단일
+  GEMM(또는 MATH) *head* + element-wise *epilogue* 단계
+  (`bias/relu/scale/add/...`), PE_SCHEDULER에 **non-blocking**으로 발행되며,
+  scheduler가 tile plan을 생성하고 타일당 DMA→GEMM→write를 스트리밍한다
+  (ADR-0014 D6; `pe_scheduler.py:104-143`). 이는 일반적 multi-op DAG가 **아니다**:
+  두 GEMM을 chain할 수 없고, 인스턴스 간 register 상태를 못 들며, IPCQ를
+  pop/wait 못 한다. 따라서 본 ADR은 두 어텐션 GEMM(Q·Kᵀ, P·V)을 **각각** 자체
+  composite로 발행하고 cross-GEMM softmax 머지 + IPCQ reduction은 커널에 유지한다
+  — 새 "flash-composite" 커맨드 종류는 **필요 없다**(§1, §8 참조).
+- **Allocation policy (SP):** multi-user 패널에서 CUBE당 query head 하나; KV
+  group의 `G` query head는 한 CUBE의 PE 안에 매핑.
+
+**직교하는 두 매핑 레이어** (round-robin placement):
+- **Layer 1 (KV-parallel, intra-request):** KV token `i`는 PE
+  `(start_pe + i) mod N`에 안착 — round-robin이라 한 긴 request가 `N` PE에 걸쳐
+  분할, ≤1 token으로 균형.
+- **Layer 2 (inter-request):** `start_pe = request_id mod N`이 request별 "PE-0
+  역할"을 회전.
+- 결합: `pe = (request_id + global_token_idx) mod N`.
+
+`N` = 하나의 (query head, request)에 대한 KV-parallel reduction group의 PE 수.
+Reduction은 **intra-CUBE** 유지(query head는 CUBE를 가로지르지 않음 — 명시적
+non-goal).
+
+---
+
+## 0.5 커널 경계, 전제조건, I/O 계약
+
+### 0.5.1 디코더 레이어 내 위치
+
+```
+1. RMSNorm
+2. QKV projection (GEMM)        ─┐  qkv_rope kernel (SEPARATE, upstream)
+3. RoPE on Q and K              ─┤
+4. write new K,V → KV cache     ─┘
+5. ===== THIS KERNEL: FlashAttention =====   (post-RoPE Q, K-cache, V-cache → O)
+6. Output projection (GEMM)        out_proj kernel (SEPARATE, downstream)
+7. residual add → FFN ...
+```
+
+**전제조건 (upstream `qkv_rope`, 본 커널 아님):**
+- **P1.** Q는 이미 RoPE-회전됨. 여기서 회전 없음.
+- **P2.** K-cache는 **post-RoPE** K 저장(어텐션 시점에 재회전 안 함 — post-RoPE
+  캐싱의 이유).
+- **P3.** decode의 경우, 새 step의 K 행은 RoPE-회전되어 그 소유 PE의 K-cache
+  슬롯에 **본 커널 launch 전에 `qkv_rope`가** 추가한다. ⇒ 본 커널은 KV cache에
+  대해 **pure read**.
+- **P4.** V는 회전 안 함; V-cache는 raw projected V 보유.
+- **P5.** upstream RoPE 위치는 token의 **절대 global 위치**
+  `global_idx = local_slot·N + ((pe_id − start_pe) mod N)`(§2.1). Round-robin
+  placement ≠ RoPE 위치.
+
+### 0.5.2 Shape 기호
+`T_q` = 이번 launch의 query 길이(decode: 1; prefill/chunk: chunk 폭).
+`S` = 전체 context 길이. `S_pe` = 이 PE 소유 key ≈ `⌈S/N⌉`.
+저장 `bf16`(numpy proxy `f16`, `memory_store.py:16`); `m,ℓ,O` 누산기는 정밀도가
+중요한 곳에서 `f32`.
+
+### 0.5.3 INPUTS (커널 launch당)
+
+| Input | Shape (per KV head) | Location | Notes |
+|---|---|---|---|
+| `Q` | `[G, T_q, d]` | per-PE HBM (loaded to TCM) | post-RoPE (P1). group의 `G` query 행 batched. |
+| `K_cache` | `[S_pe, d]` | per-PE HBM, base `K_base[pe]`, contiguous | post-RoPE (P2). Read-only. 로컬은 dense, global index는 strided (§2.1). |
+| `V_cache` | `[S_pe, d]` | per-PE HBM, base `V_base[pe]` | raw V (P4). Read-only. |
+| `global_token_counter` | scalar | launch arg | 커널이 `S_pe`, slot↔global, causal bound 도출. |
+| `start_pe` (= `request_id mod N`) | scalar | launch arg | Layer-2 회전. |
+| `pe_id`, `N` | scalars | launch / `tl.program_id` | reduction-group geometry. |
+| `q_block_meta` `{q_start, T_q}` | launch arg | prefill/SP causal masking & skip. |
+| `O_base` | address | launch arg | 최종 O가 쓰이는 곳 (root PE만). |
+| `softmax_scale` | scalar | launch arg | `1/√d`. |
+
+**Ring Attention (§5.5)**에는 ring step마다 추가: IPCQ로 들어오는 `K_block,
+V_block`을 ping-pong 버퍼에(post-RoPE), 그리고 causal step-skip용
+`step_kv_global_range`.
+
+### 0.5.4 OUTPUTS
+
+| Output | Shape | Location | Notes |
+|---|---|---|---|
+| `O` (final) | `[G, T_q, d]` | `O_base`, **root PE만** | 정규화 `O_acc / ℓ_acc`, 저장 dtype으로 cast. |
+
+**중간값 (non-root PE, IPCQ 상 — 커널-가시 출력 아님):**
+`(m_i, ℓ_i, O_i)` (`m,ℓ`: `[G, T_q]`; `O_i`: `[G, T_q, d]` 비정규화)이
+`tl.send`로 `tree_parent`에 push(§4). `O_i`가 무거운 부분.
+
+**No-SP (`N=1`):** IPCQ partial 없음; 단일 PE의 실행 중 `(m,ℓ,O)`를 제자리에서
+정규화하여 `O_base`에 기록.
+
+**명시적 비-출력:** KV-cache write(upstream, P3), output projection(downstream),
+score `S` 및 prob `P`(미생성).
+
+---
+
+## 1. 결정 (메커니즘)
+
+**커널을 *하이브리드*로 구현한다: 두 GEMM(Q·Kᵀ, P·V)을 scheduler가 관리하는
+`tl.composite(op="gemm")` 커맨드로 발행하고, online-softmax 머지와 cross-PE
+reduction은 커널 수준 `tl` op으로 유지한다.** `tl.load`은 **lazy**다(non-blocking;
+wait는 load된 데이터의 첫 사용 시점에 자동 삽입 — ADR-0062). 따라서 명시적 HBM
+load가 뒤따르는 compute와 오버랩된다. kernbench의 실행 + latency 모델에 근거한
+근거:
+
+- **GEMM tiling이 PE_SCHEDULER에 offload된다.** `CompositeCmd`는 non-blocking
+  (`kernel_runner.py:182-191`, `pe_scheduler.py:104-121`): 커널이 **하나의 coarse
+  descriptor**(M = `G·T_q`, PE당 전체 tile sweep)를 push하면 scheduler가 tile
+  plan을 생성하고 타일당 DMA→GEMM→write를 스트리밍한다(ADR-0014 D6). K/V는
+  scheduler가 HBM에서 스트리밍하는 `tl.ref` operand이므로, 타일당 **K/V prefetch는
+  scheduler의 일**이다 — 명시적 prefetch op 없음. CPU(greenlet)는 현재 composite가
+  도는 동안 **다음** composite를 낼 수 있게 풀려나, scheduler가 타일을 가로질러
+  GEMM 엔진을 saturate 유지한다.
+- **이는 하드웨어를 반영**하며 CPU issue-rate를 execution-rate로부터 decouple한다.
+  대조적으로 blocking per-op `tl.dot` 경로는 매 GEMM마다 CPU를 멈추고, 끼어드는
+  softmax MATH op 동안 GEMM 엔진 **버블**을 남긴다; CPU가 fine-grained per-tile
+  발행을 따라잡을 수 있을 때만 현실적이다.
+- **실행 중 `(m, ℓ, O)` flash 상태는 Python `TensorHandle`로** 루프를 관통한다
+  (baseline이 이미 그렇게 함); softmax 머지(max/exp/sum/rescale)는 두 GEMM
+  composite **사이**의 커널 수준 `tl` MATH다. 기존 `CompositeCmd`는 두 GEMM을
+  chain하거나 cross-tile register 상태를 못 들므로(§0), 머지는 필연적으로 커널에
+  산다 — 이것이 하이브리드 분할이지, 우회한 한계가 아니다.
+- **Cross-PE 결합**은 P·V 후 `(m, ℓ, O)`에 대한 log-sum-exp **tree**로, 커널 수준
+  `tl.send`/`tl.recv`(§4)를 통해 — 불변.
+
+그래서 per-tile 내부 파이프라인은:
+
+```
+q_g = tl.load(Q group)                         # lazy; auto-wait at first use
+per tile j:
+  Sⱼ = tl.composite("gemm", a=q_g, b=tl.ref(Kⱼ))  → Sⱼ   # scheduler streams Kⱼ DMA + GEMM
+  Sⱼ += maskⱼ                                          # kernel MATH, boundary tile only
+  online-softmax: mⱼ, m_new, P, corr, ℓ               # kernel MATH
+  Oⱼ = tl.composite("gemm", a=P, b=tl.ref(Vⱼ))   → Oⱼ   # scheduler streams Vⱼ DMA + GEMM
+  O = O*corr + Oⱼ;  m = m_new                         # kernel MATH (running merge)
+```
+
+각 타일의 MATH 임시값을 `tl.scratch_scope`(ADR-0063)로 감싸 scratch를 O(1)로
+유지하고, 다음 타일의 composite를 현재 타일 결과를 wait하기 전에 발행
+(non-blocking 핸들)하여 scheduler가 타일을 가로질러 파이프라인되게 한다.
+
+제어 흐름(tile skip, mask 생성, reduction 스케줄링, 주소 산술)은 **커널**에 산다
+(greenlet 본문의 평범한 Python `if`/산술). 이는 kernbench의 greenlet 모델이 이미
+허용하는 바 그대로다(`kernel_runner.py`, ADR-0020 D3).
+
+> **이것이 무엇을 대체하나.** 이전 iteration은 "composite는 latency 이점 없음"이라는
+> 논거로 순수 greenlet primitive 경로(전부 `tl.dot`, composite 없음)를 제안했다.
+> 그것은 **오직** 시뮬레이터가 현재 per-op CPU 발행 비용을 **0**으로 청구하기
+> 때문에만(`dispatch_cycles=0`, `pe_cpu.py`) 성립한다 — descriptor offload가 숨기려
+> 존재하는 바로 그 CPU issue-rate / DMA-program 비용을 모델이 지워버린 것이다.
+> 하이브리드가 효율 커널의 충실한 표현이다. 이점의 **측정 가능한** 크기(CPU가 많은
+> 타일에 대해 엔진을 saturate할 수 있는가?)는 op 종류별 발행 비용 모델링에
+> 좌우되며, future work로 추적된다(cost model; §9). `dispatch_cycles=0`에서도
+> non-blocking composite 경로는 blocking `tl.dot` 경로가 남기는 GEMM 엔진 버블을
+> 채운다.
+>
+> **이것이 효율적 선택인 이유.** GEMM tiling + DMA 스트리밍 + cross-tile
+> 파이프라이닝은 검증된 `CompositeCmd` scheduler 경로에 offload되고; softmax 머지와
+> 검증된 IPCQ collective는 커널에 남는다. 진짜 새 기계장치는 작고 범용적인 두
+> primitive뿐이다(ADR-0062 lazy `tl.load`, ADR-0063 `tl.scratch_scope`); reduction은
+> `tl.send`/`tl.recv`를 재사용한다. 전용 "flash-composite" 커맨드(softmax 머지 +
+> carried register 상태 + IPCQ-push epilogue를 내재화하는 한 종류)는 만들지
+> **않는다** — 크고 특수목적이며, 하이브리드 대비 유일한 delta(완전 softmax
+> offload)가 현재 모델링 충실도에서 정당화되지 않는다; §8 참조.
+
+---
+
+## 2. 메모리 레이아웃 & 드라이버 책임
+
+### 2.1 KV cache 할당
+- per-PE KV 버퍼는 per-PE 최대 context `⌈max_context / N⌉ × d × dtype`로 K와 V를
+  각각 sizing. kernbench에서는 sequence 차원에 대해 `pe`(또는 `cube`) 축이
+  `row_wise`인 `DPPolicy`로 배치된다(`policy/placement/dp.py`; baseline은 K/V에
+  `DPPolicy(pe="row_wise")`, Q에 `replicate`).
+- PE 안에서 할당된 token은 **연속 append**된다(slot 0,1,2,…). global index는
+  strided(`i, i+N, …`)이나 per-PE 버퍼는 dense ⇒ DMA가 연속 유지.
+- Slot → global: `global_idx = local_slot·N + ((pe_id − start_pe) mod N)`.
+
+### 2.2 드라이버 launch당 의무 (최소)
+드라이버는 launch마다 base + counter + rotation을 공급하고; 커널이 나머지를
+도출한다:
+
+| Launch arg | Purpose |
+|---|---|
+| `K_base[pe]`, `V_base[pe]` | per-PE KV 버퍼 base (tensor VA에서) |
+| `O_base` | 어텐션 출력 목적지 |
+| `global_token_counter` | 현재 sequence 위치 |
+| `start_pe = request_id mod N` | Layer-2 회전 |
+| `pe_id` (`tl.program_id`), `N` | geometry |
+| `q_block_meta` | prefill/SP query block 시작 + 길이 |
+
+커널 도출(평범한 산술):
+- **이번 step에 내가 쓸 차례:** `(start_pe + counter) mod N == pe_id`
+- **내 valid 길이:** `base = counter // N; rem = counter % N;
+  my_len = base + (1 if ((pe_id − start_pe) mod N) < rem else 0)`
+- **read 범위:** tile `0 .. ⌈my_len / TILE⌉`.
+- **causal bound / per-tile skip:** query block과 각 tile의 global 위치에서.
+
+드라이버 = base + counter + rotation. 단일 공식
+`(request_id + token_idx) mod N`이 placement policy 전부다.
+
+---
+
+## 3. Per-tile op 시퀀스 (greenlet `tl`)
+
+한 iteration = 한 PE의 한 KV tile. 두 GEMM은 `tl.composite(op="gemm")`
+(scheduler 관리 tiling + K/V DMA 스트리밍); softmax 머지는 그 사이의 커널 `tl`
+MATH다. 실제 `tl` 이름(`tl_context.py`), lazy `tl.load`(ADR-0062),
+`tl.scratch_scope`(ADR-0063) 사용:
+
+```python
+# running state (persistent arena — allocated once, outside the scope)
+# m: [G, T_q]  l: [G, T_q]  O: [G, T_q, d]
+q_g = tl.load(Q_group_ptr, (G*T_q, d))               # lazy; auto-wait at first use (ADR-0062)
+
+with tl.scratch_scope():                              # per-tile MATH temporaries recycled
+    Sj    = tl.composite("gemm", a=q_g,                # [G·T_q, TILE]; scheduler streams Kⱼ DMA
+                         b=tl.ref(K_base + j*TILE*d, (TILE, d))) * softmax_scale
+    if mask_j is not None:
+        Sj = Sj + mask_j                              # additive causal mask (boundary tile)
+    m_j   = tl.max(Sj, axis=-1)
+    m_new = tl.maximum(m, m_j)
+    P     = tl.exp(Sj - m_new)                        # no full-matrix softmax; streaming
+    corr  = tl.exp(m - m_new)                         # rescale factor for old accumulators
+    l     = l * corr + tl.sum(P, axis=-1)
+    Oj    = tl.composite("gemm", a=P,                  # [G·T_q, d]; scheduler streams Vⱼ DMA
+                         b=tl.ref(V_base + j*TILE*d, (TILE, d)))
+    O     = O * corr + Oj                             # running merge (kernel MATH)
+    m     = m_new
+```
+
+Notes:
+- `q_g`는 `[G·T_q, d]`로 reshape된 GQA-batched query(`G` group 행을 matmul M
+  차원에 fold; byte-보존). 한 K/V tile이 모든 `G·T_q` 행을 서비스 — GQA 재사용
+  레버 — broadcast 없이.
+- **K/V는 `tl.ref` operand**로 composite scheduler가 HBM에서 타일당 스트리밍한다
+  (`pe_scheduler.py:104-143`): 그것이 *바로* prefetch/파이프라인이므로 명시적
+  prefetch op이 없다. 타일 `j+1`의 composite를 타일 `j`를 wait하기 전에 발행
+  (non-blocking 핸들)하면 scheduler가 타일을 가로질러 파이프라인되고 GEMM 엔진이
+  saturate 유지된다.
+- `tl.trans`는 kernbench에서 **메타데이터 전용**이고(`tl_context.py:390`)
+  `MemoryStore.read`는 transpose가 아니라 *reshape*한다(`memory_store.py:73`).
+  zero/structural 실행엔 무해하나; 비자명 수치 데이터에선 reshape-not-transpose가
+  되어 transpose된 K로의 Q·Kᵀ는 주의가 필요하다(§11) — K를 transpose해 `[d, TILE]`로
+  미리 저장하거나, 실제 `tl.transpose`를 추가(후보 primitive, 시뮬레이터의
+  성능-모델링 목적상 아마 불필요).
+- Masking: 커널이 query/KV global offset에서 boundary-tile mask를 만들어 더한다
+  (`Sj + mask_j`); full-past tile은 `None` 전달; full-future tile은 **skip**(`if`이
+  enqueue 안 함).
+- **새** "flash-composite" 커맨드 종류(softmax 머지 + carried `(m,ℓ,O)`를
+  내재화하는 것)는 쓰지 **않는다**; 기존 `CompositeCmd`가 각 GEMM을 담당하고 머지는
+  커널에 남는다(§1, §8 항목 4).
+
+---
+
+## 4. Reduction (KV-parallel / SP combine) — root로의 tree
+
+tile sweep 후 각 PE는 `(m_i, ℓ_i, O_i)`(비정규화)를 가진다. 결합/교환 가능한
+log-sum-exp 머지로 결합(baseline의 fold와 동일 수학, `_attention_mesh_mlo.py:117-122`):
+
+```python
+def merge(m_a, l_a, O_a, m_b, l_b, O_b):
+    m = tl.maximum(m_a, m_b)
+    sa, sb = tl.exp(m_a - m), tl.exp(m_b - m)
+    return m, l_a*sa + l_b*sb, O_a*sa + O_b*sb
+# final: O = O_root / l_root      # normalise once at the tree root
+```
+
+- **타이밍:** 교환은 **P·V 후**(각 PE가 최종 `O_i`를 가짐).
+- **토폴로지:** mesh tree, depth `⌈log₂ N⌉`, 물리 이웃 상. `N=8`의 경우: level-0
+  `(0↔1,2↔3,4↔5,6↔7)`, level-1 `(1↔3,5↔7)`, level-2 `(3↔7)`, root = `7`. 각 tree
+  pair는 물리 `N/S/E/W` 이웃이어야 `tl.send(dir)`/`tl.recv(dir)`이 실제 link를
+  쓴다(튜닝 항목 §9; 이웃을 wiring하는 SFR install은
+  `configure_sfr_intracube_pe_ring`, baseline이 사용).
+- **왜 baseline fan-out이 아니라 tree인가:** 어텐션은 한 곳(`O_base`를 쓰는 PE)에서
+  `O`가 필요하다. tree는 `⌈log₂ N⌉` 단계(`N=8`에 3) vs baseline all-to-all의 `N−1`
+  (7). per-PE sweep이 짧은 decode에서 reduction이 지배적 비용이므로 실제 이득이다.
+  (baseline fan-out은 답을 모든 rank에 복제 — 여기선 불필요.)
+- **Payload:** `O_i`(`d`-vector)가 무겁고; `m,ℓ`은 `(g, T_q)`당 scalar. 별도
+  `tl.send`로 전송(baseline은 triplet을 세 send로 — 같은 패턴).
+
+커널 구조(greenlet):
+
+```python
+# leaf / internal node: fold children that send to me, then forward up
+for child_dir in tree_children_dirs(pe_id, N):
+    m_c = tl.recv(child_dir, m.shape); l_c = tl.recv(child_dir, l.shape)
+    O_c = tl.recv(child_dir, O.shape)
+    m, l, O = merge(m, l, O, m_c, l_c, O_c)
+if not is_root(pe_id):
+    tl.send(parent_dir(pe_id), m); tl.send(parent_dir(pe_id), l); tl.send(parent_dir(pe_id), O)
+else:
+    tl.store(O_base, O / l)
+```
+
+`tl.recv`는 데이터 도착까지 blocking(`tl_context.py:446-499`); 머지 순서가 정적
+tree로 고정되므로 data-dependent `pop`이 불필요 — **하드웨어/composite
+`pop`-as-dependency 변경이 필요 없는** 이유(§6).
+
+---
+
+## 5. Case별 커널
+
+네 case 모두 §3(tile sweep) + §4(reduction)를 공유하며; `N`, query-block 폭,
+어느 `tile_*` predicate가 발동하는지에서만 다르다.
+
+### 5.1 공통 skeleton
+
+```python
+def attn_kernel(q_ptr, k_ptr, v_ptr, o_ptr, counter, start_pe, N,
+                q_block, softmax_scale, *, tl):
+    pe_id   = tl.program_id(axis=rank_axis)
+    my_len  = valid_len(counter, start_pe, pe_id, N)
+    n_tiles = ceil(my_len / TILE)
+    q_g     = load_Q_group(q_ptr)            # [G·T_q, d]; lazy tl.load, G folded into M (no broadcast)
+    m, l, O = init_running()                 # persistent arena: -inf, 0, zeros
+    for j in range(n_tiles):                 # K/V streamed per tile by the composite scheduler (§3)
+        if tile_all_future(j, q_block):      # causal skip (kernel if)
+            continue
+        run_tile(j, mask_or_null(j, q_block))   # §3: 2 composites + softmax MATH, in tl.scratch_scope
+    if N == 1:
+        tl.store(o_ptr, O / l)               # no reduction
+    else:
+        tree_reduce_and_store(pe_id, N, o_ptr)  # §4
+```
+
+### 5.2 DECODE, no SP (`N=1`, 한 PE가 head의 KV 소유)
+- `T_q = 1`, 모든 과거 KV에 attend ⇒ future tile 없음, 마지막 ragged tile만 mask.
+- **GQA 재사용이 전부**(decode는 KV-load-bound): KV head의 `G=8` query 행을 matmul
+  **M** 차원에 fold(`q_g`를 `[G, T_q, d] → [G·T_q, d]`로 reshape, byte-보존).
+  그러면 `Q·Kᵀ`는 `composite([G·T_q, d], Kᵀ[d, TILE]) → [G·T_q, TILE]`이고
+  `P·V`는 `composite([G·T_q, TILE], V[TILE, d]) → [G·T_q, d]`. KV tile(`[TILE, d]`)이
+  공유 `K`/`V` operand — **한 번 스트리밍되어 모든 `G·T_q` 행이 자동 재사용** —
+  그것들이 GEMM의 M 행이기 때문. K/V broadcast 불필요; composite의 tile plan 내
+  `m = G·T_q`가 timing이 모든 `G` 행의 작업을 올바르게 세게 한다(leading batch
+  축은 세어지지 *않음* — §8 참조).
+- `S_pe`가 scratch에 맞으면(작은/중간 context) 이것은 **one-shot** partial
+  attention으로 축약(Q·Kᵀ용 composite 하나, softmax 하나, P·V용 composite 하나) —
+  바로 baseline `_attention_mesh_mlo`의 `_partial_attention`, 단지 GQA-batched에
+  composite 경로. Tiling(§3)은 `S_pe`가 scratch scope의 tile 예산을 초과할 때만
+  발동.
+
+### 5.3 DECODE, with SP / KV-parallel (`N=8`)
+- 한 request의 KV가 8 PE에 round-robin; 각각 ≈`my_len` token 소유. 각 PE가 자기
+  로컬 tile에 대해 `G=8` query 행을 GQA-batch.
+- `T_q=1` ⇒ 짧은 per-PE sweep ⇒ **reduction 지배** ⇒ §4 tree(3단계)가 구조적으로
+  중요한 부분. Reduction latency는 batch될 때 다른 동시 decode token으로, 혹은 긴
+  single-stream context의 긴 per-PE tile sweep으로 숨겨진다.
+
+### 5.4 PREFILL, no SP
+- 전체 prompt 상주; query는 `T_q` token 블록(chunk).
+- causality가 실재, query block `[qs, qe)` vs KV tile `[ks, ke)`:
+  - `ke ≤ qs` → `tile_all_past` → `mask=None`, full compute.
+  - `ks ≥ qe` → `tile_all_future` → **skip**(커널 `if`).
+  - overlap → `tile_partial` → 커널이 삼각 additive mask 생성, tile op 시퀀스가
+    더함(`Sj + mask_j`).
+- GQA는 group의 query head를 decode와 같은 축으로 batch.
+
+### 5.5 PREFILL, with SP (Ring KV)
+- KV가 `N` PE(ring)에 sharded; 각 step이 peer의 KV 블록을 IPCQ로 전달. 실행 중
+  `(m,ℓ,O)`는 **ring step을 가로질러** carry(Python 핸들, `_attention_mesh_kv`가
+  오늘 하는 그대로).
+- IPCQ가 다음 step의 KV 수신을 현재 step의 compute와 `tl.recv_async`/`tl.wait`로
+  오버랩(`tl_context.py:543-560` — 이미 존재).
+- **Causal ring skip:** 들어오는 step의 KV 블록이 전부 로컬 query block *이후*면 그
+  compute를 skip(recv-consume 안 함 / fold 안 함) — causal 어텐션에서 약 절반 step
+  제거.
+- 수신 버퍼는 ping-pong(persistent arena, 재활용 안 함).
+- ring 후, `(m,ℓ,O)`는 query-소유 PE별 최종; KV-parallel 분할이 남으면 §4가
+  tree-merge한 뒤 정규화 + store.
+
+baseline `_attention_mesh_kv`가 이미 ring fold를 구현한다; 본 ADR은 GQA 재사용,
+causal step-skip, `recv_async` 오버랩을 추가한다.
+
+---
+
+## 6. 왜 하드웨어 / composite 변경이 불필요한가
+
+- HW/composite `pop`-as-dependency가 도움 될 유일한 곳은 poll을 숨길 다른 작업이
+  없는 외로운 reduction — 즉 batch=1 **그리고** 짧은 context. 타깃은 **agentic =
+  낮은 batch, 긴 context** ⇒ 각 PE가 많은 KV tile을 가짐; §4 tree의 `tl.recv`
+  blocking은 동시 token / 긴 context의 sweep 작업으로 가려진다.
+- §4 reduction은 **정적** tree를 쓰므로 recv 순서가 컴파일 타임에 고정 —
+  data-dependent pop 없음. `tl.recv`(blocking)로 충분.
+- **결정: greenlet `tl.send`/`tl.recv` collective로 출시.** 짧은-context,
+  single-stream, latency-critical 타깃이 나타날 때만 HW `pop` 재검토.
+
+---
+
+## 7. 제어 vs 실행 분할 (load-bearing 원칙)
+
+| Concern | Owner |
+|---|---|
+| tile skip (future), mask 생성, causal bound | **커널** (greenlet 본문의 Python `if` + 산술) |
+| 주소 / offset / valid-length 산술 | **커널** (counter에서) |
+| reduction 스케줄링, IPCQ send/recv 순서 | **커널** (정적 tree) |
+| Q·Kᵀ, P·V (per-tile K/V DMA 스트리밍 + tiling 포함) | **`tl.composite`** → PE_SCHEDULER |
+| Q load, mask add, softmax math, 실행 중 `(m,ℓ,O)` 머지 | **`tl` op** PE 엔진 상 (커널 발행) |
+
+커널이 결정하고; GEMM은 composite로 scheduler에 offload되며; 나머지 `tl` op은
+이미 결정된 작업을 실행한다. 이는 kernbench가 이미 지원하는 greenlet + composite
+모델 그대로 — 새 제어 추상화 없음.
+
+---
+
+## 8. 필요한 kernbench 변경
+
+**설계 iteration의 정정:** 진짜 GQA(`h_q > h_kv`)는 **신규 primitive 불필요** —
+§5.2의 커널 재구조화(KV head별, `G`를 M에 fold, byte-보존 reshape)만 필요. 보조
+ADR들은 *efficiency / scale* enabler이지 GQA blocker가 아니다.
+
+**커널 내 알고리즘 작업 (신규 primitive 없음; 기존 `tl` API):**
+
+- **GQA Q축 batching**(재사용 레버) — KV head별로 `G·T_q`를 matmul M 차원에 fold
+  (§5.2); `_view` 식 byte-보존 reshape; GEMM은 M = `G·T_q`인
+  `tl.composite(op="gemm")`. 오늘 timing/data 모드 모두 동작.
+- **composite를 통한 GEMM**(§1/§3) — Q·Kᵀ와 P·V를 각각 non-blocking
+  `tl.composite(op="gemm")`로 발행; PE_SCHEDULER가 tiling하고 `tl.ref` K/V
+  operand의 DMA를 스트리밍(기존 `CompositeCmd`; 새 커맨드 종류 없음).
+- root로의 tree reduction(§4)이 baseline all-to-all fan-out 대체 — `tl.send`/
+  `tl.recv` 상의 순수 커널 제어 흐름.
+- causal tile skip + additive boundary mask(§3/§5.4) — 커널 `if` + `+`로 더한
+  mask 텐서.
+- round-robin KV placement / valid-length 산술(§2) — launch-arg 산술 +
+  `DPPolicy(pe="row_wise")`.
+
+**efficiency / scale을 위한 신규 primitive (각각 보조 ADR 있음):**
+
+1. **per-tile scratch 재활용** — **ADR-0063**(`tl.scratch_scope`). *스케일에 필요*:
+   `S=16` 상한(1 MiB bump allocator)을 제거해 현실적 context 길이가 돌게 함. 셋 중
+   최고 가치.
+2. **Lazy `tl.load`** — **ADR-0062**(non-blocking load + 첫 사용 시점 auto-wait; API
+   표면 불변). *효율*: 명시적 load(Q group, 비-composite 커널)를 뒤따르는 compute와
+   오버랩. 타일당 **K/V** prefetch는 composite scheduler가 처리(§1)하므로, 이것은
+   나머지 명시적 load를 담당. 전역 시맨틱 변경 → 기존 골든 재생성(ADR-0062 D3).
+3. **GQA head / mask broadcast** — **ADR-0061**(`tl.broadcast`). *선택적 편의*, GQA
+   blocker 아님(위 정정 참조). `G·T_q` 행에 걸친 additive-mask 구성과 범용 커널에
+   유용; `np.matmul`이 data 모드에서 이미 broadcast하므로 정확성엔 불필요. 최저
+   우선순위.
+
+**명시적 REJECTED (효율적 대안 선택):**
+
+4. ~~내부 루프 전체를 내재화하는 전용 "flash-composite" 커맨드 종류 —
+   DMA→MM→VEC→DMA→MM→VEC + carried `(m,ℓ,O)` register 상태 + 꼬리 IPCQ push.~~ 두
+   GEMM은 기존 `CompositeCmd`를 **쓴다**(§1/§3) — 그것이 scheduler 관리 tiling, K/V
+   DMA 스트리밍, cross-tile 파이프라이닝을 준다. 기각되는 것은 softmax 머지 +
+   cross-tile register 수명 + IPCQ-push epilogue까지 흡수하는 **새** 커맨드 종류다:
+   크고 특수목적이며, 하이브리드 대비 유일한 delta(완전 softmax offload)가 현재
+   모델링 충실도(per-op CPU 발행 비용 = 0; §1 참조)에서 정당화되지 않는다. cost
+   model(§9)이 완전 offload를 측정 가능하게 가치 있게 만들면 재검토.
+5. ~~하드웨어 `pop`-as-dependency.~~ 범위 밖(§6).
+6. ~~본 커널 내 RoPE / QKV projection / KV-cache write.~~ Upstream `qkv_rope`
+   (P1–P5). RoPE를 fold하면 매 decode step마다 과거 tile을 재회전해야 하고 Ring
+   Attention의 post-RoPE pass-through를 깬다.
+
+---
+
+## 9. Open 튜닝 항목 (kernbench에서 측정, blocking 아님)
+
+1. **Composite tile-pipeline depth** — KV-load-bound에 지배적; 커널이 wait 전에
+   non-blocking composite를 얼마나 앞서 발행하는지, 그리고 scheduler의 타일당
+   스트리밍 depth.
+2. **reduction tree의 PE↔mesh-이웃 매핑** — 각 depth-`⌈log₂ N⌉` pair가 물리
+   `N/S/E/W` 이웃인지 보장; 나쁜 매핑은 hop 추가. SFR install 대조 검증.
+3. **TILE 크기** — scratch 상주(`S/P tile + O_acc + G-way GQA`)와 DMA 효율 균형;
+   ADR-0063 및 scheduler의 `TILE_M/K/N`(`pe_scheduler.py`)과 상호작용.
+4. §5.5의 **Ring 버퍼 ping-pong vs `recv_async` depth**.
+5. **decode의 one-shot vs tiled 교차점**(§5.2) — tiling이 단일 composite를 이기는
+   `S_pe` 임계값.
+6. **per-op CPU 발행 비용 (cost model)** — 현재 `dispatch_cycles=0`(`pe_cpu.py`)이라
+   composite-vs-primitive 발행 오버헤드가 보이지 않음. op 종류별 차등 발행 비용
+   (`tl.composite` descriptor push ≫ primitive op)이 하이브리드의 CPU-saturation
+   이점을 **측정 가능**하게 한다(§1). **ADR-0064**에 명세; 별도 future work로 추적.
+
+---
+
+## 10. Coverage 요약
+
+| Case | KV placement | Inner loop | Reduction | Masking |
+|---|---|---|---|---|
+| Decode, no SP | 1 PE, all KV | one-shot or tiled, GQA-batched | none | last tile only |
+| Decode, SP | round-robin `N` PEs | tiled, GQA-batched | §4 tree (`⌈log₂ N⌉`) | last tile only |
+| Prefill, no SP | resident | tiled per q-block | none | skip future / triangular boundary |
+| Prefill, SP (Ring) | ring-sharded | per ring step (recv_async overlap) | running state + §4 tree | causal step-skip + boundary |
+
+**Case별 I/O** (전체 계약 §0.5):
+
+| Case | Inputs | Output | IPCQ partials |
+|---|---|---|---|
+| Decode, no SP | `Q[G,1,d]`, full `K/V[S,d]` on 1 PE | `O[G,1,d]` at `O_base` | none |
+| Decode, SP | `Q[G,1,d]`, per-PE `K/V[S_pe,d]` | `O[G,1,d]` (root) | `(m,ℓ,O_i)` per PE → tree |
+| Prefill, no SP | `Q[G,T_q,d]`, `K/V[≤end,d]` | `O[G,T_q,d]` at `O_base` | none |
+| Prefill, SP (Ring) | `Q[G,T_q,d]`, ring-delivered `K/V` blocks | `O[G,T_q,d]` (root) | running `(m,ℓ,O)` + tree |
+
+모든 case에서: KV-cache write와 RoPE는 **upstream**에서; output projection은
+**downstream**에서; score `S`와 prob `P`는 미생성.
+
+---
+
+## 11. 검증 계획 (Phase 1 테스트 개요)
+
+SPEC/ADR coverage: R5 (PE↔PE IPCQ, PE↔HBM), R2 (traversal에 의한 latency),
+ADR-0023/0025 (IPCQ), ADR-0046 (`tl` 계약), ADR-0054 (eval bench).
+
+시뮬레이터의 계약은 bit-exact 수치가 아니라 **traversal에 의한 latency + 결정성 +
+구조적 정확성**(SPEC §0, §0.1)이다 — Phase 2 데이터는 주로 data path를
+exercise하려 존재하고, `tl.trans`는 reshape-not-transpose이며, `bf16`은 `f16`으로
+모델된다. 따라서 검증은 **구조/타이밍 우선**, 수치 parity는 제한된 2차 체크.
+
+1. **data 모드 실행(`enable_data=True`):** GQA 커널(`h_q = G·h_kv`)이 baseline
+   head-packing이 부딪히는 byte-conservation 에러 없이 — 네 case 모두 — 완료. (이는
+   §5.2 재구조화가 필요하지, 신규 primitive가 *아님*.)
+   **수치 parity(2차):** reshape-as-transpose가 정확한 symmetric/identity 입력에
+   대해, 커널 `O`가 numpy FlashAttention reference와 fp tolerance 내 일치. 완전
+   asymmetric parity는 실제 `tl.transpose`에 좌우됨(범위 밖; 플래그됨).
+2. **GQA 재사용:** `h_q = G·h_kv`에서 K/V `dma_read_count`가 `G`와 무관(타일당 한
+   load, group에 걸쳐 재사용), 반면 GEMM 작업은 `G`로 스케일. 레버가 실제 발동함을
+   assert.
+3. **tree reduction 단계 수:** decode-SP가 `⌈log₂ N⌉` reduction 라운드(`N−1`이
+   아님)를 발행; op_log `ipcq_send`/`recv` 수가 tree와 일치.
+4. **Causal skip:** prefill이 모든 `tile_all_future` tile을 skip — GEMM 수가 full
+   grid가 아니라 하삼각 tile 수와 일치.
+5. **Long context (ADR-0063):** scope 없이 1 MiB를 넘기는 `S` sweep이 완료하고
+   reference와 일치.
+6. **Load/compute 오버랩:** tiled sweep의 end-to-end latency가 직렬
+   `Σ(load+compute)`보다 낮음 — composite scheduler의 타일당 K/V 스트리밍(§1/§3)과
+   Q load를 오버랩하는 lazy `tl.load`(ADR-0062)에서. (오버랩은 실제 모델 동시성,
+   빼기가 아님.)
+7. **Composite GEMM offload (구조적):** 각 tile의 Q·Kᵀ와 P·V가 blocking `tl.dot`이
+   아니라 `CompositeCmd`(non-blocking)를 PE_SCHEDULER에 emit; op_log가 composite
+   tile plan을 보이고 커널이 다음 tile의 composite를 wait 전에 발행(cross-tile
+   파이프라이닝).
+8. **결정성:** 동일 입력 → 동일 op_log + latency (SPEC §0.1).
+
+---
+
+## B. 하이브리드 전환에서 나온 Open 설계 항목 (추후 검토)
+
+이들은 결정이 순수 greenlet primitive 경로에서 **composite hybrid + lazy
+`tl.load`**(이번 개정)로 옮겨갈 때 생겼다. 설계를 막는 것은 없으며; 각각 구현 중
+검증 패스가 필요하다. (묻지 않고) 작업 합의에 따라 여기 기록한다 — 권고는 내가
+예측한 기본값이며; 검토 시 수정.
+
+1. **DDD-0060이 아직 미동기화.** Detailed Design Document는 여전히 옛
+   `tl.load_async` 더블버퍼 경로와 primitive `tl.dot` 내부 루프를 기술한다(그
+   §4.3/§5/§10). 하이브리드(composite GEMM, lazy load, K 사전 transpose)로
+   갱신해야 한다. DDD가 파생 how-to이고 큰 재작성이라 *검토용으로 남김*; ADR이 이제
+   권위 있는 기록이다. **권고:** 구현 시작 전 후속으로 DDD 동기화.
+
+2. **composite GEMM의 K operand 방향.** Q·Kᵀ는 `b = [d, TILE]`가 필요하나 KV
+   cache는 K를 `[S_pe, d]`로 저장한다. `tl.trans`는 메타데이터 전용이고
+   `MemoryStore.read`는 transpose가 아니라 reshape(`memory_store.py:73`) — 비자명
+   데이터엔 런타임 transpose가 틀리다. **권고:** cache에 K를 **사전 transpose**
+   `[d, S_pe]`로 저장(pseudocode와 §3이 이를 가정)하여 `tl.ref(k_tile, (d, TILE))`이
+   연속 slice가 되게. upstream `qkv_rope` write 레이아웃이 이를 지원하는지 검증,
+   아니면 실제 `tl.transpose` 추가(더 무거움; 연기).
+
+3. **composite 출력 버퍼 vs `tl.scratch_scope`.** 각 Q·Kᵀ composite는 `Sj`를
+   `out_addr`에 쓰고; 커널이 softmax MATH용으로 그것을 읽는다. 그 출력 버퍼와
+   in-flight composite의 타깃은 per-tile `scratch_scope`(ADR-0063)가 소비 전에
+   재활용하지 **않는** 곳에 살아야 한다 — in-flight lazy load와 같은 규율(ADR-0062
+   D-Negative). **권고:** *현재* tile의 composite 출력은 scoped arena에(같은
+   iteration에 소비); persistent `(m,ℓ,O)`는 바깥에. 다음 tile의 composite를 일찍
+   발행할 때(cross-tile 파이프라이닝) use-after-recycle이 없음을 검증.
+
+4. **composite 스트리밍 하의 GQA `dma_read_count` 레버.** 레버(§11.2: K/V
+   `dma_read_count`가 `G`와 무관)는 composite가 모든 `G·T_q` M-행에 재사용되는
+   **하나의** K/V tile DMA를 emit한다고 가정한다. scheduler의 `generate_gemm_plan`은
+   `TILE_M/K/N`(`pe_scheduler.py:35-37`, 32/64/32)으로 tiling한다 — `G·T_q`에 대한
+   M-tiling이 M-tile마다 공유 K/V tile DMA를 재발행하지 **않음**을 확인(즉 operand
+   DMA가 M-tile에 걸쳐 공유되거나, 아니면 레버가 약해짐). **권고:** levers 테스트에서
+   assert; 위반 시 GQA 이득은 DMA가 아니라 compute에만 — 여전히 정확하나 헤드라인이
+   바뀜.
+
+5. **커널 TILE vs scheduler `TILE_M/K/N`.** 커널은 논리적 KV `TILE`을 다루고;
+   scheduler는 고정 `TILE_M/K/N`으로 내부 재tiling한다. 두 tiling 레이어가
+   상호작용(scratch 상주, 파이프라인 depth). **권고:** 커널 TILE을 K/V 스트리밍
+   granularity로 보고 scheduler가 GEMM을 sub-tile하게; DDD에 관계를 문서화하고 둘 다
+   sweep(§9 항목 1, 3).
+
+6. **cost model은 별도 ADR.** 하이브리드의 CPU-saturation 이점은
+   `dispatch_cycles=0`인 동안 보이지 않는다. op 종류별 issue-cost 모델은
+   **ADR-0064**에 명세; 본 ADR의 §1/§9는 *측정 가능한*(구조적뿐 아닌) 이점을 위해
+   그것에 의존. **권고:** eval에서 하이브리드 latency 이점을 주장하기 전에 ADR-0064
+   모델을 도입.
+
+7. **Ring 경로(§5.5) GEMM.** §5.5는 여전히 primitive op + `recv_async`로 ring
+   fold를 기술한다. 일관성을 위해 ring의 per-step Q·Kᵀ / P·V도 composite여야 하고;
+   IPCQ `recv_async` 오버랩은 직교하며 유지. **권고:** 구현 중 ring step에 같은
+   하이브리드 형태 적용; 낮은 리스크, §3 미러.

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

@@ -0,0 +1,817 @@
+# ADR-0060: AHBM GQA Fused Attention Kernel (Llama3-70B)
+
+## Status
+
+Proposed
+
+**Context model:** Llama3-70B.
+**Decision drivers:** agentic workload → low batch, long context;
+KV-load-bound decode; sequence-parallel (Ring KV) for long-context prefill.
+
+**Supersedes / extends:** the existing mesh-native attention kernels
+`_attention_mesh_kv` (prefill) and `_attention_mesh_mlo` (decode) and the
+`milestone-gqa-llama70b` eval bench. See *§A. Relationship to existing
+kernbench work* — that code is the baseline this ADR upgrades to a real
+GQA, causal, long-context kernel.
+
+**Supporting ADRs** (efficiency / scale enablers — *not* GQA blockers;
+see §8 correction): **ADR-0063** `tl.scratch_scope` (per-tile scratch
+recycling — required for realistic context length), **ADR-0062** lazy
+`tl.load` (non-blocking load with auto-wait on first use → load/compute
+overlap), **ADR-0061** `tl.broadcast` (optional mask/general
+convenience). The two GEMMs (Q·Kᵀ, P·V) are issued as scheduler-managed
+`tl.composite` commands (the existing `CompositeCmd`; no new command
+kind); real GQA itself needs only kernel restructuring (§5.2).
+
+**Algorithm lineage.** This kernel is **FlashAttention** (tiling +
+online/streaming softmax with fused P·V — no full score matrix
+materialised). The KV-parallel split-and-combine in §4 is
+**FlashDecoding** (split-KV with log-sum-exp merge). The Ring path in
+§5.5 is **Ring Attention** (KV blocks rotated around the mesh, folded by
+the same online softmax). No new math is introduced; this ADR maps those
+known algorithms onto the kernbench **greenlet `tl` programming model**
+(ADR-0020, ADR-0046) and the **IPCQ** PE↔PE collective (ADR-0023/0025).
+
+---
+
+## TL;DR — final GQA kernel (composite hybrid) pseudocode
+
+The decision (§1) in one place: **GEMMs → `tl.composite`; softmax merge +
+tree reduction → kernel; `tl.load` is lazy.** Per-KV-head, `G` folded into
+the matmul M dim (real GQA, no broadcast). This is the reference shape; the
+prose sections elaborate each piece.
+
+```python
+def gqa_attention(q_ptr, k_ptr, v_ptr, o_ptr,
+                  counter, start_pe, N, q_block, scale, *, tl):
+    # ---- geometry (kernel arithmetic, §2) ----
+    pe_id   = tl.program_id(axis=rank_axis)
+    G       = H_q // H_kv                       # query heads per KV head (=8)
+    causal  = q_block.is_prefill
+
+    for kv in range(H_kv):                       # one KV head per iteration (§5.2)
+        # Q group: G rows folded into M → [G·T_q, d]. Lazy load (ADR-0062):
+        # issue now, auto-wait at first use inside the first composite.
+        q_g     = tl.load(q_base(kv), (G * q_block.T_q, d))      # [G·T_q, d]
+        my_len  = valid_len(counter, start_pe, pe_id, N) if not causal else S_kv
+        n_tiles = ceil(my_len / TILE)
+
+        # persistent arena (outside scratch_scope, ADR-0063): -inf, 0, zeros
+        m, l, O = init_running(G * q_block.T_q, d)
+
+        for j in range(n_tiles):
+            if causal and tile_all_future(j, q_block):
+                continue                                         # causal skip (kernel if)
+            with tl.scratch_scope():                             # per-tile temporaries (ADR-0063)
+                # --- GEMM #1 on the scheduler: Q·Kⱼᵀ; K streamed by composite ---
+                Sj = tl.composite("gemm", a=q_g,
+                                  b=tl.ref(k_tile(kv, j), (d, TILE))) * scale   # Kᵀ pre-stored [d,TILE]
+                if causal and tile_partial(j, q_block):
+                    Sj = Sj + causal_mask(j, q_block)            # additive boundary mask
+                # --- online softmax merge in the kernel (MATH ops) ---
+                m_new = tl.maximum(m, tl.max(Sj, axis=-1))
+                P     = tl.exp(Sj - m_new)
+                corr  = tl.exp(m - m_new)
+                l     = l * corr + tl.sum(P, axis=-1)
+                # --- GEMM #2 on the scheduler: P·Vⱼ; V streamed by composite ---
+                Oj    = tl.composite("gemm", a=P,
+                                     b=tl.ref(v_tile(kv, j), (TILE, d)))
+                O     = O * corr + Oj                            # running merge
+                m     = m_new
+
+        # ---- cross-PE combine (§4): log-sum-exp tree to root, or store ----
+        if N == 1:
+            tl.store(o_base(kv), O / l)
+        else:
+            tree_reduce_and_store(m, l, O, pe_id, N, o_base(kv))  # tl.send/tl.recv
+```
+
+> **Why this shape:** the two `tl.composite` calls offload tiling + K/V DMA
+> streaming + cross-tile pipelining to PE_SCHEDULER (CPU issues coarse
+> descriptors and runs ahead → engines stay saturated, §1); the softmax
+> merge and the IPCQ tree reduction stay in the kernel because the existing
+> `CompositeCmd` cannot carry cross-tile state. `K` is pre-stored
+> transposed `[d, TILE]` to sidestep the reshape-not-transpose caveat
+> (§3, §B). A bespoke "flash-composite" command kind is **not** introduced
+> (§8 item 4).
+
+---
+
+## A. Relationship to existing kernbench work (read first)
+
+kernbench **already runs FlashAttention with an online-softmax `(m, ℓ, O)`
+merge over IPCQ today.** Two kernels exist:
+
+| File | Role | Mechanism |
+|---|---|---|
+| `src/kernbench/benches/_attention_mesh_kv.py` | prefill (Ring K/V) | per-rank partial attention, bidirectional K/V fan-out, online-softmax fold |
+| `src/kernbench/benches/_attention_mesh_mlo.py` | decode (split-KV) | per-rank one-shot partial attention, bidirectional `(m,ℓ,O)` fan-out, log-sum-exp merge |
+
+Both are driven by `milestone-gqa-llama70b` (4 panels:
+`{single,multi}_user × {prefill,decode}`,
+`src/kernbench/benches/milestone_gqa_llama70b.py`) and tested in
+`tests/attention/test_milestone_gqa_llama70b.py`.
+
+**They are written in the greenlet `tl` API:** `tl.load`, `tl.dot`,
+`tl.softmax`/`tl.max`/`tl.sum`/`tl.exp`, `tl.send`/`tl.recv`, and Python
+`-`/`*`/`/` on `TensorHandle` (each emits a `MathCmd`) — the GEMMs as
+**blocking `tl.dot`**, not composites. The running `(m, ℓ, O)` is just
+Python `TensorHandle`s threaded through the loop. This ADR keeps the
+running state and the softmax merge in the kernel but **moves the two
+GEMMs onto the scheduler-managed `tl.composite` path** (see §1) — this
+**matters for the design**.
+
+**Three deliberate limitations of the baseline** — exactly what an
+*efficient GQA* kernel must lift:
+
+1. **No GQA reuse.** `h_q == h_kv == 1`
+   (`test_milestone_gqa_llama70b.py:137-142`). The test attributes this to
+   a MemoryStore byte-conservation failure on a *broadcast view*, but that
+   failure is a property of the baseline's **head-packing hack**
+   (`_view(K, (h_q·d, S_kv))`, which conflates all heads into one matmul
+   dim and only conserves bytes when `h_q == h_kv`). The correct fix is
+   **kernel restructuring**, not a broadcast op: process **one KV head at
+   a time** and fold the `G` group rows into the matmul **M** dimension
+   (§5.2). That uses only byte-conserving reshapes, so real GQA
+   (`h_q = G·h_kv`) runs with **no new primitive** — see §8.
+2. **O(N) reduction.** The baseline does an all-to-all bidirectional
+   fan-out so *every* rank ends with the full answer (`n_ranks − 1`
+   steps). Attention only needs `O` at the query owner → a **tree
+   reduction** to root is `⌈log₂ N⌉` steps (§4).
+3. **Validation scale only.** `S = 16` because the 1 MiB scratch bump
+   allocator leaks per-tile temporaries
+   (`test_milestone_gqa_llama70b.py:123-148`) and there is no causal
+   masking / tiling → fixed by **ADR-0063** (recycling) + §5 (tiling,
+   causal skip) + composite K/V streaming (§3) + **ADR-0062** (lazy load
+   overlap).
+
+**Documentation debt (out of scope but recorded):** the baseline cites
+ADR-0055/0056/0057/0058/0059, **none of which exist as files** — they are
+ghost references. This ADR does not retro-write them; see the Detailed
+Design Document's *Open Decisions* for the recommendation.
+
+---
+
+## 0. Reference dimensions (Llama3-70B)
+
+| Symbol | Meaning | Value |
+|---|---|---|
+| `H_q` | query heads | 64 |
+| `H_kv` | KV heads | 8 |
+| `G` | GQA group size = `H_q / H_kv` | 8 |
+| `d` | head dim | 128 |
+| `L` | layers | 80 |
+| `D` | model dim | 8192 |
+
+Hardware recap:
+- **AHBM (chip)** = set of **CUBE**s (memory cubes, each with a logic die
+  containing **PE**s) + **IO die** (ADR-0003).
+- **IPCQ**: PE↔PE queues, 4 mesh-direction queue-pairs per PE
+  (`N/S/E/W`, ADR-0023 D3; `global_*` for inter-SIP, ADR-0032). Kernel
+  API: `tl.send(dir, src)` / `tl.recv(dir, shape, dtype)`
+  (`tl_context.py:402-499`).
+- **Composite command** (`CompositeCmd`, `pe_commands.py:144-162`): a
+  single GEMM (or MATH) *head* plus element-wise *epilogue* stages
+  (`bias/relu/scale/add/...`), issued **non-blocking** to PE_SCHEDULER,
+  which generates a tile plan and streams DMA→GEMM→write per tile
+  (ADR-0014 D6; `pe_scheduler.py:104-143`). It is **not** a general
+  multi-op DAG: it cannot chain two GEMMs, cannot carry register state
+  across instances, and cannot pop/wait on IPCQ. This ADR therefore issues
+  **each** of the two attention GEMMs (Q·Kᵀ, P·V) as its own composite and
+  keeps the cross-GEMM softmax merge + the IPCQ reduction in the kernel —
+  it does **not** need a new "flash-composite" command kind (see §1, §8).
+- **Allocation policy (SP):** one query head per CUBE in the multi-user
+  panels; the `G` query heads of a KV group map within a CUBE's PEs.
+
+**Two orthogonal mapping layers** (round-robin placement):
+- **Layer 1 (KV-parallel, intra-request):** KV token `i` lands on PE
+  `(start_pe + i) mod N` — round-robin so one long request is split
+  across `N` PEs, balanced to ≤1 token.
+- **Layer 2 (inter-request):** `start_pe = request_id mod N` rotates the
+  "PE-0 role" per request.
+- Combined: `pe = (request_id + global_token_idx) mod N`.
+
+`N` = number of PEs in the KV-parallel reduction group for one
+(query head, request). Reduction stays **intra-CUBE** (a query head never
+spans CUBEs — explicit non-goal).
+
+---
+
+## 0.5 Kernel boundary, preconditions, I/O contract
+
+### 0.5.1 Position in the decoder layer
+
+```
+1. RMSNorm
+2. QKV projection (GEMM)        ─┐  qkv_rope kernel (SEPARATE, upstream)
+3. RoPE on Q and K              ─┤
+4. write new K,V → KV cache     ─┘
+5. ===== THIS KERNEL: FlashAttention =====   (post-RoPE Q, K-cache, V-cache → O)
+6. Output projection (GEMM)        out_proj kernel (SEPARATE, downstream)
+7. residual add → FFN ...
+```
+
+**Preconditions (upstream `qkv_rope`, NOT this kernel):**
+- **P1.** Q is already RoPE-rotated. No rotation here.
+- **P2.** K-cache stores **post-RoPE** K (never re-rotated at attention
+  time — the reason for post-RoPE caching).
+- **P3.** For decode, the new step's K row is RoPE-rotated and appended
+  to its owning PE's K-cache slot **by `qkv_rope` before this kernel
+  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·N + ((pe_id − start_pe) mod N)`
+  (§2.1). Round-robin placement ≠ RoPE position.
+
+### 0.5.2 Shape symbols
+`T_q` = query length this launch (decode: 1; prefill/chunk: chunk width).
+`S` = total context length. `S_pe` = keys owned by this PE ≈ `⌈S/N⌉`.
+Storage `bf16` (numpy proxy `f16`, `memory_store.py:16`); `m,ℓ,O`
+accumulators `f32` where precision matters.
+
+### 0.5.3 INPUTS (per kernel launch)
+
+| Input | Shape (per KV head) | Location | Notes |
+|---|---|---|---|
+| `Q` | `[G, T_q, d]` | per-PE HBM (loaded to TCM) | post-RoPE (P1). `G` query rows of the group batched. |
+| `K_cache` | `[S_pe, d]` | per-PE HBM, base `K_base[pe]`, contiguous | post-RoPE (P2). Read-only. Dense locally, strided in global index (§2.1). |
+| `V_cache` | `[S_pe, d]` | per-PE HBM, base `V_base[pe]` | raw V (P4). Read-only. |
+| `global_token_counter` | scalar | launch arg | kernel derives `S_pe`, slot↔global, causal bounds. |
+| `start_pe` (= `request_id mod N`) | scalar | launch arg | Layer-2 rotation. |
+| `pe_id`, `N` | scalars | launch / `tl.program_id` | reduction-group geometry. |
+| `q_block_meta` `{q_start, T_q}` | launch arg | prefill/SP causal masking & skip. |
+| `O_base` | address | launch arg | where final O is written (root PE only). |
+| `softmax_scale` | scalar | launch arg | `1/√d`. |
+
+For **Ring Attention (§5.5)** add per ring step: incoming `K_block,
+V_block` via IPCQ into ping-pong buffers (post-RoPE), plus
+`step_kv_global_range` for the causal step-skip.
+
+### 0.5.4 OUTPUTS
+
+| Output | Shape | Location | Notes |
+|---|---|---|---|
+| `O` (final) | `[G, T_q, d]` | `O_base`, **root PE only** | normalised `O_acc / ℓ_acc`, cast to storage dtype. |
+
+**Intermediate (non-root PE, on IPCQ — not a kernel-visible output):**
+`(m_i, ℓ_i, O_i)` (`m,ℓ`: `[G, T_q]`; `O_i`: `[G, T_q, d]` unnormalised)
+pushed to `tree_parent` via `tl.send` (§4). `O_i` is the heavy part.
+
+**No-SP (`N=1`):** no IPCQ partials; the single PE's running `(m,ℓ,O)` is
+normalised in place and written to `O_base`.
+
+**Explicitly NOT outputs:** KV-cache writes (upstream, P3), output
+projection (downstream), score `S` and probs `P` (never materialised).
+
+---
+
+## 1. Decision (mechanism)
+
+**Implement the kernel as a *hybrid*: issue the two GEMMs (Q·Kᵀ and P·V)
+as scheduler-managed `tl.composite(op="gemm")` commands, and keep the
+online-softmax merge and the cross-PE reduction as kernel-level `tl`
+ops.** `tl.load` is **lazy** (non-blocking; the wait is auto-inserted at
+first use of the loaded data — ADR-0062), so explicit HBM loads overlap
+the compute that follows. Rationale grounded in kernbench's execution +
+latency model:
+
+- **GEMM tiling is offloaded to PE_SCHEDULER.** A `CompositeCmd` is
+  non-blocking (`kernel_runner.py:182-191`, `pe_scheduler.py:104-121`):
+  the kernel pushes **one coarse descriptor** (M = `G·T_q`, the whole
+  per-PE tile sweep) and the scheduler generates the tile plan and streams
+  DMA→GEMM→write per tile (ADR-0014 D6). K/V are `tl.ref` operands the
+  scheduler streams from HBM, so per-tile **K/V prefetch is the
+  scheduler's job** — no explicit prefetch op. The CPU (greenlet) is freed
+  to issue the **next** composite while the current one runs, so the
+  scheduler keeps the GEMM engine saturated across tiles.
+- **This reflects the hardware** and decouples CPU issue-rate from
+  execution-rate. The blocking per-op `tl.dot` path, by contrast, stalls
+  the CPU on every GEMM and leaves GEMM-engine **bubbles** during the
+  interleaved softmax MATH ops; it is realistic only if the CPU can keep
+  up with fine-grained per-tile issue.
+- **The running `(m, ℓ, O)` flash state stays Python `TensorHandle`s**
+  threaded through the loop (the baseline already does this); the softmax
+  merge (max/exp/sum/rescale) is kernel-level `tl` MATH **between** the two
+  GEMM composites. The existing `CompositeCmd` cannot chain two GEMMs or
+  carry cross-tile register state (§0), so the merge necessarily lives in
+  the kernel — this is the hybrid split, not a limitation worked around.
+- **Cross-PE combination** is a log-sum-exp **tree** over `(m, ℓ, O)`
+  after P·V, via kernel-level `tl.send`/`tl.recv` (§4) — unchanged.
+
+So the per-tile inner pipeline is:
+
+```
+q_g = tl.load(Q group)                         # lazy; auto-wait at first use
+per tile j:
+  Sⱼ = tl.composite("gemm", a=q_g, b=tl.ref(Kⱼ))  → Sⱼ   # scheduler streams Kⱼ DMA + GEMM
+  Sⱼ += maskⱼ                                          # kernel MATH, boundary tile only
+  online-softmax: mⱼ, m_new, P, corr, ℓ               # kernel MATH
+  Oⱼ = tl.composite("gemm", a=P, b=tl.ref(Vⱼ))   → Oⱼ   # scheduler streams Vⱼ DMA + GEMM
+  O = O*corr + Oⱼ;  m = m_new                         # kernel MATH (running merge)
+```
+
+with each tile's MATH temporaries wrapped in `tl.scratch_scope`
+(ADR-0063) so scratch stays O(1), and the next tile's composites issued
+before the current tile's results are waited on (non-blocking handles) so
+the scheduler pipelines across tiles.
+
+Control flow (tile skip, mask generation, reduction scheduling, address
+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
+> "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
+> issue-rate / DMA-program cost that descriptor offload exists to hide.
+> The hybrid is the faithful representation of an efficient kernel. The
+> **measurable** size of the win (can the CPU saturate the engines for
+> many tiles?) is gated on modelling an op-type-differentiated issue cost,
+> tracked as future work (cost model; §9). Even at `dispatch_cycles=0` the
+> non-blocking composite path fills the GEMM-engine bubbles the blocking
+> `tl.dot` path leaves.
+>
+> **Why this is the efficient choice.** GEMM tiling + DMA streaming +
+> cross-tile pipelining are offloaded to the proven `CompositeCmd`
+> scheduler path; the softmax merge and the proven IPCQ collective stay in
+> the kernel. The only genuinely new machinery is two small, general
+> primitives (ADR-0062 lazy `tl.load`, ADR-0063 `tl.scratch_scope`); the
+> reduction reuses `tl.send`/`tl.recv`. A bespoke "flash-composite"
+> command (one kind internalising the softmax merge + carried register
+> state + an IPCQ-push epilogue) is **not** built — large, special-purpose,
+> and its only delta over this hybrid (full softmax offload) is not
+> justified at the current modelling fidelity; see §8.
+
+---
+
+## 2. Memory layout & driver responsibilities
+
+### 2.1 KV cache allocation
+- Per-PE KV buffers sized to per-PE max context `⌈max_context / N⌉ × d ×
+  dtype`, for K and V separately. In kernbench these are deployed with a
+  `DPPolicy` whose `pe` (or `cube`) axis is `row_wise` over the sequence
+  dimension (`policy/placement/dp.py`; the baseline uses
+  `DPPolicy(pe="row_wise")` for K/V and `replicate` for Q).
+- Within a PE, assigned tokens are **appended contiguously** (slot
+  0,1,2,…). Global indices are strided (`i, i+N, …`) but the per-PE
+  buffer is dense ⇒ DMA stays contiguous.
+- Slot → global: `global_idx = local_slot·N + ((pe_id − start_pe) mod N)`.
+
+### 2.2 Driver per-launch duties (minimal)
+The driver supplies, per launch, the bases + counter + rotation; the
+kernel derives the rest:
+
+| Launch arg | Purpose |
+|---|---|
+| `K_base[pe]`, `V_base[pe]` | per-PE KV buffer bases (from tensor VA) |
+| `O_base` | attention output destination |
+| `global_token_counter` | current sequence position |
+| `start_pe = request_id mod N` | Layer-2 rotation |
+| `pe_id` (`tl.program_id`), `N` | geometry |
+| `q_block_meta` | prefill/SP query block start + length |
+
+Kernel-derived (plain arithmetic):
+- **my turn to write this step:** `(start_pe + counter) mod N == pe_id`
+- **my valid length:** `base = counter // N; rem = counter % N;
+  my_len = base + (1 if ((pe_id − start_pe) mod N) < rem else 0)`
+- **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 N` is the entire placement policy.
+
+---
+
+## 3. Per-tile op sequence (greenlet `tl`)
+
+One iteration = one KV tile on one PE. The two GEMMs are
+`tl.composite(op="gemm")` (scheduler-managed tiling + K/V DMA streaming);
+the softmax merge is kernel `tl` MATH between them. Real `tl` names
+(`tl_context.py`), with lazy `tl.load` (ADR-0062), `tl.scratch_scope`
+(ADR-0063):
+
+```python
+# running state (persistent arena — allocated once, outside the scope)
+# m: [G, T_q]  l: [G, T_q]  O: [G, T_q, d]
+q_g = tl.load(Q_group_ptr, (G*T_q, d))               # lazy; auto-wait at first use (ADR-0062)
+
+with tl.scratch_scope():                              # per-tile MATH temporaries recycled
+    Sj    = tl.composite("gemm", a=q_g,                # [G·T_q, TILE]; scheduler streams Kⱼ DMA
+                         b=tl.ref(K_base + j*TILE*d, (TILE, d))) * softmax_scale
+    if mask_j is not None:
+        Sj = Sj + mask_j                              # additive causal mask (boundary tile)
+    m_j   = tl.max(Sj, axis=-1)
+    m_new = tl.maximum(m, m_j)
+    P     = tl.exp(Sj - m_new)                        # no full-matrix softmax; streaming
+    corr  = tl.exp(m - m_new)                         # rescale factor for old accumulators
+    l     = l * corr + tl.sum(P, axis=-1)
+    Oj    = tl.composite("gemm", a=P,                  # [G·T_q, d]; scheduler streams Vⱼ DMA
+                         b=tl.ref(V_base + j*TILE*d, (TILE, d)))
+    O     = O * corr + Oj                             # running merge (kernel MATH)
+    m     = m_new
+```
+
+Notes:
+- `q_g` is the GQA-batched query reshaped to `[G·T_q, d]` (the `G` group
+  rows folded into the matmul M dim; byte-conserving). One K/V tile serves
+  all `G·T_q` rows — the GQA reuse lever — with no broadcast.
+- **K/V are `tl.ref` operands** the composite scheduler streams from HBM
+  per tile (`pe_scheduler.py:104-143`): that *is* the prefetch/pipeline,
+  so there is no explicit prefetch op. Issuing tile `j+1`'s composites
+  before waiting on tile `j` (non-blocking handles) keeps the scheduler
+  pipelined across tiles and the GEMM engine saturated.
+- `tl.trans` is **metadata-only** in kernbench (`tl_context.py:390`) and
+  `MemoryStore.read` *reshapes* rather than transposes
+  (`memory_store.py:73`). For zero/structural runs this is harmless; for
+  non-trivial numeric data it yields a reshape-not-transpose, so Q·Kᵀ via
+  a transposed K needs care (§11) — store K pre-transposed `[d, TILE]`, or
+  add a real `tl.transpose` (a candidate further primitive, likely
+  unnecessary given the simulator's performance-modeling purpose).
+- Masking: the kernel builds the boundary-tile mask from query/KV global
+  offsets and adds it (`Sj + mask_j`); full-past tiles pass `None`;
+  full-future tiles are **skipped** (the `if` never enqueues them).
+- A **new** "flash-composite" command kind (one that internalises the
+  softmax merge + carried `(m,ℓ,O)`) is **not** used; the existing
+  `CompositeCmd` covers each GEMM and the merge stays in the kernel
+  (§1, §8 item 4).
+
+---
+
+## 4. Reduction (KV-parallel / SP combine) — tree to root
+
+After its tile sweep each PE holds `(m_i, ℓ_i, O_i)` (unnormalised).
+Combine via the associative/commutative log-sum-exp merge (identical math
+to the baseline's fold, `_attention_mesh_mlo.py:117-122`):
+
+```python
+def merge(m_a, l_a, O_a, m_b, l_b, O_b):
+    m = tl.maximum(m_a, m_b)
+    sa, sb = tl.exp(m_a - m), tl.exp(m_b - m)
+    return m, l_a*sa + l_b*sb, O_a*sa + O_b*sb
+# final: O = O_root / l_root      # normalise once at the tree root
+```
+
+- **Timing:** exchange is **after P·V** (each PE has its final `O_i`).
+- **Topology:** mesh tree, depth `⌈log₂ N⌉`, over physical neighbours.
+  For `N=8`: level-0 `(0↔1,2↔3,4↔5,6↔7)`, level-1 `(1↔3,5↔7)`, level-2
+  `(3↔7)`, root = `7`. Each tree pair must be a physical `N/S/E/W`
+  neighbour so `tl.send(dir)`/`tl.recv(dir)` use a real link
+  (tuning item §9; the SFR install that wires neighbours is
+  `configure_sfr_intracube_pe_ring`, used by the baseline).
+- **Why a tree, not the baseline fan-out:** attention needs `O` at one
+  place (the PE that writes `O_base`). A tree is `⌈log₂ N⌉` steps
+  (3 for `N=8`) vs the baseline all-to-all's `N−1` (7). For decode, where
+  the per-PE sweep is short, the reduction is the dominant cost, so this
+  is a real win. (The baseline's fan-out replicates the answer to all
+  ranks — unnecessary here.)
+- **Payload:** `O_i` (`d`-vector) is heavy; `m,ℓ` are scalars per `(g,
+  T_q)`. Sent as separate `tl.send`s (the baseline sends the triplet as
+  three sends — same pattern).
+
+Kernel structure (greenlet):
+
+```python
+# leaf / internal node: fold children that send to me, then forward up
+for child_dir in tree_children_dirs(pe_id, N):
+    m_c = tl.recv(child_dir, m.shape); l_c = tl.recv(child_dir, l.shape)
+    O_c = tl.recv(child_dir, O.shape)
+    m, l, O = merge(m, l, O, m_c, l_c, O_c)
+if not is_root(pe_id):
+    tl.send(parent_dir(pe_id), m); tl.send(parent_dir(pe_id), l); tl.send(parent_dir(pe_id), O)
+else:
+    tl.store(O_base, O / l)
+```
+
+`tl.recv` blocks until data arrives (`tl_context.py:446-499`); the merge
+order is fixed by the static tree, so no data-dependent `pop` is needed —
+which is why **no hardware/composite `pop`-as-dependency change is
+required** (§6).
+
+---
+
+## 5. Per-case kernels
+
+All four cases share §3 (tile sweep) + §4 (reduction); they differ only
+in `N`, query-block width, and which `tile_*` predicates fire.
+
+### 5.1 Common skeleton
+
+```python
+def attn_kernel(q_ptr, k_ptr, v_ptr, o_ptr, counter, start_pe, N,
+                q_block, softmax_scale, *, tl):
+    pe_id   = tl.program_id(axis=rank_axis)
+    my_len  = valid_len(counter, start_pe, pe_id, N)
+    n_tiles = ceil(my_len / TILE)
+    q_g     = load_Q_group(q_ptr)            # [G·T_q, d]; lazy tl.load, G folded into M (no broadcast)
+    m, l, O = init_running()                 # persistent arena: -inf, 0, zeros
+    for j in range(n_tiles):                 # K/V streamed per tile by the composite scheduler (§3)
+        if tile_all_future(j, q_block):      # causal skip (kernel if)
+            continue
+        run_tile(j, mask_or_null(j, q_block))   # §3: 2 composites + softmax MATH, in tl.scratch_scope
+    if N == 1:
+        tl.store(o_ptr, O / l)               # no reduction
+    else:
+        tree_reduce_and_store(pe_id, N, o_ptr)  # §4
+```
+
+### 5.2 DECODE, no SP (`N=1`, one PE owns the head's KV)
+- `T_q = 1`, attends to all past KV ⇒ no future tiles, mask only on the
+  final ragged tile.
+- **GQA reuse is the whole game** (decode is KV-load-bound): the `G=8`
+  query rows of the KV head are folded into the matmul **M** dimension
+  (`q_g` reshaped `[G, T_q, d] → [G·T_q, d]`, a byte-conserving reshape).
+  Then `Q·Kᵀ` is `composite([G·T_q, d], Kᵀ[d, TILE]) → [G·T_q, TILE]` and
+  `P·V` is `composite([G·T_q, TILE], V[TILE, d]) → [G·T_q, d]`. The KV tile
+  (`[TILE, d]`) is the shared `K`/`V` operand — **streamed once, reused by
+  all `G·T_q` rows automatically** because they are the M rows of the
+  GEMM. No broadcast of K/V is needed; `m = G·T_q` in the composite's tile
+  plan also makes the timing count all `G` rows' work correctly (a leading
+  batch axis would *not* be counted — see §8).
+- If `S_pe` fits in scratch (small/medium context) this degenerates to a
+  **one-shot** partial attention (one composite for Q·Kᵀ, one softmax, one
+  composite for P·V) — exactly the baseline `_attention_mesh_mlo`
+  `_partial_attention`, just GQA-batched and on the composite path. Tiling
+  (§3) only kicks in when `S_pe` exceeds the scratch scope's tile budget.
+
+### 5.3 DECODE, with SP / KV-parallel (`N=8`)
+- One request's KV is round-robin across 8 PEs; each owns ≈`my_len`
+  tokens. Each PE GQA-batches its `G=8` query rows over its local tiles.
+- `T_q=1` ⇒ short per-PE sweep ⇒ **reduction dominates** ⇒ the §4 tree
+  (3 steps) is the structurally important part. Reduction latency is
+  hidden by other concurrent decode tokens when batched, or by the long
+  per-PE tile sweep for long single-stream context.
+
+### 5.4 PREFILL, no SP
+- Whole prompt resident; query is a block of `T_q` tokens (chunked).
+- Causality is real, per query block `[qs, qe)` vs KV tile `[ks, ke)`:
+  - `ke ≤ qs` → `tile_all_past` → `mask=None`, full compute.
+  - `ks ≥ qe` → `tile_all_future` → **skip** (kernel `if`).
+  - overlap → `tile_partial` → kernel builds a triangular additive mask,
+    the tile op sequence adds it (`Sj + mask_j`).
+- GQA batches the group's query heads along the same axis as decode.
+
+### 5.5 PREFILL, with SP (Ring KV)
+- KV sharded across `N` PEs (the ring); each step delivers a peer's KV
+  block over IPCQ. Running `(m,ℓ,O)` is carried **across ring steps**
+  (Python handles, exactly as `_attention_mesh_kv` does today).
+- IPCQ overlaps next step's KV receive with current step's compute via
+  `tl.recv_async`/`tl.wait` (`tl_context.py:543-560` — already exists).
+- **Causal ring skip:** if the incoming step's KV block is entirely
+  *after* the local query block, skip its compute (don't recv-consume /
+  don't fold) — eliminates ≈half the steps for causal attention.
+- Receive buffers ping-pong (persistent arena, not recycled).
+- After the ring, `(m,ℓ,O)` is final per query-owning PE; if KV-parallel
+  splits remain, §4 tree-merges them, then normalise + store.
+
+The baseline `_attention_mesh_kv` already implements the ring fold; this
+ADR adds GQA reuse, causal step-skip, and `recv_async` overlap.
+
+---
+
+## 6. Why no hardware / composite change is needed
+
+- The only place a HW/composite `pop`-as-dependency would help is a lone
+  reduction with no other work to hide a poll — i.e. batch=1 **and** short
+  context. The target is **agentic = low batch, long context** ⇒ each PE
+  has many KV tiles; the §4 tree's `tl.recv` blocks are covered by the
+  sweep work of concurrent tokens / long context.
+- The §4 reduction uses a **static** tree, so the recv order is fixed at
+  compile time — no data-dependent pop. `tl.recv` (blocking) suffices.
+- **Decision: ship with the greenlet `tl.send`/`tl.recv` collective.**
+  Revisit a HW `pop` only if a short-context, single-stream,
+  latency-critical target emerges.
+
+---
+
+## 7. Control vs execution split (the load-bearing principle)
+
+| Concern | Owner |
+|---|---|
+| tile skip (future), mask generation, causal bounds | **kernel** (Python `if` + arithmetic in greenlet body) |
+| address / offset / valid-length arithmetic | **kernel** (from counter) |
+| reduction scheduling, IPCQ send/recv ordering | **kernel** (static tree) |
+| Q·Kᵀ, P·V (incl. per-tile K/V DMA streaming + tiling) | **`tl.composite`** → PE_SCHEDULER |
+| Q load, mask add, softmax math, running `(m,ℓ,O)` merge | **`tl` ops** on PE engines (kernel-issued) |
+
+The kernel decides; the GEMMs are offloaded to the scheduler as
+composites; the remaining `tl` ops execute already-decided work. This is
+exactly the greenlet + composite model kernbench already supports — no new
+control abstraction.
+
+---
+
+## 8. Required kernbench changes
+
+**Correction from design iteration:** real GQA (`h_q > h_kv`) needs **no
+new primitive** — only the kernel restructuring in §5.2 (per KV head,
+`G` folded into M, byte-conserving reshapes). The supporting ADRs are
+*efficiency / scale* enablers, not GQA blockers.
+
+**Algorithm work in the kernel (no new primitive; existing `tl` API):**
+
+- **GQA Q-axis batching** (the reuse lever) — fold `G·T_q` into the matmul
+  M dim per KV head (§5.2); `_view`-style byte-conserving reshape; the
+  GEMM is a `tl.composite(op="gemm")` with M = `G·T_q`. Runs today in both
+  timing and data mode.
+- **GEMMs via composite** (§1/§3) — Q·Kᵀ and P·V each issued as a
+  non-blocking `tl.composite(op="gemm")`; PE_SCHEDULER tiles them and
+  streams the `tl.ref` K/V operands' DMA (existing `CompositeCmd`; no new
+  command kind).
+- Tree reduction to root (§4) replacing the baseline all-to-all fan-out —
+  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 `+`.
+- Round-robin KV placement / valid-length arithmetic (§2) — launch-arg
+  arithmetic + `DPPolicy(pe="row_wise")`.
+
+**New primitives for efficiency / scale (each has a supporting ADR):**
+
+1. **Per-tile scratch recycling** — **ADR-0063** (`tl.scratch_scope`).
+   *Required for scale*: removes the `S=16` ceiling (1 MiB bump
+   allocator) so realistic context lengths run. Highest-value of the
+   three.
+2. **Lazy `tl.load`** — **ADR-0062** (non-blocking load + auto-wait on
+   first use; API surface unchanged). *Efficiency*: overlaps explicit
+   loads (the Q group, non-composite kernels) with following compute. The
+   per-tile **K/V** prefetch is handled by the composite scheduler (§1),
+   so this covers the remaining explicit loads. Global semantics change →
+   existing goldens regenerate (ADR-0062 D3).
+3. **GQA head / mask broadcast** — **ADR-0061** (`tl.broadcast`).
+   *Optional convenience*, not a GQA blocker (see correction above).
+   Useful for additive-mask construction across the `G·T_q` rows and for
+   general kernels; `np.matmul` already broadcasts in data mode, so it is
+   not needed for correctness. Lowest priority.
+
+**Explicitly REJECTED (efficient alternative chosen):**
+
+4. ~~A bespoke "flash-composite" command kind that internalises the whole
+   inner loop — DMA→MM→VEC→DMA→MM→VEC with carried `(m,ℓ,O)` register
+   state and a tail IPCQ push.~~ The two GEMMs **do** use the existing
+   `CompositeCmd` (§1/§3) — that gives scheduler-managed tiling, K/V DMA
+   streaming, and cross-tile pipelining. What is rejected is a **new**
+   command kind that also absorbs the softmax merge + cross-tile register
+   lifetime + an IPCQ-push epilogue: it is large and special-purpose, and
+   its only delta over the hybrid (full softmax offload) is not justified
+   at the current modelling fidelity (per-op CPU issue cost = 0; see §1).
+   Revisit if the cost model (§9) makes full offload measurably worthwhile.
+5. ~~Hardware `pop`-as-dependency.~~ Out of scope (§6).
+6. ~~RoPE / QKV projection / KV-cache write inside this kernel.~~ Upstream
+   `qkv_rope` (P1–P5). Folding RoPE in would force re-rotating past tiles
+   every decode step and break Ring Attention's post-RoPE pass-through.
+
+---
+
+## 9. Open tuning items (measured in kernbench, not blocking)
+
+1. **Composite tile-pipeline depth** — dominant for KV-load-bound; how far
+   ahead the kernel issues non-blocking composites before waiting, and the
+   scheduler's per-tile streaming depth.
+2. **PE↔mesh-neighbour mapping for the reduction tree** — ensure each
+   depth-`⌈log₂ N⌉` pair is a physical `N/S/E/W` neighbour; bad mapping
+   adds hops. Verify against the SFR install.
+3. **TILE size** — balance scratch residency (`S/P tiles + O_acc + G-way
+   GQA`) against DMA efficiency; interacts with ADR-0063 and the
+   scheduler's `TILE_M/K/N` (`pe_scheduler.py`).
+4. **Ring buffer ping-pong vs `recv_async` depth** in §5.5.
+5. **One-shot vs tiled crossover for decode** (§5.2) — the `S_pe`
+   threshold where tiling beats a single composite.
+6. **Per-op CPU issue cost (cost model)** — currently `dispatch_cycles=0`
+   (`pe_cpu.py`), so composite-vs-primitive issue overhead is invisible.
+   An op-type-differentiated issue cost (a `tl.composite` descriptor push
+   ≫ a primitive op) is what makes the hybrid's CPU-saturation win
+   **measurable** (§1). Specified in **ADR-0064**; tracked as separate
+   future work.
+
+---
+
+## 10. Coverage summary
+
+| Case | KV placement | Inner loop | Reduction | Masking |
+|---|---|---|---|---|
+| Decode, no SP | 1 PE, all KV | one-shot or tiled, GQA-batched | none | last tile only |
+| Decode, SP | round-robin `N` PEs | tiled, GQA-batched | §4 tree (`⌈log₂ N⌉`) | last tile only |
+| Prefill, no SP | resident | tiled per q-block | none | skip future / triangular boundary |
+| Prefill, SP (Ring) | ring-sharded | per ring step (recv_async overlap) | running state + §4 tree | causal step-skip + boundary |
+
+**I/O per case** (full contract §0.5):
+
+| Case | Inputs | Output | IPCQ partials |
+|---|---|---|---|
+| Decode, no SP | `Q[G,1,d]`, full `K/V[S,d]` on 1 PE | `O[G,1,d]` at `O_base` | none |
+| Decode, SP | `Q[G,1,d]`, per-PE `K/V[S_pe,d]` | `O[G,1,d]` (root) | `(m,ℓ,O_i)` per PE → tree |
+| Prefill, no SP | `Q[G,T_q,d]`, `K/V[≤end,d]` | `O[G,T_q,d]` at `O_base` | none |
+| Prefill, SP (Ring) | `Q[G,T_q,d]`, ring-delivered `K/V` blocks | `O[G,T_q,d]` (root) | running `(m,ℓ,O)` + tree |
+
+In all cases: KV-cache writes and RoPE happen **upstream**; output
+projection **downstream**; score `S` and probs `P` are never materialised.
+
+---
+
+## 11. Verification plan (Phase 1 test outline)
+
+SPEC/ADR coverage: R5 (PE↔PE IPCQ, PE↔HBM), R2 (latency by traversal),
+ADR-0023/0025 (IPCQ), ADR-0046 (`tl` contract), ADR-0054 (eval bench).
+
+The simulator's contract is **latency by traversal + determinism +
+structural correctness** (SPEC §0, §0.1), not bit-exact numerics — Phase 2
+data exists mainly to exercise the data path, `tl.trans` is
+reshape-not-transpose, and `bf16` is modelled as `f16`. Verification is
+therefore **structural/timing-first**, with numeric parity as a bounded
+secondary check.
+
+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
+   restructuring, *not* a new primitive.)
+   **Numeric parity (secondary):** for symmetric/identity inputs where
+   reshape-as-transpose is exact, kernel `O` matches a numpy
+   FlashAttention reference within fp tolerance. Full asymmetric parity is
+   gated on a real `tl.transpose` (out of scope; flagged).
+2. **GQA reuse:** with `h_q = G·h_kv`, the K/V `dma_read_count` is
+   independent of `G` (one load per tile, reused across the group), while
+   GEMM work scales with `G`. Asserts the lever actually fires.
+3. **Tree reduction step count:** decode-SP issues `⌈log₂ N⌉` reduction
+   rounds (not `N−1`); op_log `ipcq_send`/`recv` counts match the tree.
+4. **Causal skip:** prefill skips all `tile_all_future` tiles — GEMM
+   count equals the lower-triangular tile count, not the full grid.
+5. **Long context (ADR-0063):** a sweep at `S` that overflows 1 MiB
+   without scopes completes and matches the reference.
+6. **Load/compute overlap:** end-to-end latency of the tiled sweep is
+   below the serial `Σ(load+compute)` — from the composite scheduler
+   streaming K/V per tile (§1/§3) and lazy `tl.load` (ADR-0062) overlapping
+   the Q load. (Overlap is real modelled concurrency, not a subtraction.)
+7. **Composite GEMM offload (structural):** each tile's Q·Kᵀ and P·V emit a
+   `CompositeCmd` (non-blocking) to PE_SCHEDULER, not a blocking `tl.dot`;
+   op_log shows the composite tile plan and the kernel issues the next
+   tile's composites before waiting (cross-tile pipelining).
+8. **Determinism:** identical inputs → identical op_log + latency
+   (SPEC §0.1).
+
+---
+
+## 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
+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.
+
+1. **DDD-0060 is not yet synced.** The Detailed Design Document still
+   describes the old `tl.load_async` double-buffer path and primitive
+   `tl.dot` inner loop (its §4.3/§5/§10). It must be updated to the hybrid
+   (composite GEMMs, lazy load, K pre-transposed). *Left for review*
+   because the DDD is a derived how-to and a large rewrite; the ADR is now
+   the authoritative record. **Recommend:** sync DDD as a follow-up before
+   implementation starts.
+
+2. **K operand orientation for the composite GEMM.** Q·Kᵀ needs `b =
+   [d, TILE]`, but the KV cache stores K as `[S_pe, d]`. `tl.trans` is
+   metadata-only and `MemoryStore.read` reshapes, not transposes
+   (`memory_store.py:73`) — so a runtime transpose is wrong for non-trivial
+   data. **Recommend:** store K **pre-transposed** `[d, S_pe]` in the cache
+   (the pseudocode and §3 assume this), making `tl.ref(k_tile, (d, TILE))`
+   a contiguous slice. Verify the upstream `qkv_rope` write layout supports
+   this, or add a real `tl.transpose` (heavier; deferred).
+
+3. **Composite output buffer vs `tl.scratch_scope`.** Each Q·Kᵀ composite
+   writes `Sj` to an `out_addr`; the kernel then reads it for the softmax
+   MATH. That output buffer, and the in-flight composites' targets, must
+   live where the per-tile `scratch_scope` (ADR-0063) will **not** recycle
+   them before they are consumed — same discipline as in-flight lazy loads
+   (ADR-0062 D-Negative). **Recommend:** composite outputs for the *current*
+   tile live in the scoped arena (consumed same iteration); the persistent
+   `(m,ℓ,O)` stays outside. Verify no use-after-recycle when the next
+   tile's composites are issued early (cross-tile pipelining).
+
+4. **GQA `dma_read_count` lever under composite streaming.** The lever
+   (§11.2: K/V `dma_read_count` independent of `G`) assumes the composite
+   emits **one** K/V tile DMA reused across all `G·T_q` M-rows. The
+   scheduler's `generate_gemm_plan` tiles by `TILE_M/K/N`
+   (`pe_scheduler.py:35-37`, 32/64/32) — confirm the M-tiling over `G·T_q`
+   does **not** re-issue the shared K/V tile DMA per M-tile (i.e. operand
+   DMA is shared across M-tiles, or the lever weakens). **Recommend:**
+   assert it in the levers test; if violated, the GQA win is in compute
+   only, not DMA — still correct, but the headline changes.
+
+5. **Kernel TILE vs scheduler `TILE_M/K/N`.** The kernel reasons about a
+   logical KV `TILE`; the scheduler re-tiles internally at fixed
+   `TILE_M/K/N`. Two tiling layers interact (scratch residency, pipeline
+   depth). **Recommend:** treat the kernel TILE as the K/V streaming
+   granularity and let the scheduler sub-tile the GEMM; document the
+   relationship in the DDD and sweep both (§9 items 1, 3).
+
+6. **Cost model is a separate ADR.** The hybrid's CPU-saturation benefit is
+   invisible while `dispatch_cycles=0`. The per-op-type issue-cost model is
+   specified in **ADR-0064**; this ADR's §1/§9 depend on it for the
+   *measurable* (not just structural) win. **Recommend:** land ADR-0064's
+   model before claiming hybrid latency wins in the eval.
+
+7. **Ring path (§5.5) GEMMs.** §5.5 still describes the ring fold with
+   primitive ops + `recv_async`. For consistency the ring's per-step Q·Kᵀ /
+   P·V should also be composites; the IPCQ `recv_async` overlap is
+   orthogonal and stays. **Recommend:** apply the same hybrid shape to the
+   ring step during implementation; low risk, mirrors §3.

docs/adr-proposed/ADR-0061-prog-tl-broadcast.md

@@ -0,0 +1,203 @@
+# ADR-0061: `tl.broadcast` / `tl.repeat` — data-faithful GQA head reuse
+
+## Status
+
+Proposed
+
+> Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention).
+>
+> **⚠ Re-evaluated during design iteration — this is now OPTIONAL, not a
+> GQA blocker.** The original framing (below) held that GQA `h_q > h_kv`
+> is blocked by the MemoryStore byte-conservation check on a broadcast
+> view. That is true *only of the baseline's head-packing hack*
+> (`_view(K, (h_q·d, S_kv))`). The correct GQA design (ADR-0060 §5.2)
+> processes **one KV head at a time** and folds the `G` group rows into
+> the matmul **M** dimension — using only byte-conserving reshapes — so it
+> needs **no broadcast op at all**, and `np.matmul` already broadcasts a
+> shared operand in Phase-2 data mode. This ADR therefore drops to a
+> *convenience* primitive: it cleans up additive-mask construction across
+> the `G·T_q` rows and serves general kernels. **Lowest priority of the
+> three supporting ADRs.** The rest of this document is retained as the
+> design for that convenience op.
+
+## Context
+
+### The GQA reuse lever
+
+Grouped-Query Attention amortises K/V loads: `G = H_q / H_kv` query heads
+share one KV head. For Llama3-70B `G = 8`. In a kernel this means a KV
+tile of shape `[S_kv, d]` (for one KV head) must be matmul'd against the
+`G` query rows of its group. Loading K/V **once** and reusing it across
+the group is the decode-time efficiency win (ADR-0060 §5.2).
+
+### Why it does not work today
+
+The existing mesh kernels reshape tensors with a metadata-only helper
+`_view` (`src/kernbench/benches/_attention_mesh_kv.py:25-36`,
+`_attention_mesh_mlo.py:25-36`):
+
+```python
+def _view(handle, new_shape):
+    return TensorHandle(..., shape=new_shape, nbytes=handle.nbytes, ...)
+    #                          ^^^ new shape          ^^^ OLD nbytes (unchanged)
+```
+
+This rewrites `shape` but keeps the original `nbytes`. It is purely
+symbolic. It survives **Phase 1 (timing)** because Phase 1 never reads
+tensor bytes. It fails **Phase 2 (`enable_data=True`)** because the
+data executor materialises every op through `MemoryStore.read`, which
+enforces byte conservation
+(`src/kernbench/sim_engine/memory_store.py:67-73`):
+
+```python
+if shape is not None and arr.shape != shape:
+    if arr.nbytes != np.prod(shape) * arr.dtype.itemsize:
+        raise ValueError(f"Shape mismatch: stored {arr.shape} ({arr.nbytes}B) "
+                         f"vs requested {shape} ({...}B)")
+    arr = arr.reshape(shape)
+```
+
+A reshape conserves bytes; a **broadcast does not** (`[S_kv, 1, d]` →
+`[S_kv, G, d]` multiplies the element count by `G`). So the moment a
+kernel tries to view a `h_kv`-headed K as if it had `h_q` heads, the
+stored array has `1/G` of the requested bytes and `read` raises.
+
+The current test suite documents this as a deliberate limitation
+(`tests/attention/test_milestone_gqa_llama70b.py:137-142`):
+
+> "v1 uses `h_q == h_kv == 1` to avoid … GQA broadcast view (which is
+> symbolic and does not survive MemoryStore's nbytes check under
+> simulator data execution). Real GQA (`h_q > h_kv`) is deferred."
+
+There is **no** broadcast / repeat / expand primitive in the `tl` API
+today (`src/kernbench/triton_emu/tl_context.py` exposes `trans`,
+`zeros`, `full`, `arange` as metadata-only ops; none replicate data).
+
+## Decision
+
+Add a **data-faithful broadcast op** to the `tl` API and the data
+executor. Unlike `_view`, it updates `nbytes` and, in Phase 2, actually
+replicates the underlying ndarray. The latency it charges models the
+on-chip register/SRAM fan-out, not an HBM re-load (the KV tile is read
+from HBM once; broadcast is an on-chip replication).
+
+### D1. `tl` surface
+
+```python
+def broadcast(self, x: TensorHandle, new_shape: tuple[int, ...]) -> TensorHandle:
+    """Replicate x along size-1 axes to new_shape (numpy broadcast rules).
+
+    new_shape must be broadcast-compatible with x.shape: every axis is
+    either equal to x's, or x's is 1 (the axis being expanded). nbytes of
+    the result reflect new_shape (unlike the metadata-only _view).
+    """
+```
+
+`tl.repeat(x, axis, n)` is offered as a thin convenience wrapper over
+`broadcast` for the common "insert a group axis and expand it to G"
+pattern, but `broadcast` is the primitive.
+
+### D2. Command + op_log
+
+Emit a new `BroadcastCmd` (sibling of `MathCmd`):
+
+```python
+@dataclass(frozen=True)
+class BroadcastCmd:
+    src: TensorHandle
+    out: TensorHandle          # out.nbytes reflects new_shape
+    new_shape: tuple[int, ...]
+    data_op: bool = True
+```
+
+op_log entry: `op_kind="math"`, `op_name="broadcast"` (it runs on the
+MATH/vector engine, like the other element-wise ops). The output handle
+is allocated from PE scratch via the existing `_make_compute_out`
+(`tl_context.py:148-156`), so `out.nbytes` is computed from `new_shape`
+and the byte-conservation check passes in Phase 2.
+
+### D3. Phase 2 data execution
+
+Add `_execute_broadcast` to the data executor
+(`src/kernbench/sim_engine/data_executor.py`):
+
+```python
+def _execute_broadcast(self, op):
+    p = op.params
+    src = self._resolve_read(p["src_space"], p["src_addr"],
+                             p["src_shape"], p["dtype"], op.t_start)
+    result = np.broadcast_to(src, p["new_shape"]).copy()   # .copy() → writable, owns bytes
+    self.store.write(p["dst_space"], p["dst_addr"], result)
+```
+
+`.copy()` is required: `np.broadcast_to` returns a zero-stride view, and
+MemoryStore stores references — a later consumer's nbytes check must see
+a real `prod(new_shape)`-element array.
+
+### D4. Latency model
+
+Broadcast latency = vector-engine cost of writing `prod(new_shape)`
+elements (reuse `pe_math._compute_ns(num_elements)`,
+`src/kernbench/components/builtin/pe_math.py:46-51`). This models the
+on-chip replication into the register file / SRAM that real GQA hardware
+performs. It is **not** an HBM transaction — no `dma_read` is logged,
+consistent with "K/V tile loaded once" (ADR-0060 §5.2).
+
+## Alternatives
+
+### A1. Indexed per-head dot (no broadcast op)
+
+Loop the `G` query heads, each doing its own `tl.dot(Q_g, Kᵀ)` against
+the *same* (re-`ref`'d, not re-loaded) KV tile:
+
+```python
+for g in range(G):
+    s_g = tl.dot(Q[g], K_T)     # G separate GEMMs, KV reused via tl.ref
+```
+
+- Pro: no new primitive; works in Phase 2 today.
+- Con: `G×` more GEMM commands (op_log inflation), no batched-GEMM
+  modelling, and the per-call overhead is charged `G` times. It also
+  obscures the "single batched matmul over the group" the hardware does.
+  Rejected as the primary path; acceptable as a fallback if the batched
+  2D dot's shape handling proves awkward.
+
+### A2. Keep `_view`, relax the nbytes check
+
+Make `MemoryStore.read` tolerate broadcast (stride-0) shapes. Rejected:
+it silently breaks byte conservation for *all* readers and would mask
+genuine shape bugs elsewhere. The check is a correctness invariant
+(ADR-0052), not an obstacle to route around.
+
+### A3. Materialise the broadcast at deploy time
+
+Store K/V already expanded to `h_q` heads in HBM. Rejected: defeats the
+entire point of GQA (it `G×`'s KV-cache HBM footprint and bandwidth).
+
+## Consequences
+
+### Positive
+- Unblocks real GQA (`h_q > h_kv`) in Phase 2 — the headline blocker for
+  ADR-0060 and the deferred milestone work.
+- A general, reusable primitive (any kernel needing numpy-style
+  broadcast benefits, not just attention).
+- Byte conservation stays intact; no weakening of MemoryStore.
+
+### Negative
+- One new command type + one data-executor handler + one `tl` method.
+- Phase 2 `.copy()` allocates the expanded array in the store; for very
+  large groups this is real memory. Bounded by tile size in the tiled
+  design (ADR-0060), so acceptable.
+
+## Test Requirements
+
+1. **Phase 1 timing**: `tl.broadcast([S,1,d] → [S,G,d])` logs one
+   `math/broadcast` op with `num_elements == S*G*d`; no `dma_read`.
+2. **Phase 2 data**: broadcasting a known array yields the numpy
+   `broadcast_to(...).copy()` result; a downstream `tl.dot` consuming it
+   passes the MemoryStore nbytes check (the exact case that fails today).
+3. **GQA end-to-end**: re-run a reduced `milestone-gqa-llama70b` panel
+   with `h_q=8, h_kv=1` under `enable_data=True`; it must complete (today
+   it raises `ValueError: Shape mismatch`).
+4. **Broadcast-incompatible shape** (e.g. axis size 3 → 8) raises a
+   clear error at emit time, not deep in Phase 2.

docs/adr-proposed/ADR-0062-prog-tl-async-load-ko.md

@@ -0,0 +1,148 @@
+# ADR-0062: Lazy `tl.load` — 첫 사용 시점 auto-wait를 갖는 non-blocking HBM load
+
+## Status
+
+Proposed
+
+> **ADR-0060**(AHBM GQA Fused Attention)의 보조 ADR. Decode와 long-context
+> 어텐션은 **KV-load-bound**이며, load/compute 오버랩이 지배적 레버다. 본
+> ADR은 별도의 `load_async` op을 추가하는 대신 `tl.load` 자체를 **lazy**
+> (non-blocking, wait를 첫 사용 시점에 자동 삽입)로 만든다. 오늘날 유일한
+> async primitive는 IPCQ 통신용이지 HBM load용이 아니다.
+
+## Context
+
+### 오버랩이 요구하는 것
+
+FlashAttention은 operand를 스트리밍한다: GEMM/MATH 엔진이 현재 타일을 처리하는
+동안 DMA 엔진은 이미 다음 operand를 당기고 있어야 한다. 그 오버랩이 있으면
+대역폭-bound 커널은 타일당 `compute + dma` 대신 대략 `max(compute, dma)`로
+돈다.
+
+ADR-0060의 하이브리드 설계에서 두 GEMM은 `tl.composite(op="gemm")`로 발행되고,
+그 K/V operand는 `tl.ref`(HBM 상주, PE_SCHEDULER가 타일당 스트리밍)이므로 **타일당
+K/V prefetch는 composite scheduler가 처리**한다. lazy `tl.load`는 나머지 명시적
+load(Q group, 그리고 비-composite 커널)를 담당하여 그것들도 greenlet을 멈추는
+대신 뒤따르는 compute와 오버랩되게 한다.
+
+### 현재 존재하는 것
+
+- `tl.load(ptr, shape, dtype)`은 **blocking**이다: `DmaReadCmd`를 emit하고
+  greenlet 커널은 PE_DMA가 완료를 알릴 때까지 suspend된다
+  (`tl_context.py:177-203`; greenlet 구동 `kernel_runner.py:146-153`). 한
+  커널에서 두 `tl.load`가 in-flight일 수 없다.
+- async 패턴이 **존재**하긴 하나 IPCQ 전용이다:
+  `tl.recv_async(dir, ...) -> RecvFuture` + 지연 wait
+  (`kernel_runner.py:248-285`). 이는 메커니즘을 입증한다 — future를 반환하고
+  나중에 wait 체크(`if not future.event.triggered: yield future.event`)로
+  resolve되는 non-blocking 커맨드가 greenlet 모델에서 동작한다.
+- DMA 엔진은 read 채널을 write 채널과 분리된 SimPy resource(capacity 1,
+  `pe_dma.py:45`)로 모델한다. 따라서 in-flight read는 표현 가능하며, 단일 read
+  채널에서 직렬화되면서 PE_GEMM/PE_MATH의 compute와 오버랩된다. 오직
+  *커널-대면 API*만이 현재 load-vs-compute를 직렬화한다.
+
+`tl.load`는 오늘날 뒤따르는 compute와 오버랩될 수 없다.
+
+## Decision
+
+**`tl.load`를 lazy로 만든다: `DmaReadCmd`를 발행하고 핸들을 즉시 반환
+(non-blocking); 런타임은 load된 데이터가 실제로 처음 소비되는 지점에 wait를
+자동 삽입한다.** 커널-대면 API는 불변이다 — 작성자는 계속 `tl.load`를 쓴다 —
+따라서 이는 새 op이 아니라 *시맨틱* 변경이다. 기존 `recv_async`/wait 기계장치를
+(1) HBM-load 경로로, (2) 명시적 `tl.wait` 호출에서 암묵적·의존성 기반 first-use
+wait로 일반화한다.
+
+### D1. `tl` 표면 — 불변
+
+`tl.load(ptr, shape, dtype)`은 시그니처와 `TensorHandle` 반환을 유지한다.
+바뀌는 것은 *언제* blocking하느냐다: 발행 시점에는 결코 아니고, 결과가 소비 op에
+의해 처음 읽힐 때만 암묵적으로.
+
+### D2. 메커니즘 — non-blocking 발행 + 사용 시점 auto-wait
+
+- `tl.load`는 완료 event를 yield하지 않고 `DmaReadCmd`를 PE_DMA에 post하며,
+  반환 핸들에 pending event를 기록한다(`recv_async` 패턴을 load에 적용).
+- 소비 op(`tl.dot`, MATH op, `tl.store`, `tl.composite` operand, …)이
+  dispatch될 때, 런타임은 각 입력 핸들의 pending load event를 확인하고 아직
+  triggered가 아니면 먼저 yield한다 — 즉 wait가 가장 늦은 올바른 지점(첫 사용)에
+  자동 삽입된다.
+- op_log 엔트리는 불변(`memory/dma_read`): asynchrony는 스케줄링 속성이지 새 op
+  종류가 아니므로 `dma_read_count` 및 기존 op_log 소비자가 계속 동작한다.
+
+### D3. 범위 — 전역
+
+`tl.load`는 opt-in 플래그 뒤가 아니라 **모든 곳에서** lazy다. 이것이 충실한
+모델이다: 매 load마다 blocking하는 것은 *naive* 커널의 속성이고, 효율 커널(그리고
+실제 컴파일러)은 load를 hoist하고 사용 시점에만 wait한다. 결과: `tl.load`와 첫
+사용 사이에 독립 작업이 있는 기존 커널은 **더 낮은(빠른) latency**를 본다 —
+동작의 회귀가 아니라 모델의 정확도 개선. 바뀌는 골든 latency는 **재생성**해야
+하며; load 직후 곧바로 사용하는 커널은 변화가 없다(auto-wait이 즉시 발동하여
+blocking과 동일).
+
+### D4. Latency / 오버랩 시맨틱
+
+- `tl.load`는 **발행**(descriptor push)만 청구하고 커널은 진행한다. (per-op
+  발행 비용은 `dispatch_cycles`, 현재 0; op 종류별 차등 발행 비용은 별도로
+  추적 — ADR-0060 §1, §9, ADR-0064.)
+- DMA 전송은 read 채널(capacity 1)을 모델 지속시간만큼 점유하며 커널이 다음에
+  내는 compute와 병렬; 다중 in-flight load는 채널에서 직렬화되나 compute와
+  오버랩.
+- 자동 삽입된 wait는 전송이 끝나지 않았을 때만 blocking.
+- 결정성 보존: wait 지점은 프로그램 순서(첫 사용)로 고정되고, 완료는 modeled DMA
+  채널의 scheduled event다(SPEC §0.1, R8). latency 빼기 없음 — 오버랩은 실제
+  모델 동시성.
+
+> **모델링 가정.** auto-wait-at-first-use는 *잘 스케줄된* 커널(컴파일러가
+> wait를 가장 늦은 올바른 지점에 둠)을 모델한다. 실제 컴파일러는 이를 근사하며,
+> 일부 load는 hoist 불가(register pressure, aliasing). 성능 시뮬레이터(SPEC
+> §0)에서 잘 스케줄된 경우를 모델하는 것이 의도된 동작이다.
+
+## Alternatives
+
+### A1. 별도 `tl.load_async` op (이전 제안)
+
+커널이 손으로 호출하는 명시적 `load_async`/`wait` 쌍 추가(더블버퍼 댄스). 기각:
+커널-대면 API를 키우고 버퍼 수명 부기를 모든 커널 작성자에게 떠넘긴다. lazy
+`tl.load`는 **API 표면 변경 없이** 동일 오버랩과 컴파일러식 auto-wait를 준다.
+
+### A2. 커널별 opt-in lazy load
+
+`tl.load`를 기본 blocking으로 두고 새 커널만 opt-in. 기각: `tl.load` 시맨틱을 두
+변종으로 쪼개고 기존 bench로부터 모델 개선을 숨긴다; 전역 lazy가 더 깔끔하고
+골든 재생성 비용은 1회성(D3).
+
+### A3. composite 스트리밍에만 의존 (lazy load 없이)
+
+ADR-0060의 composite는 `tl.ref` K/V operand를 스트리밍하므로 GEMM operand DMA는
+이미 오버랩된다. 그러나 명시적 load(Q group, 비-composite 커널)는 lazy `tl.load`
+없이 여전히 stall. 단독으로는 불충분.
+
+## Consequences
+
+### Positive
+- 커널-대면 API 변경 **제로**로 load/compute 오버랩; 작성자는 계속 `tl.load`.
+- `recv_async`/wait와 대칭 — 개념적 표면적 작음.
+- 일반적: 어떤 대역폭-bound 커널도 자동 prefetch.
+
+### Negative
+- load와 사용 사이에 독립 작업이 있는 커널의 기존 골든 latency가 이동(빨라짐)
+  → 1회성 재생성(D3).
+- auto-wait는 런타임이 핸들별 pending event를 추적하고 소비-op dispatch에서
+  확인하도록 요구(데이터 의존성 추적).
+- scratch 수명(ADR-0063)과 상호작용: in-flight load의 대상 버퍼는 auto-wait이
+  발동하기 전에 recycle되면 안 된다. recycling scope는 살아있는(미-wait된) load
+  버퍼를 제외해야 한다.
+
+## Test Requirements
+
+1. **오버랩이 실재**: `tl.load` 후 비슷한 지속시간의 독립 GEMM을 내는 커널이
+   `load+gemm`이 아니라 ≈`max(load, gemm)`으로 완료(end-to-end latency가 직렬
+   합보다 엄격히 작음).
+2. **auto-wait 정확성**: load된 `TensorHandle`이 처음 소비될 때, 같은 주소에
+   대해 오늘날 blocking `tl.load`와 동일한 바이트를 보유(Phase 2).
+3. **둘 in-flight**: 서로 다른 주소로의 두 `tl.load`가 나중에 소비되어 둘 다
+   올바른 독립 텐서로 resolve; 그 DMA는 read 채널에서 직렬화.
+4. **op_log 호환성**: 각 `tl.load`은 여전히 정확히 하나의 `memory/dma_read`를
+   로그; `dma_read_count`는 blocking 버전 대비 불변.
+5. **무-오버랩 무-변화**: load 직후 곧바로 사용하는 커널은 blocking 모델과 동일한
+   latency(auto-wait 즉시 발동).

docs/adr-proposed/ADR-0062-prog-tl-async-load.md

@@ -0,0 +1,169 @@
+# ADR-0062: Lazy `tl.load` — non-blocking HBM load with auto-wait on first use
+
+## Status
+
+Proposed
+
+> Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention). Decode and
+> long-context attention are **KV-load-bound**; load/compute overlap is a
+> dominant lever. This ADR makes `tl.load` itself **lazy** (non-blocking,
+> with the wait automatically inserted at first use) rather than adding a
+> separate `load_async` op. Today the only async primitive is for IPCQ
+> comms, not for HBM loads.
+
+## Context
+
+### What overlap requires
+
+FlashAttention streams operands: while the GEMM/MATH engine works on the
+current tile, the DMA engine should already be pulling the next operand.
+With that overlap a bandwidth-bound kernel runs at roughly
+`max(compute, dma)` per tile instead of `compute + dma`.
+
+In ADR-0060's hybrid design the two GEMMs are issued as
+`tl.composite(op="gemm")` whose K/V operands are `tl.ref` (HBM-resident,
+streamed per tile by PE_SCHEDULER), so the **per-tile K/V prefetch is
+handled by the composite scheduler**. Lazy `tl.load` covers the remaining
+explicit loads (the Q group, and any non-composite kernel) so those also
+overlap the compute that follows instead of stalling the greenlet.
+
+### What exists
+
+- `tl.load(ptr, shape, dtype)` is **blocking**: it emits `DmaReadCmd` and
+  the greenlet kernel suspends until PE_DMA signals completion
+  (`tl_context.py:177-203`; greenlet drive `kernel_runner.py:146-153`).
+  No two `tl.load`s can be in flight from one kernel.
+- An async pattern **does** exist, but only for IPCQ:
+  `tl.recv_async(dir, ...) -> RecvFuture` + a deferred wait
+  (`kernel_runner.py:248-285`). It proves the machinery — a non-blocking
+  command that returns a future, resolved later by a wait check
+  (`if not future.event.triggered: yield future.event`) — works in the
+  greenlet model.
+- The DMA engine models a read channel as a SimPy resource (capacity 1,
+  `pe_dma.py:45`) separate from the write channel, so in-flight reads are
+  representable; they serialise on the single read channel while
+  overlapping compute on PE_GEMM/PE_MATH. Only the *kernel-facing API*
+  currently serialises load-vs-compute.
+
+`tl.load` cannot today overlap with the compute that follows it.
+
+## Decision
+
+**Make `tl.load` lazy: it issues the `DmaReadCmd` and returns a handle
+immediately (non-blocking); the runtime auto-inserts the wait at the
+first point the loaded data is actually consumed.** The kernel-facing API
+is unchanged — authors keep writing `tl.load` — so this is a *semantics*
+change, not a new op. It generalises the existing `recv_async`/wait
+machinery (1) to the HBM-load path and (2) from an explicit `tl.wait`
+call to an implicit, dependency-driven wait at first use.
+
+### D1. `tl` surface — unchanged
+
+`tl.load(ptr, shape, dtype)` keeps its signature and `TensorHandle`
+return. What changes is *when* it blocks: never at issue, only implicitly
+when its result is first read by a consuming op.
+
+### D2. Mechanism — non-blocking issue + auto-wait on use
+
+- `tl.load` posts the `DmaReadCmd` to PE_DMA without yielding its
+  completion event, and records the pending event on the returned handle
+  (the `recv_async` pattern, applied to loads).
+- When a consuming op (`tl.dot`, a MATH op, `tl.store`, a `tl.composite`
+  operand, …) is dispatched, the runtime checks each input handle for a
+  pending load event and yields it first if not yet triggered — i.e. the
+  wait is inserted automatically at the latest correct point (first use).
+- The op_log entry is unchanged (`memory/dma_read`): asynchrony is a
+  scheduling property, not a new op kind, so `dma_read_count` and existing
+  op_log consumers keep working.
+
+### D3. Scope — global
+
+`tl.load` is lazy **everywhere**, not behind an opt-in flag. This is the
+faithful model: blocking on every load is a property of a *naive* kernel;
+an efficient kernel (and a real compiler) hoists the load and waits only
+at use. Consequence: existing kernels that have independent work between a
+`tl.load` and its first use see **lower (faster) latency** — a
+correctness improvement of the model, not a behaviour regression. Golden
+latencies that change must be **regenerated**; kernels that load then
+immediately use see no change (the auto-wait fires at once, identical to
+blocking).
+
+### D4. Latency / overlap semantics
+
+- `tl.load` charges the **issue** (descriptor push) only; the kernel
+  proceeds. (Per-op issue cost is `dispatch_cycles`, currently 0; an
+  op-type-differentiated issue cost is tracked separately — ADR-0060 §1,
+  §9.)
+- The DMA transfer occupies the read channel (capacity 1) for its
+  modelled duration in parallel with whatever compute the kernel issues
+  next; multiple in-flight loads serialise on the channel but overlap
+  compute.
+- The auto-inserted wait blocks only if the transfer has not finished.
+- Determinism is preserved: the wait point is fixed by program order
+  (first use), and completion is a scheduled event on the modelled DMA
+  channel (SPEC §0.1, R8). No latency subtraction — the overlap is real
+  modelled concurrency.
+
+> **Modelling assumption.** Auto-wait-at-first-use models a *well-
+> scheduled* kernel (the compiler places the wait at the latest correct
+> point). Real compilers approximate this; some loads cannot be hoisted
+> (register pressure, aliasing). For a performance simulator (SPEC §0)
+> modelling the well-scheduled case is the intended behaviour.
+
+## Alternatives
+
+### A1. Separate `tl.load_async` op (earlier proposal)
+
+Add an explicit `load_async`/`wait` pair the kernel calls by hand (the
+double-buffer dance). Rejected: it enlarges the kernel-facing API and
+pushes buffer-lifetime bookkeeping onto every kernel author, when lazy
+`tl.load` gives the same overlap with **no** API-surface change and a
+compiler-style auto-wait.
+
+### A2. Per-kernel opt-in lazy load
+
+Keep `tl.load` blocking by default; make new kernels opt in. Rejected:
+splits `tl.load` semantics into two variants and hides the model
+improvement from existing benches; global lazy is cleaner and the
+golden-regeneration cost is one-time (D3).
+
+### A3. Rely on composite streaming only (no lazy load)
+
+ADR-0060's composites stream their `tl.ref` K/V operands, so the GEMM
+operand DMA already overlaps. But explicit loads (the Q group, and any
+non-composite kernel) still stall without lazy `tl.load`. Insufficient on
+its own.
+
+## Consequences
+
+### Positive
+- Load/compute overlap with **zero** kernel-facing API change; authors
+  keep writing `tl.load`.
+- Symmetric with `recv_async`/wait — low conceptual surface area.
+- General: any bandwidth-bound kernel prefetches automatically.
+
+### Negative
+- Existing golden latencies shift (faster) for kernels with independent
+  work between load and use → one-time regeneration (D3).
+- Auto-wait requires the runtime to track per-handle pending events and
+  check them at consuming-op dispatch (data-dependency tracking).
+- Interacts with scratch lifetime (ADR-0063): an in-flight load's target
+  buffer must not be recycled before its auto-wait fires. The recycling
+  scope must exclude live (un-waited) load buffers.
+
+## Test Requirements
+
+1. **Overlap is real**: a kernel that issues `tl.load` then an
+   independent GEMM of comparable duration completes in ≈`max(load,
+   gemm)`, not `load+gemm` (end-to-end latency strictly below the serial
+   sum).
+2. **Auto-wait correctness**: the loaded `TensorHandle`, when first
+   consumed, carries the same bytes as today's blocking `tl.load` for the
+   same address (Phase 2).
+3. **Two in flight**: two `tl.load`s to distinct addresses, consumed
+   later, both resolve to correct, independent tensors; their DMAs
+   serialise on the read channel.
+4. **op_log compatibility**: each `tl.load` still logs exactly one
+   `memory/dma_read`; `dma_read_count` unchanged vs the blocking version.
+5. **No-overlap no-change**: a kernel that loads then immediately uses has
+   identical latency to the blocking model (auto-wait fires at once).

docs/adr-proposed/ADR-0063-prog-tl-scratch-scope.md

@@ -0,0 +1,162 @@
+# ADR-0063: `tl.scratch_scope` — per-tile scratch recycling for long context
+
+## Status
+
+Proposed
+
+> Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention). Long-context
+> attention sweeps many K/V tiles; each tile's intermediates
+> (`scores`, `P`, `exp`, partial `O`, …) currently allocate fresh
+> scratch that is never freed within a kernel invocation. The 1 MiB
+> per-PE scratch budget is exhausted long before a realistic context
+> length.
+
+## Context
+
+### The bump allocator
+
+`TLContext` allocates every math/compute output handle from a per-PE
+scratch pool with a **linear bump cursor**
+(`src/kernbench/triton_emu/tl_context.py`; `_scratch_alloc`):
+
+```python
+def _scratch_alloc(self, nbytes):
+    aligned = (nbytes + 15) & ~15
+    addr = self._scratch_base + self._scratch_cursor
+    self._scratch_cursor += aligned
+    if self._scratch_cursor > self._scratch_size:      # default 1 << 20 = 1 MiB
+        raise RuntimeError("TLContext scratch overflow: ...")
+    return addr
+```
+
+The docstring states the cursor **"resets on every kernel invocation"** —
+i.e. only at kernel entry, never *within* the kernel body. Every
+`tl.dot`, `tl.softmax`, `tl.exp`, `tl.sum`, `a - b`, `a * b`, … grabs a
+fresh slice and nothing is reclaimed until the kernel returns.
+
+### Why this bites the GQA kernel
+
+The current milestone bench keeps `S_q = S_kv_per_rank = 16` **explicitly
+because** of this limit
+(`tests/attention/test_milestone_gqa_llama70b.py:123-148`):
+
+> "S_q_prefill and S_kv_per_rank are deliberately small (16 each) so the
+> simulator's 1 MB per-PE TCM kernel scratch is not exhausted by the
+> bump-allocated handle outputs of softmax/exp/dot/sum chains over
+> n_ranks ring steps."
+
+A FlashAttention sweep over `n_tiles` tiles allocates O(`n_tiles` ×
+per-tile-intermediates). For any realistic context this overflows. The
+*math* of flash attention needs only **O(1)** live scratch — the running
+`(m, l, O)` plus the current tile's working set — because each tile's
+temporaries are dead once that tile is folded into the running state. The
+allocator just doesn't know they're dead.
+
+## Decision
+
+Add a **scratch scope** that lets a kernel mark a region of allocations
+as reclaimable, so per-tile temporaries are recycled while live running
+state (and in-flight prefetch buffers) are preserved.
+
+### D1. `tl` surface — context manager
+
+```python
+with tl.scratch_scope():
+    s   = tl.dot(q, k_t)        # all handles allocated inside the `with`
+    p   = tl.softmax(s)         # share a region that is rewound on exit
+    o_j = tl.dot(p, v)
+    # ... fold o_j into running (m,l,O) which live OUTSIDE the scope ...
+# on __exit__: cursor rewinds to its value at __enter__
+```
+
+Semantics:
+- `__enter__` records the current `_scratch_cursor` as a save-point.
+- `__exit__` restores the cursor to the save-point, freeing everything
+  allocated inside.
+- Handles allocated **outside** the scope (running `m,l,O`, prefetch
+  buffers held by `LoadFuture` from ADR-0062) keep their addresses —
+  they were allocated before the save-point or in an enclosing scope.
+
+### D2. Safety contract
+
+A handle allocated inside a scope **must not** be read after the scope
+exits — its bytes may be overwritten by the next scope's allocations.
+The flash loop respects this naturally: the only values that survive a
+tile iteration are the running accumulators, which are allocated outside
+the per-tile scope and updated by ops *inside* it writing to outside
+addresses (the merge writes new running state — see D3).
+
+### D3. Interaction with running accumulators
+
+The online-softmax merge reads the old running `(m, l, O)` and the
+current tile's `(m_j, l_j, O_j)`, producing new running values. To keep
+the new running values outside the recycled region, the merge writes them
+to **stable scratch** allocated once before the loop (a small fixed
+"running-state" arena, distinct from the per-tile scope). Concretely the
+kernel keeps two arenas:
+
+- **persistent arena** (allocated once): `m, l, O` (and their
+  double-buffer if needed for the merge).
+- **scoped arena** (rewound each tile): `scores, P, exp, O_j, scale_*`.
+
+This mirrors how real flash-attention SRAM budgeting works: a small
+persistent accumulator region + a recycled tile working set.
+
+### D4. Nesting
+
+Scopes nest (stack of save-points). Inner scope exit rewinds to the inner
+save-point; outer exit rewinds further. This supports an outer
+"per-query-block" scope around an inner "per-KV-tile" scope for prefill.
+
+## Alternatives
+
+### A1. Round-trip temporaries through HBM
+
+Store intermediates to HBM and reload to "free" TCM scratch. Rejected:
+turns a TCM-resident streaming kernel into an HBM-bandwidth-bound one —
+the opposite of the goal, and it pollutes op_log with spurious DMA.
+
+### A2. Tiled `tl.composite` (scheduler-managed scratch)
+
+`tl.composite` recycles per-tile scratch inside PE_SCHEDULER
+automatically. As in ADR-0062 A1, this is attractive but blocked on a
+flash-capable composite kind (two GEMMs + carried `(m,l,O)`), a much
+larger change. `scratch_scope` gives the same memory behaviour for the
+greenlet kernel with a tiny, general primitive. The two can coexist.
+
+### A3. Grow the scratch budget
+
+Bump `pe_tcm.kernel_scratch_mb`. Rejected as a fix: it only pushes the
+context-length ceiling out linearly while still leaking O(`n_tiles`)
+scratch, and it misrepresents the hardware (real TCM is small; the point
+of flash attention is O(1) working set). Useful only as a coarse knob,
+not a substitute for recycling.
+
+## Consequences
+
+### Positive
+- Removes the artificial `S = 16` validation-scale ceiling; enables
+  realistic context lengths in both timing and data modes.
+- Faithfully models flash attention's O(1) working-set property.
+- Small, general primitive (any tiled kernel benefits).
+
+### Negative
+- A use-after-scope bug silently reads stale bytes. Mitigated by the D2
+  contract, by keeping the scope discipline inside a shared attention
+  helper, and (optionally) by a debug build that poisons rewound regions.
+- Must coordinate with ADR-0062: prefetch buffers are live across tile
+  iterations, so they belong to the persistent arena, not the scoped one.
+
+## Test Requirements
+
+1. **Recycling**: a loop of `N` tiles inside `tl.scratch_scope()` keeps
+   peak `_scratch_cursor` bounded by one tile's footprint, independent of
+   `N` (today it grows linearly and overflows).
+2. **Correctness**: a flash sweep with scopes produces the same `O` as
+   the same sweep without scopes at a small `N` that fits without
+   recycling (Phase 2).
+3. **Long context**: a sweep at an `N` that would overflow 1 MiB without
+   scopes completes (the exact failure the `S=16` cap avoids today).
+4. **Persistent-vs-scoped isolation**: running `(m,l,O)` allocated
+   outside the scope retains correct values across `__exit__`.
+5. **Nesting**: nested scopes rewind to the correct save-points.

docs/adr-proposed/ADR-0064-perf-cpu-issue-cost-model-ko.md

@@ -0,0 +1,166 @@
+# ADR-0064: op 종류별 CPU 발행 비용 모델 (command construct + dispatch)
+
+## Status
+
+Proposed
+
+> **ADR-0060**(AHBM GQA Fused Attention)의 보조 ADR. 거기서의 하이브리드
+> 결정(GEMM은 `tl.composite`, softmax 머지는 커널)이 이기는 이유는
+> **타일링을 PE_SCHEDULER에 offload하여 CPU가 coarse descriptor를 내리고
+> 앞서 나가 엔진을 saturate**하기 때문이다. 그 이점이 현재 시뮬레이터에서는
+> per-op CPU 발행 비용이 0이라 **보이지 않는다**. 본 ADR은 발행 비용을
+> 실재화하고 **op 종류별로 차등**화하여, composite-vs-primitive 트레이드오프
+> (그리고 CPU saturation)가 측정 가능하도록 한다.
+
+## Context
+
+### 현재 상태
+
+- 모든 `tl.*` op은 커맨드를 emit하기 전에 `_emit_dispatch_overhead()`를
+  호출하며(`tl_context.py:123-125, 190, 227, 235, 612, …`), 이는
+  `dispatch_cycles > 0`일 때**만** `PeCpuOverheadCmd(cycles=dispatch_cycles)`를
+  emit한다. 즉 발행 비용은 `tl.load`, `tl.dot`, MATH op, `tl.composite`에
+  동일하게 적용되는 **단일 균일 노브**다.
+- 그 노브는 두 실행 경로 모두에서 **0**으로 하드코딩되어 있다
+  (`pe_cpu.py:101` greenlet runner, `:195` legacy replay). ⇒ 커맨드 발행
+  — *descriptor를 construct하고 scheduler 큐에 push* — 이 현재 PE_CPU에서
+  **0 ns**다.
+- `PeCpuOverheadCmd`는 PE_CPU(`kernel_runner.py:131-132`)와
+  PE_SCHEDULER(`pe_scheduler.py:97-100`)에서 `yield env.timeout(cmd.cycles)`로
+  소비된다. 수동 `tl.cycles(n)`도 존재한다(`tl_context.py:695`).
+
+### 본 ADR을 규정하는 두 가지 발견 (ADR-0060 검토에서)
+
+- **Q1 — composite 발행 비용.** `tl.composite`를 construct + push 하는 비용은
+  CPU 시간 기준 **~40 ns** 정도(descriptor 빌드 + 큐 push)로 예상된다. 현재는
+  **0**이다. 훅은 존재하며, 값(그리고 op 종류별 차등)만 빠져 있다.
+- **Q2 — scheduler dispatch vs DMA latency.** PE_SCHEDULER의 composite
+  dispatch는 **non-blocking**이다: `_dispatch_composite`는 tile plan을 생성하고
+  feeder에 enqueue한 뒤 즉시 반환한다(`pe_scheduler.py:104-121`). 실제 **DMA
+  latency는 타일이 흐를 때 PE_DMA에서 부과**되며(`drain_ns =
+  compute_drain_ns(path, nbytes)`, `pe_dma.py:89`), scheduler dispatch에
+  덩어리로 들어가지 **않는다** ⇒ **중복 계상 없음**, DMA는 modeled component에
+  머문다(SPEC §0.1). scheduler의 plan 생성 자체는 현재 sim 시간 **0**이다.
+
+### 균일-그리고-0이 하이브리드에 틀린 이유
+
+하이브리드의 논거 전체는 **하나의** composite descriptor가 `N_tiles`만큼의
+GEMM 타일링을 offload하므로, CPU가 `O(N_tiles × ops/tile)`개의 fine 커맨드
+대신 `O(1)`개의 coarse 커맨드를 낸다는 것이다. 발행 비용 = 0(그리고 균일)이면
+모델은 다음을 보여줄 수 없다:
+- CPU가 충분히 빨리 못 밀어넣을 때 primitive 경로가 엔진을 **saturate 못 할 수
+  있음**(핵심 ADR-0060 §1 주장), 그리고
+- composite가 단일 primitive보다 construct 비용이 **더 크지만** 그것이 대체하는
+  다수의 primitive보다는 훨씬 작음.
+
+단일 균일 `dispatch_cycles`로는 이를 표현할 수 없다: 실제 construct 비용에서
+`tl.load` ≪ `tl.composite`이기 때문이다.
+
+## Decision
+
+### D1. PE_CPU의 op 종류별 발행 비용 테이블
+
+단일 `dispatch_cycles` 스칼라를 **커맨드 종류별 비용 테이블**로 교체하고,
+발행 시점(커맨드가 scheduler로 dispatch되기 전)에 PE_CPU에서 부과한다.
+greenlet이 이 시간을 지불하며, 이것이 바로 CPU가 얼마나 빨리 work를 밀어넣을
+수 있는지를 gate한다 — saturation 레버.
+
+시작 추정치(추후 calibrate — 검토 항목 참조):
+
+| 발행 op | CPU 발행 비용 (construct + push) |
+|---|---|
+| `tl.composite` (GEMM/MATH descriptor) | ~40 ns (Q1 추정) |
+| `tl.load` / `tl.store` (DMA descriptor) | 작음 (수 ns) |
+| `tl.dot` / MATH / IPCQ send/recv | 작음 (수 ns) |
+
+요점은 **비율**이다: composite construct ≫ primitive 발행, 그러나 그것이
+대체하는 primitive 발행들의 합 ≪. 정확한 ns는 설정 가능하다.
+
+### D2. 실행 latency는 엔진에 유지 (Q2 — 변경 없음)
+
+DMA/GEMM/MATH 실행 latency는 오늘처럼 PE_DMA / PE_GEMM / PE_MATH에서 부과된다.
+발행 비용(D1)은 **CPU 측에만** 추가되며, `drain_ns`를 건드리지 않으므로 중복
+계상이 없다. 이는 SPEC §0.1(latency는 modeled component에서)을 보존한다.
+
+### D3. scheduler plan 생성 비용 — 0에서 시작, 재검토
+
+PE_SCHEDULER의 tile-plan 생성은 오늘 sim 시간 0이다. 일단 유지한다(지배적
+레버는 CPU 발행 비용 D1); scheduler 측 비용이 유의미하다고 판명되면 기존
+`overhead_ns` node attr로 노출한다. 여기서 결정하지 않고 검토 항목으로 둔다.
+
+### D4. 설정 가능한 값; 골든 재생성
+
+비용 테이블은 설정 가능하다(per-topology / node attrs, 기본값은 D1 추정치).
+발행 비용을 0이 아니게 하면 **모든** bench latency가 바뀌므로, 골든 latency를
+한 번 **재생성**한다 — ADR-0062의 전역 lazy-load 회귀와 동일한 자세의 모델
+충실도 개선이다.
+
+## Alternatives
+
+### A1. 단일 균일 `dispatch_cycles > 0` 유지
+
+기각: op construct 비용은 자릿수 단위로 다르다(`tl.load` vs `tl.composite`).
+균일값은 primitive를 과대 청구하거나 composite를 과소 청구하며, 어느 쪽이든
+하이브리드가 의존하는 composite-vs-primitive 트레이드오프를 왜곡한다.
+
+### A2. 발행 비용을 PE_CPU 대신 PE_SCHEDULER에 부과
+
+기각: saturation 질문은 *"CPU가 엔진을 바쁘게 유지할 만큼 descriptor를 빨리
+밀어넣을 수 있는가?"*이며 — 이는 **PE_CPU**의 issue-bandwidth 속성이다.
+scheduler에 부과하면 CPU back-pressure를 모델하지 못한다.
+
+### A3. DMA program/setup 시간을 별도의 고정 per-descriptor 비용으로 모델
+
+실제 HW는 전송 시간과 별개로 DMA descriptor program 비용을 지불한다. 연기:
+초기에는 descriptor-program 비용을 **발행 op**의 D1 비용에 합친다(DMA를
+유발하는 `tl.load` / composite). calibration이 유의미함을 보이면 PE_DMA 고정
+setup으로 분리한다. 검토 항목.
+
+## Consequences
+
+### Positive
+- 하이브리드의 CPU-offload / saturation 이점(ADR-0060 §1)이 구조적일 뿐
+  아니라 **측정 가능**해진다.
+- composite-vs-primitive 및 타일링 granularity 트레이드오프가 latency에
+  보여 eval 서사를 가능케 한다.
+- 하드웨어에 더 충실(issue bandwidth는 실제 병목).
+
+### Negative
+- **모든** bench 골든이 이동 → 한 번의 재생성(D4); CI 골든 fixture 갱신.
+- op 종류별 값은 calibration 필요; ~40 ns composite 수치는 추정이고
+  primitive는 미지정 — 결과는 숫자만큼만 정확하다(calibrate 전까지 절대
+  latency를 과대 주장 금지).
+- 발행 경로에 비용 테이블 lookup 추가(런타임 영향 미미).
+
+## Open review items (자율 결정; 검토 시 수정)
+
+1. **op 종류별 값의 calibration 출처.** composite ≈ 40 ns는 작업 추정치;
+   primitive 발행 비용은 placeholder. *권고:* 문서화된 가정(instruction-issue +
+   큐 push)에서 값을 택하고, 실제 reference가 생기기 전까지 절대 latency를
+   잠정으로 취급; **비율**을 방어 가능하게 유지.
+2. **scheduler plan-gen 비용 (D3).** *권고:* 초기에는 0 유지; scheduler-bound
+   동작을 보이는 workload가 나오면 `overhead_ns`로 노출.
+3. **DMA program 시간 (A3).** *권고:* 먼저 발행 op의 비용에 합치고; 필요 시
+   PE_DMA setup으로 분리.
+4. **테이블 위치.** *권고:* 흩어진 node attr이 아니라, 커맨드 종류로 키잉되고
+   per-topology 오버라이드 가능한 작은 중앙 cost 모듈 — 값을 한 곳에서 검토
+   가능하도록.
+5. **회귀 롤아웃.** *권고:* lazy-load(ADR-0062)와 본 ADR의 cost model을 한
+   번의 골든 재생성 패스로 함께 도입하여 두 번의 churn을 피함.
+6. **legacy replay 경로와의 상호작용**(`pe_cpu.py:_execute_legacy`) — greenlet과
+   replay 경로가 동일한 cost 테이블을 읽어 결과가 일치하도록 보장. 검증 요.
+
+## Test Requirements
+
+1. **composite는 한 번, primitive는 op마다 청구.** `N_tiles`에 걸쳐 하나의
+   `tl.composite`를 내는 커널은 PE_CPU에서 composite 발행 비용을 한 번 청구하고;
+   동등한 `N_tiles × ops` primitive 커널은 op당 비용을 `N_tiles × ops`번 청구.
+   PE_CPU busy time이 그에 따라 다름을 assert.
+2. **DMA latency 불변 (Q2).** 고정된 전송에 대해 PE_DMA `drain_ns`는 ADR 이전
+   값과 동일 — 발행 비용은 DMA에 합쳐지지 않고 PE_CPU에 가산(중복 계상 없음).
+3. **saturation 관측 가능.** per-op 발행 비용이 0이 아니면, many-tile primitive
+   sweep는 GEMM 엔진 idle(CPU-bound 발행)을 보이고 composite sweep는 바쁘게
+   유지 — ADR-0060 §1 레버.
+4. **결정성:** 동일 입력 → 동일 op_log + latency (SPEC §0.1).
+5. **경로 일치:** greenlet과 legacy-replay 경로가 동일 커널에 대해 동일한 발행
+   비용 회계를 산출.

docs/adr-proposed/ADR-0064-perf-cpu-issue-cost-model.md

@@ -0,0 +1,179 @@
+# ADR-0064: Per-op-type CPU issue cost model (command construct + dispatch)
+
+## Status
+
+Proposed
+
+> Supporting ADR for **ADR-0060** (AHBM GQA Fused Attention). The hybrid
+> decision there (GEMMs via `tl.composite`, softmax merge in the kernel)
+> wins by **offloading tiling to PE_SCHEDULER so the CPU issues coarse
+> descriptors and runs ahead, keeping the engines saturated**. That win is
+> currently **invisible in the simulator** because per-op CPU issue cost is
+> zero. This ADR makes the issue cost real and **op-type-differentiated**,
+> so composite-vs-primitive trade-offs (and CPU saturation) are measurable.
+
+## Context
+
+### What exists today
+
+- Every `tl.*` op calls `_emit_dispatch_overhead()` before emitting its
+  command (`tl_context.py:123-125, 190, 227, 235, 612, …`), which emits
+  `PeCpuOverheadCmd(cycles=dispatch_cycles)` **only if** `dispatch_cycles
+  > 0`. So the issue cost is a **single uniform knob** applied identically
+  to `tl.load`, `tl.dot`, a MATH op, and `tl.composite`.
+- That knob is hardcoded to **0** in both live execution paths
+  (`pe_cpu.py:101` greenlet runner, `:195` legacy replay). ⇒ issuing a
+  command — *constructing the descriptor and pushing it to the scheduler
+  queue* — currently costs **0 ns** on PE_CPU.
+- `PeCpuOverheadCmd` is consumed as `yield env.timeout(cmd.cycles)` on
+  PE_CPU (`kernel_runner.py:131-132`) and on PE_SCHEDULER
+  (`pe_scheduler.py:97-100`); a manual `tl.cycles(n)` also exists
+  (`tl_context.py:695`).
+
+### Two findings that frame this ADR (from ADR-0060 review)
+
+- **Q1 — composite issue cost.** Constructing + pushing a `tl.composite`
+  is expected to cost on the order of **~40 ns** of CPU time (descriptor
+  build + queue push). Today it is **0**. The hook exists; only the value
+  (and its per-op-type differentiation) is missing.
+- **Q2 — scheduler dispatch vs DMA latency.** `PE_SCHEDULER`'s composite
+  dispatch is **non-blocking**: `_dispatch_composite` generates the tile
+  plan and enqueues to the feeder, returning immediately
+  (`pe_scheduler.py:104-121`). The actual **DMA latency is charged on
+  PE_DMA** as each tile flows through (`drain_ns = compute_drain_ns(path,
+  nbytes)`, `pe_dma.py:89`), **not** lumped into the scheduler dispatch ⇒
+  **no double-counting**, and DMA stays on its modelled component
+  (SPEC §0.1). The scheduler's own plan-generation currently costs **0**
+  sim time.
+
+### Why uniform-and-zero is wrong for the hybrid
+
+The hybrid's whole argument is that **one** composite descriptor offloads
+`N_tiles` worth of GEMM tiling, so the CPU issues `O(1)` coarse commands
+instead of `O(N_tiles × ops/tile)` fine ones. With issue cost = 0 (and
+uniform), the model cannot show:
+- that the primitive path may **fail to saturate** the engines when the
+  CPU cannot push fast enough (the core ADR-0060 §1 claim), nor
+- that a composite **costs more to construct** than a single primitive but
+  far less than the many primitives it replaces.
+
+A single uniform `dispatch_cycles` cannot express this: `tl.load` ≪
+`tl.composite` in real construct cost.
+
+## Decision
+
+### D1. Per-op-type issue-cost table on PE_CPU
+
+Replace the single `dispatch_cycles` scalar with a **cost table keyed by
+command type**, charged on PE_CPU at issue (before the command is
+dispatched to the scheduler). The greenlet pays this time, which is
+exactly what gates how fast the CPU can push work — the saturation lever.
+
+Starting estimates (calibrate later — see review items):
+
+| Issued op | CPU issue cost (construct + push) |
+|---|---|
+| `tl.composite` (GEMM/MATH descriptor) | ~40 ns (Q1 estimate) |
+| `tl.load` / `tl.store` (DMA descriptor) | small (a few ns) |
+| `tl.dot` / MATH / IPCQ send/recv | small (a few ns) |
+
+The point is the **ratio**: a composite construct ≫ a primitive issue, but
+≪ the sum of the primitive issues it replaces. Exact ns are configurable.
+
+### D2. Execution latency stays on the engines (Q2 — no change)
+
+DMA/GEMM/MATH execution latency remains charged on PE_DMA / PE_GEMM /
+PE_MATH as today. The issue cost (D1) is **CPU-side only** and additive; it
+does **not** touch `drain_ns`, so there is no double-count. This preserves
+SPEC §0.1 (latency on modelled components).
+
+### D3. Scheduler plan-generation cost — start at 0, revisit
+
+`PE_SCHEDULER`'s tile-plan generation costs 0 sim time today. Keep that for
+now (the dominant lever is CPU issue cost, D1); expose it later via the
+existing `overhead_ns` node attr if a scheduler-side cost proves material.
+Marked as a review item, not decided here.
+
+### D4. Configurable values; goldens regenerate
+
+The cost table is configurable (per-topology / node attrs, default to the
+D1 estimates). Turning issue cost non-zero changes **every** bench's
+latency, so golden latencies are **regenerated** once — a model-fidelity
+improvement, same posture as ADR-0062's global lazy-load regression.
+
+## Alternatives
+
+### A1. Keep a single uniform `dispatch_cycles > 0`
+
+Rejected: op construct costs differ by an order of magnitude (`tl.load` vs
+`tl.composite`); a uniform value either over-charges primitives or
+under-charges composites, and in both cases misrepresents the
+composite-vs-primitive trade-off the hybrid depends on.
+
+### A2. Charge the issue cost on PE_SCHEDULER instead of PE_CPU
+
+Rejected: the saturation question is *"can the CPU push descriptors fast
+enough to keep the engines busy?"* — that is a **PE_CPU** issue-bandwidth
+property. Charging it on the scheduler would not model CPU back-pressure.
+
+### A3. Model DMA program/setup time as a separate fixed per-descriptor cost
+
+Real HW pays a DMA descriptor program cost distinct from transfer time.
+Deferred: initially fold the descriptor-program cost into the **issuing
+op's** D1 cost (a `tl.load` / composite that triggers DMA). Split it out to
+a PE_DMA fixed setup only if calibration shows it matters. Review item.
+
+## Consequences
+
+### Positive
+- The hybrid's CPU-offload / saturation win (ADR-0060 §1) becomes
+  **measurable**, not just structural.
+- Composite-vs-primitive and tiling-granularity trade-offs are visible in
+  latency, enabling the eval narrative.
+- More faithful to hardware (issue bandwidth is a real bottleneck).
+
+### Negative
+- **All** bench goldens shift → one-time regeneration (D4); CI golden
+  fixtures update.
+- Per-op-type values need calibration; the ~40 ns composite figure is an
+  estimate, primitives are unspecified — results are only as good as the
+  numbers (do not over-claim absolute latencies until calibrated).
+- Adds a cost-table lookup on the issue path (negligible runtime).
+
+## Open review items (decided autonomously; revise on review)
+
+1. **Calibration source for the per-op-type values.** Composite ≈ 40 ns is
+   a working estimate; primitive issue costs are placeholders. *Recommend:*
+   pick values from a documented assumption (instruction-issue + queue
+   push) and treat absolute latency as provisional until a real reference
+   exists; keep the **ratios** defensible.
+2. **Scheduler plan-gen cost (D3).** *Recommend:* keep 0 initially; expose
+   via `overhead_ns` if a workload shows scheduler-bound behaviour.
+3. **DMA program time (A3).** *Recommend:* fold into the issuing op's cost
+   first; split to PE_DMA setup only if needed.
+4. **Where the table lives.** *Recommend:* a small central cost module
+   keyed by command type, overridable per topology — not scattered node
+   attrs — so the values are reviewable in one place.
+5. **Regression rollout.** *Recommend:* land lazy-load (ADR-0062) and this
+   ADR's cost model in one golden-regeneration pass to avoid two churns.
+6. **Interaction with the legacy replay path** (`pe_cpu.py:_execute_legacy`)
+   — ensure both the greenlet and replay paths read the same cost table so
+   results match. Verify.
+
+## Test Requirements
+
+1. **Composite charges once; primitives charge per-op.** A kernel issuing
+   one `tl.composite` over `N_tiles` charges one composite issue cost on
+   PE_CPU; the equivalent `N_tiles × ops` primitive kernel charges the
+   per-op cost `N_tiles × ops` times. Assert the PE_CPU busy time differs
+   accordingly.
+2. **DMA latency unchanged (Q2).** For a fixed transfer, PE_DMA `drain_ns`
+   is identical to the pre-ADR value — the issue cost is additive on
+   PE_CPU, not folded into DMA (no double-count).
+3. **Saturation is observable.** With non-zero per-op issue cost, a
+   many-tile primitive sweep shows GEMM-engine idle (CPU-bound issue)
+   whereas the composite sweep keeps it busy — the ADR-0060 §1 lever.
+4. **Determinism:** identical inputs → identical op_log + latency
+   (SPEC §0.1).
+5. **Path parity:** greenlet and legacy-replay paths produce identical
+   issue-cost accounting for the same kernel.

docs/adr-proposed/DDD-0060-gqa-fused-attention-detailed-design.md

@@ -0,0 +1,532 @@
+# Detailed Design Document — AHBM GQA Fused Attention (ADR-0060)
+
+**Status:** Draft (companion to ADR-0060, Proposed)
+**Scope:** implementation-ready design for an efficient GQA FlashAttention
+kernel on AHBM in the kernbench simulator.
+**Audience:** reviewer resuming with *"GQA 검토 시작하자"*. Read ADR-0060
+first (decision record), then this (the how), then the *Open Decisions*
+section (§10) for the choices made autonomously that need your sign-off.
+
+> **How this document was produced.** It is the result of several
+> design↔evaluation iterations against the real kernbench source
+> (`src/kernbench/triton_emu/tl_context.py`, `sim_engine/*`,
+> `components/builtin/pe_*`, the `_attention_mesh_*` kernels, and the GQA
+> tests). Every claim about what exists is grounded in a file path; every
+> claim about what is new is tied to a supporting ADR (0061/0062/0063).
+
+---
+
+## 1. Goal and success criteria
+
+**Goal:** an *efficient* GQA fused attention kernel that runs on AHBM in
+kernbench, across the four operating points decode/prefill ×
+single-/multi-user, at realistic (not just validation) scale, in both
+timing (`enable_data=False`) and data (`enable_data=True`) modes.
+
+**Success criteria:**
+1. Correct: kernel output matches a numpy FlashAttention reference within
+   fp tolerance, with **real GQA** (`H_q=64, H_kv=8, G=8`).
+2. Efficient: the three levers actually fire and are observable in op_log:
+   GQA K/V-load amortisation, KV prefetch overlap, `⌈log₂ N⌉` reduction.
+3. Scales: context length not capped by the 1 MiB scratch bump allocator.
+4. Deterministic and SPEC-compliant (all latency from modelled events).
+
+**Non-goals:** RoPE / QKV projection / KV-cache writes (upstream
+`qkv_rope`); output projection (downstream); cross-CUBE query-head split;
+cycle-accurate microarchitecture (SPEC §5).
+
+---
+
+## 2. Where the design sits in the current code
+
+```
+runtime_api (host)
+  RuntimeContext.zeros/empty/launch        context.py  (tensor deploy, DPPolicy, kernel launch)
+        │  KernelLaunchMsg
+        ▼
+sim_engine
+  GraphEngine(enable_data=…)               engine.py
+  DataExecutor (Phase 2)                   data_executor.py   ← ADR-0061 adds _execute_broadcast
+  MemoryStore (nbytes check)               memory_store.py    ← the GQA blocker lives here
+        │
+        ▼
+components / PE pipeline (per PE)
+  PE_CPU → PE_SCHEDULER → PE_{DMA,GEMM,MATH,FETCH_STORE,TCM,IPCQ}
+  greenlet kernel ↔ SimPy                  kernel_runner.py, pe_cpu.py
+        │
+        ▼
+tl programming model                       tl_context.py
+  load/store/dot/softmax/max/sum/exp/maximum/send/recv/trans/…
+  ← ADR-0061 adds broadcast/repeat
+  ← ADR-0062 adds load_async (+ wait arm)
+  ← ADR-0063 adds scratch_scope
+```
+
+The attention kernels are **bench-side Python** (`src/kernbench/benches/`)
+invoked via `ctx.launch(name, kernel_fn, *args)`; they run as greenlet
+kernels that emit `tl.*` ops. No component (device-side) code changes are
+needed for the algorithm — only the three `tl`/sim_engine primitives.
+
+---
+
+## 3. File plan
+
+### 3.1 New / modified production files
+
+| File | Change | ADR |
+|---|---|---|
+| `src/kernbench/triton_emu/tl_context.py` | add `broadcast`/`repeat`; `load_async` + `wait` arm; `scratch_scope` ctx-mgr | 0061/0062/0063 |
+| `src/kernbench/common/pe_commands.py` | add `BroadcastCmd`; `DmaReadCmd.blocking` flag (or `DmaReadAsyncCmd`) | 0061/0062 |
+| `src/kernbench/sim_engine/data_executor.py` | add `_execute_broadcast`; handle async dma_read resolution | 0061/0062 |
+| `src/kernbench/components/builtin/pe_dma.py` | non-blocking read path returns a future-resolving event | 0062 |
+| `src/kernbench/sim_engine/event_log.py` / op_log | recognise `broadcast` op_name (math) | 0061 |
+
+### 3.2 New kernel + bench files
+
+| File | Role |
+|---|---|
+| `src/kernbench/benches/_gqa_attention.py` | the unified GQA kernel (4 cases via flags), with a shared `_flash_tile_sweep` helper and `_tree_reduce` helper |
+| `src/kernbench/benches/_gqa_helpers.py` | `valid_len`, `tile_*` predicates, mask builders, tree topology (children/parent dirs), GQA-batched Q load |
+| extend `src/kernbench/benches/milestone_gqa_llama70b.py` | add headline-scale panels with real `G`, prefetch, tiling |
+
+### 3.3 New tests
+
+| File | Covers |
+|---|---|
+| `tests/attention/test_gqa_correctness.py` | numpy reference parity, all 4 cases, real GQA |
+| `tests/attention/test_gqa_levers.py` | GQA load amortisation, tree step count, causal skip, prefetch overlap |
+| `tests/test_tl_broadcast.py` | ADR-0061 unit |
+| `tests/test_tl_load_async.py` | ADR-0062 unit |
+| `tests/test_tl_scratch_scope.py` | ADR-0063 unit |
+
+---
+
+## 4. Data model and launch contract
+
+### 4.1 Tensor placement (DPPolicy)
+
+Using `DPPolicy(cube, pe, num_cubes, num_pes)` (`policy/placement/dp.py`):
+
+| Tensor | single-user (PE ring in 1 cube) | multi-user (cube ring) |
+|---|---|---|
+| `Q` (group, replicated) | `cube=replicate, pe=replicate, num_pes=N` | `cube=replicate, pe=replicate, num_cubes=N` |
+| `K`, `V` (sequence-sharded) | `cube=replicate, pe=row_wise, num_pes=N` | `cube=row_wise, pe=replicate, num_cubes=N` |
+| `O` (group, replicated) | as `Q` | as `Q` |
+
+This matches the baseline `_PANEL_DISPATCH`
+(`milestone_gqa_llama70b.py:64-81`). `row_wise` over the sequence (M) axis
+realises the round-robin shard (§2.1 of ADR-0060); the contiguous per-PE
+buffer is the row-wise slice.
+
+`rank_axis` selects which `tl.program_id` carries the ring rank: `0` for
+single-user (PE id in cube), `1` for multi-user (cube level; only PE 0 of
+each cube participates) — kept from the baseline so the SFR installs
+(`configure_sfr_intracube_pe_ring` / `configure_sfr_intercube_multisip`)
+apply unchanged.
+
+### 4.2 Launch signature
+
+```python
+ctx.launch(
+    f"{panel}_gqa", gqa_attention_kernel,
+    q, k, v, o,                       # Tensors (TensorArg)
+    S_q, S_kv_per_rank, H_q, H_kv, D_HEAD, n_ranks,   # scalars
+    rank_axis, mode, counter, start_pe,               # scalars (mode: 0=decode,1=prefill)
+    softmax_scale,                                     # float scalar
+    _auto_dim_remap=False,            # keep d_head from colliding (baseline note)
+)
+```
+
+`H_q > H_kv` is the change from the baseline's `H_q=H_kv=1`; it requires
+ADR-0061 to survive Phase 2.
+
+### 4.3 GQA-batched Q load — fold `G` into the matmul M dim (no broadcast)
+
+For a KV head `kv`, its group is query heads `[kv·G, (kv+1)·G)`. Load
+those `G` rows and **reshape the group into the matmul M dimension** —
+this is the key to GQA reuse and it needs only a byte-conserving reshape,
+**not** a broadcast:
+
+```python
+def load_q_group(q_ptr, kv, G, T_q, d, tl):
+    # Q laid out [H_q, T_q, d]; the group is contiguous in the head axis.
+    qg = tl.load(q_ptr + kv*G*T_q*d, shape=(G, T_q, d), dtype="f16")
+    return _view(qg, (G*T_q, d))            # byte-conserving reshape → M = G*T_q
+```
+
+Then per KV tile (`Kj`, `Vj` are `[TILE, d]`, one KV head, loaded once):
+
+```python
+Sj  = tl.dot(q_g, tl.trans(Kj)) * scale     # [G*T_q, d]·[d, TILE] -> [G*T_q, TILE]
+Oj  = tl.dot(P,   Vj)                        # [G*T_q, TILE]·[TILE, d] -> [G*T_q, d]
+```
+
+K/V are the **shared operands** (the matmul's `b`), loaded once and reused
+across all `G·T_q` M-rows automatically. The emitted `GemmCmd` carries
+`m = G·T_q`, so Phase-1 timing counts all `G` rows' work correctly.
+
+> **Why not a leading batch axis** (`[G,T_q,d]·[G,d,TILE]`)? `tl.dot`
+> derives `m,k,n` from the **last two dims only** (`tl_context.py:220`)
+> and `GemmCmd` carries no batch dim — so a leading `G` axis is computed
+> correctly by `np.matmul` in Phase 2 (it broadcasts) but **under-counted
+> by `G×` in Phase-1 timing**. Folding into M fixes both. This supersedes
+> the earlier batched-vs-loop question.
+
+> **Transpose caveat.** `tl.trans` is metadata-only and `MemoryStore.read`
+> *reshapes* (not transposes) (`memory_store.py:73`,
+> `data_executor.py:153-168` does `np.matmul` on the reshaped operands).
+> Fine for timing/structure and for the baseline's zero data; for
+> non-trivial numeric parity, store K pre-transposed as `[d, S_pe]` or add
+> a real `tl.transpose` (§10.10).
+
+---
+
+## 5. Kernel design (concrete, real `tl`)
+
+### 5.1 Shared tile sweep helper
+
+```python
+def flash_tile_sweep(q_g, k_base, v_base, n_tiles, TILE, d, scale,
+                     q_block, causal, tl, PREFETCH=2):
+    """One PE's FlashAttention sweep over its local KV tiles.
+    Returns running (m, l, O) in the persistent arena."""
+    m, l, O = _init_running(q_g.shape, d, tl)           # persistent arena
+    f_k, f_v = {}, {}
+    for j in range(min(PREFETCH, n_tiles)):             # prime (ADR-0062)
+        if causal and tile_all_future(j, q_block): continue
+        f_k[j] = tl.load_async(k_base + j*TILE*d, (TILE, d))
+        f_v[j] = tl.load_async(v_base + j*TILE*d, (TILE, d))
+    for j in range(n_tiles):
+        if causal and tile_all_future(j, q_block): continue
+        with tl.scratch_scope():                        # recycle temporaries (ADR-0063)
+            Kj = tl.wait(f_k[j])
+            Sj = tl.dot(q_g, tl.trans(Kj)) * scale
+            mask = make_mask_tile(j, q_block) if (causal and tile_partial(j, q_block)) else None
+            if mask is not None: Sj = Sj + mask
+            m_j   = tl.max(Sj, axis=-1)
+            m_new = tl.maximum(m, m_j)
+            P     = tl.exp(Sj - m_new)
+            corr  = tl.exp(m - m_new)
+            l     = l * corr + tl.sum(P, axis=-1)
+            Vj    = tl.wait(f_v[j])
+            O     = O * corr + tl.dot(P, Vj)
+            m     = m_new
+            nxt = j + PREFETCH                           # refill pipeline
+            if nxt < n_tiles and not (causal and tile_all_future(nxt, q_block)):
+                f_k[nxt] = tl.load_async(k_base + nxt*TILE*d, (TILE, d))
+                f_v[nxt] = tl.load_async(v_base + nxt*TILE*d, (TILE, d))
+    return m, l, O
+```
+
+> Note: `m, l, O` are reassigned inside the `with` but their **storage**
+> must be the persistent arena, not the recycled scope. Concretely the
+> merge ops write to pre-allocated persistent handles (double-buffered);
+> ADR-0063 §D3 describes the persistent-vs-scoped arena split. This is the
+> one subtlety the implementer must get right — see §10.5.
+
+### 5.2 Tree reduction helper
+
+```python
+def tree_reduce_and_store(m, l, O, pe_id, N, o_ptr, tl):
+    for cdir in tree_children_dirs(pe_id, N):           # static tree (§4 ADR-0060)
+        m_c = tl.recv(cdir, m.shape); l_c = tl.recv(cdir, l.shape)
+        O_c = tl.recv(cdir, O.shape)
+        mn  = tl.maximum(m, m_c); sa = tl.exp(m - mn); sb = tl.exp(m_c - mn)
+        l   = l*sa + l_c*sb; O = O*sa + O_c*sb; m = mn
+    if is_root(pe_id, N):
+        tl.store(o_ptr, O / l)
+    else:
+        pdir = parent_dir(pe_id, N)
+        tl.send(pdir, m); tl.send(pdir, l); tl.send(pdir, O)
+```
+
+`tree_children_dirs` / `parent_dir` map the binary tree (level pairs in
+ADR-0060 §4) onto physical `N/S/E/W` neighbours — this mapping is the
+tuning item ADR-0060 §9.2 and an Open Decision (§10.2).
+
+### 5.3 Top-level kernel (all four cases)
+
+```python
+def gqa_attention_kernel(q_ptr,k_ptr,v_ptr,o_ptr,
+                         S_q,S_kv_per_rank,H_q,H_kv,d,n_ranks,
+                         rank_axis,mode,counter,start_pe,scale,*,tl):
+    if rank_axis != 0 and tl.program_id(axis=0) != 0:   # multi-user: only PE0 per cube
+        return
+    pe_id = tl.program_id(axis=rank_axis)
+    G = H_q // H_kv
+    causal = (mode == 1)                                # prefill
+    q_block = (counter, S_q) if causal else (counter, 1)
+    # decode small-S fast path = one tile; otherwise tiled (TILE from config)
+    for kv in range(H_kv):                              # one KV head per iter (or PE-mapped)
+        q_g  = load_q_group(q_ptr, kv, G, S_q, d, tl)
+        my_len  = valid_len(counter, start_pe, pe_id, n_ranks) if mode==0 else S_kv_per_rank
+        n_tiles = ceil(my_len / TILE)
+        m,l,O = flash_tile_sweep(q_g, k_base(kv), v_base(kv), n_tiles, TILE, d,
+                                 scale, q_block, causal, tl)
+        if n_ranks == 1:
+            tl.store(o_ptr_for(kv), O / l)
+        elif mode == 1 and is_ring(...):                # prefill ring already folded across steps
+            ring_fold_then_store(...)                   # §5.4
+        else:
+            tree_reduce_and_store(m,l,O, pe_id, n_ranks, o_ptr_for(kv), tl)
+```
+
+This collapses the baseline's two separate kernels (`_attention_mesh_kv`,
+`_attention_mesh_mlo`) into one parameterised kernel; the `mode`/`n_ranks`
+flags select the case. (Keeping them separate is also fine — Open
+Decision §10.6.)
+
+### 5.4 Ring (prefill SP) specifics
+
+The Ring path keeps the baseline's structure (`_attention_mesh_kv.py`)
+with three upgrades:
+- `recv_async` for the next step's K/V (overlap) — `tl_context.py:543`.
+- causal step-skip: `if step_all_future(step, q_block): continue` before
+  folding (skips ≈half the steps).
+- GQA: `q_g` is the `G`-row group; `tl.broadcast` for the KV head.
+
+The fold across ring steps is the same online-softmax merge as §5.2's
+inner update — running `(m,l,O)` carried across steps as Python handles
+(exactly as the baseline does today).
+
+---
+
+## 6. The three new primitives — integration points
+
+### 6.1 ADR-0061 `tl.broadcast` (OPTIONAL — not on the GQA critical path)
+GQA reuse is achieved by folding `G` into M (§4.3), so broadcast is **not**
+required to unblock `h_q > h_kv`. It remains a clean primitive for additive
+causal-mask construction across the `G·T_q` rows and for general kernels.
+- `tl_context.py`: new `broadcast(x, new_shape)` → emits `BroadcastCmd`,
+  output handle from `_make_compute_out` (so `nbytes` is correct).
+- `data_executor.py`: `_execute_broadcast` = `np.broadcast_to(src,
+  shape).copy()` → `store.write`.
+- Latency: `pe_math._compute_ns(prod(new_shape))`; logs `math/broadcast`,
+  **no** `dma_read`.
+- Lowest priority; deliver after the efficiency/scale features land.
+
+### 6.2 ADR-0062 `tl.load_async`
+- `tl_context.py`: `load_async(ptr, shape, dtype) → LoadFuture`; extend
+  `wait` to resolve `LoadFuture`.
+- `pe_dma.py`: non-blocking read path; transfer occupies the read channel
+  in parallel; completion is a scheduled event.
+- op_log unchanged (`memory/dma_read`) ⇒ `dma_read_count` metric stable.
+
+### 6.3 ADR-0063 `tl.scratch_scope`
+- `tl_context.py`: context manager saving/restoring `_scratch_cursor`.
+- Persistent arena: `m,l,O` (+ double-buffer) and live `LoadFuture`
+  buffers allocated outside the scope.
+- Removes the `S=16` cap (`test_milestone_gqa_llama70b.py:123-148`).
+
+---
+
+## 7. Phased implementation plan (test-first per CLAUDE.md)
+
+Each phase is an independent Phase-1→Phase-2 cycle. Ordered so each builds
+on a green predecessor; every phase keeps the existing baseline green.
+
+| Phase | Deliverable | Gate |
+|---|---|---|
+| **P1** | Real GQA by restructuring (per KV head, `G`→M fold); one-shot, no tiling. **No new primitive.** | data-mode run completes for `h_q=G·h_kv`; K/V `dma_read_count` independent of `G`; `gemm` m=`G·T_q` |
+| **P2** | Tree reduction replacing fan-out (decode SP) | `⌈log₂ N⌉` reduction rounds; structure preserved |
+| **P3** | ADR-0063 `tl.scratch_scope` + tiled sweep | long-`S` decode completes (today caps at S=16) |
+| **P4** | ADR-0062 `tl.load_async` + prefetch in sweep | latency < serial Σ(load+compute) |
+| **P5** | Causal masking + tile/step skip (prefill, ring) | GEMM count = triangular tile count |
+| **P6** | Headline-scale milestone panels (real `G`, tiling, prefetch) + figures | sweep.json headline rows; figures render |
+| **P7** *(opt)* | ADR-0061 `tl.broadcast` (mask convenience); optional `tl.transpose` for numeric parity | broadcast unit test; asymmetric numeric parity |
+
+**P1 alone delivers "GQA actually runs" — and it needs no new feature**,
+only the §4.3/§5 restructuring (the iteration-2 finding). P3 is the most
+important *new* feature (removes the S=16 scale ceiling). P4 is the
+efficiency lever. P2/P5 are algorithmic. P7 is optional polish.
+
+---
+
+## 8. Verification plan (concrete)
+
+Grounded in SPEC R2/R5, ADR-0023/0025 (IPCQ), ADR-0046 (`tl`),
+ADR-0054 (eval bench). Mirrors ADR-0060 §11.
+
+**Runs + numeric (Phase 2, `enable_data=True`):**
+- `test_gqa_correctness.py`: for `(G∈{1,8}, T_q∈{1,16}, S, N∈{1,8})` the
+  GQA kernel **completes** (no byte-conservation error) — the `G=8` cases
+  that the baseline head-packing cannot express. This is gated on the
+  §4.3 restructuring, **not** on a new primitive.
+- *Numeric parity (secondary):* `O ≈ numpy FlashAttention` for
+  symmetric/identity inputs where reshape-as-transpose is exact; full
+  asymmetric parity is gated on an optional `tl.transpose` (§10.10).
+
+**Levers (op_log assertions):**
+- GQA amortisation: K/V `dma_read_count` constant as `G` grows 1→8;
+  `gemm` work grows with `G`.
+- Tree: decode-SP `ipcq_send`/`recv` counts = tree edges, reduction
+  rounds = `⌈log₂ N⌉` (not `N−1`).
+- Causal skip: prefill `gemm_count` = lower-triangular tile count.
+- Prefetch overlap: tiled-sweep latency strictly < `Σ(load+compute)`.
+- Scratch: peak `_scratch_cursor` bounded by one tile, independent of
+  `n_tiles`.
+
+**Invariants:**
+- Determinism: identical inputs → identical op_log + latency (SPEC §0.1).
+- Every routed request latency > 0 (SPEC §0.1).
+
+---
+
+## 9. Performance model (what we expect to see)
+
+Rough per-PE decode cost (KV-load-bound), with prefetch overlap:
+
+```
+t_decode_pe ≈ max( DMA(K+V over S_pe tiles),  compute(QKᵀ + PV over S_pe, ×G amortised) )
+            + t_tree_reduce(⌈log₂ N⌉ hops, payload O[G,d])
+```
+
+Levers and their expected signature in op_log / latency:
+- **GQA reuse** → KV bytes moved are `H_kv·S·d` not `H_q·S·d` (8× less),
+  with compute still `H_q`-scaled. Decode being KV-load-bound, this is the
+  dominant decode win.
+- **Prefetch** → DMA hidden behind compute ⇒ `t ≈ max(...)` not sum.
+- **Tree** → reduction `⌈log₂ N⌉` (3 for N=8) vs baseline `N−1` (7).
+- **Causal skip** (prefill) → ≈½ the tile/step work.
+
+These are the quantities the milestone bench's op_log summary + a new
+latency column should report (Open Decision §10.3 on adding latency to the
+sweep JSON).
+
+---
+
+## 10. Open Decisions / Assumptions made autonomously (REVIEW THESE)
+
+Choices I made to keep progressing without your input. Each is reversible;
+none is yet implemented in production code (docs only). Marked
+**[recommend]** = my default if you don't object.
+
+### 10.1 Greenlet `tl` model over composites **[recommend: greenlet]**
+The biggest design call. ADR-0060 §1/§8 argues the greenlet path is
+equally efficient in kernbench's per-op latency model and far smaller to
+build than a flash-composite. **Risk:** if you specifically want the
+*composite* abstraction exercised for the eval narrative (e.g. to study
+scheduler-managed tiling), we'd instead invest in a flash-composite
+command kind. I judged that out of proportion to the goal. *Confirm the
+greenlet direction.*
+
+### 10.2 Tree reduction topology / neighbour mapping **[recommend: binary tree on N/S/E/W]**
+ADR-0060 §4 fixes a binary tree (N=8: root=7). The mapping of tree edges
+to physical mesh directions depends on the SFR install's neighbour
+wiring. I assumed `configure_sfr_intracube_pe_ring` gives a ring on which
+the tree pairs are 1-hop neighbours. **Needs verification** against the
+actual installed neighbour table; if pairs aren't 1-hop, either relabel
+PEs or accept multi-hop sends. *Confirm or hand me the intended mapping.*
+
+### 10.3 Keep the baseline all-to-all as an option? **[recommend: replace with tree]**
+The baseline fan-out gives the answer on *every* rank; the tree gives it
+on root only. Attention needs root-only, so tree wins. If any downstream
+consumer needs O replicated (none found), we'd keep fan-out. *Assuming
+root-only is fine.*
+
+### 10.4 GQA matmul shape — RESOLVED in iteration 2 **[fold G into M]**
+Earlier this asked "batched leading axis vs per-`g` loop." Resolved: do
+**neither** — fold `G·T_q` into the matmul M dimension (§4.3). A leading
+`G` axis would be computed by `np.matmul` in data mode but **under-counted
+`G×` in Phase-1 timing** (`GemmCmd` carries no batch dim,
+`tl_context.py:220`). Folding into M gives correct timing *and* correct
+data with a single byte-conserving reshape and one shared K/V operand. No
+new primitive, no per-`g` loop.
+
+### 10.5 Persistent-vs-scoped arena split **[recommend: two explicit arenas]**
+ADR-0063 requires `m,l,O` to live outside the recycled scope. The cleanest
+implementation reserves a small fixed persistent region at kernel start
+and double-buffers the merge. I assumed this is acceptable extra
+bookkeeping inside the shared helper. *Alternative:* make `scratch_scope`
+take an explicit "keep-alive" handle list. I picked the two-arena model
+for clarity.
+
+### 10.6 One unified kernel vs two (decode/prefill) **[recommend: unify]**
+§5.3 unifies into one parameterised kernel; the baseline has two. Unifying
+reduces duplication but makes the kernel branch-heavy. *Alternative:* keep
+`_gqa_decode` and `_gqa_prefill` separate, sharing helpers. Low stakes;
+either is fine.
+
+### 10.7 Ghost ADRs 0055–0059 **[recommend: backfill as Draft, separate task]**
+The baseline cites ADR-0055/0056/0057/0058/0059 which **do not exist**
+(`grep` finds refs in code/tests, no files; no git history of them). This
+is pre-existing documentation debt, not created by this work. I did **not**
+backfill them (out of scope, and reverse-engineering intent is risky).
+*Recommendation:* a separate `Draft` ADR-backfill task to document the
+existing mesh kernels + milestone bench + SFR install, so ADR-0060 has a
+real predecessor to "supersede". *Confirm you want this as follow-up.*
+
+### 10.8 Numbering **[recommend: as chosen]**
+GQA ADR = **0060**; supporting = **0061/0062/0063**. I avoided 0055–0059
+(claimed by ghost refs) to prevent collisions. The file was renamed from
+`ADR-XXXX-…` to `ADR-0060-algo-…` (title was erroneously "ADR-001").
+
+### 10.9 bf16 modelled as f16 **[assumption, pre-existing]**
+`memory_store.py:16` maps `bf16→float16` (numpy has no bf16). Numeric
+parity tests must use tolerances that accommodate this, and we should not
+claim bf16 *accuracy* results — only timing/structure. Pre-existing
+simulator limitation; flagged so results aren't over-interpreted.
+
+### 10.10 `tl.trans` is reshape-not-transpose **[decision: structural-first verify]**
+Discovered in iteration 2: `tl.trans` only swaps shape metadata
+(`tl_context.py:390`) and `MemoryStore.read` *reshapes* the bytes
+(`memory_store.py:73`); `data_executor._execute_gemm` then `np.matmul`s
+the reshaped operands (`data_executor.py:153-168`). So for **non-trivial
+numeric data**, `Q·Kᵀ` via `tl.trans(K)` is mathematically a reshape, not
+a transpose — wrong numbers (the baseline never notices: it runs on
+all-zeros and only asserts op_log structure). **Decision:** treat the
+simulator as a *performance* model (SPEC §0) — verify structure + timing +
+determinism first; treat numeric parity as a bounded secondary check
+(symmetric/identity inputs, or store K pre-transposed). A real
+data-materialising `tl.transpose` is a *candidate* further primitive if
+true numeric validation is ever required — **recommend deferring it**; it
+is not needed for the performance-modeling goal. *Confirm you're content
+with structural-first verification.*
+
+### 10.11 GQA reuse = fold `G` into M, NOT a broadcast op **[finding, iteration 2]**
+The headline mechanism finding. The deferred-GQA limitation
+(`test_milestone_gqa_llama70b.py:137-142`) was blamed on a missing
+broadcast / the MemoryStore nbytes check. Iteration 2 shows the real
+cause is the baseline's head-packing reshape, and the fix is kernel
+restructuring (§4.3), which needs **no new primitive**. This demoted
+ADR-0061 from "the blocker" to "optional convenience" and reordered the
+phase plan (§7). Recorded explicitly because it changes the project's
+critical path: *GQA itself is not gated on any new feature; scale
+(ADR-0063) and efficiency (ADR-0062) are the features that matter.*
+
+---
+
+## 11. Risks
+
+| Risk | Likelihood | Mitigation |
+|---|---|---|
+| Numeric parity blocked by trans-as-reshape | med | structural-first verify (§10.10); pre-transpose K or defer `tl.transpose` |
+| Tree pairs not 1-hop on installed mesh | med | verify SFR neighbour table (§10.2); relabel PEs |
+| scratch_scope use-after-scope bug | med | two-arena discipline in shared helper; debug poison |
+| prefetch buffer recycled while in flight | med | live `LoadFuture` buffers in persistent arena (§10.5) |
+| GQA broadcast `.copy()` memory blow-up | low | bounded by TILE in tiled design |
+| Effort underestimated (6 phases) | med | P0–P1 alone deliver the core unblock; later phases independent |
+
+---
+
+## 12. Glossary & references
+
+- **GQA / G** — grouped-query attention; `G = H_q/H_kv` query heads share
+  one KV head. Llama3-70B: `G=8`.
+- **FlashAttention / FlashDecoding / Ring Attention** — see ADR-0060
+  lineage note.
+- **IPCQ** — inter-PE queues; `tl.send`/`tl.recv` over `N/S/E/W`
+  (ADR-0023/0025/0032).
+- **Greenlet `tl` model** — kernel interleaves with SimPy; ops emit
+  commands scheduled on PE engines (ADR-0020, ADR-0046).
+- **Persistent vs scoped arena** — ADR-0063 scratch split.
+
+Key source anchors: `tl_context.py` (tl API), `memory_store.py:67-73`
+(nbytes check), `data_executor.py` (Phase 2), `pe_dma.py`/`pe_gemm.py`/
+`pe_math.py` (latency), `_attention_mesh_kv.py` / `_attention_mesh_mlo.py`
+(baseline kernels), `milestone_gqa_llama70b.py` (eval bench),
+`tests/attention/test_milestone_gqa_llama70b.py` (current limits).
+
+ADRs: **0060** (this design), **0061** (broadcast), **0062** (async load),
+**0063** (scratch scope); related accepted **0014/0020/0023/0025/0042/
+0045/0046/0052/0054**.

docs/adr/ADR-0001-mem-physaddr-layout.md

@@ -0,0 +1,362 @@
+# ADR-0001: 51-bit Physical Address Layout & Decoding Contract
+
+## Status
+
+Accepted (Revision 2 — 2026-04-27: concrete bit layout, rack_id removal,
+Tray->SIP / SIP->DIE renaming, PE/MCPU/IOCPU sub-unit tables.
+Supersedes ADR-0031.)
+
+## Date
+
+2026-04-27 (original: 2026-02-27)
+
+## Context
+
+KernBench requires a stable, parsable physical address scheme that:
+
+- can be decoded into routing domains (SIP / die / HBM / PE-resource / IOCPU)
+- remains topology-agnostic (no hardcoded counts)
+- supports swappable policy and DI-first components
+- covers multiple SIPs, AHBM dies, and IO chiplet dies in a unified space
+
+### History
+
+- Original ADR-0001 defined a 51-bit layout with `rack_id(4) + sip_id(4) +
+  sip_seg(5) + local_offset(38)`. `rack_id` was never used in practice.
+- ADR-0031 (stub) requested PE-resource range partition but was never
+  implemented.
+
+Revision 2 removes `rack_id`, renames `sip_seg -> die_id`, and provides
+concrete sub-unit tables for PE, MCPU, CUBE_SRAM, and IOCPU resources.
+ADR-0031 is superseded.
+
+## Decision
+
+We define a **PhysAddr value object** and an **address decoding contract**
+that converts an integer address into routing domains.
+
+### D1. PhysAddr is an immutable value object
+
+- PhysAddr is immutable and comparable as a pure value.
+- Any allocator returns a **fully specified PhysAddr** (not partial metadata).
+- No global state may be required to interpret a PhysAddr.
+
+### D2. 51-bit Physical Address Layout
+
+A 51-bit physical address is adopted.
+
+#### 2.1 Top-Level Address Map
+
+```text
+[50:47] sip_id        (4)     -- 16 SIPs
+[46:42] die_id        (5)     -- 32 dies per SIP
+[41: 0] local_offset  (42)    -- 4 TB per die
+```
+
+```text
+50      47 46      42 41                      0
++---------+----------+-------------------------+
+| sip_id  | die_id   |      local_offset       |
++---------+----------+-------------------------+
+```
+
+#### 2.2 die_id Allocation
+
+| die_id | Meaning |
+|--------|---------|
+| 0..15  | AHBM dies |
+| 16..20 | IOCHIPLET dies |
+| 21..31 | Reserved |
+
+#### 2.3 AHBM Die Layout
+
+Only lower 256 GB of the 4 TB die-local window is assigned.
+
+```text
+[41:38] MBZ            (4)
+[37]    addr_space      (1)    -- 0 = local resource, 1 = HBM memory
+[36: 0] sub-address    (37)
+```
+
+| addr_space | Meaning |
+|------------|---------|
+| 0 | Local resource |
+| 1 | HBM memory |
+
+##### 2.3.1 HBM Window (addr_space = 1)
+
+```text
+[36:0] hbm_offset     (37)    -- 128 GB decode window
+```
+
+The architectural decode window is fixed at 128 GB. Implemented capacity
+may be smaller depending on SKU/topology (see D4).
+
+##### 2.3.2 Resource Window (addr_space = 0)
+
+```text
+[36:34] resource_kind  (3)
+[33: 0] kind_local    (34)    -- 16 GB per kind
+```
+
+| resource_kind | Meaning |
+|---------------|---------|
+| 000 | PE_LOCAL |
+| 001 | MCPU_LOCAL |
+| 010 | CUBE_SRAM |
+| 011..111 | Reserved |
+
+Each kind gets a 16 GB decode region.
+
+##### 2.3.3 PE_LOCAL (resource_kind = 000)
+
+```text
+[33]    MBZ            (1)
+[32:29] pe_id          (4)     -- 0..15
+[28:25] pe_sub_unit    (4)
+[24: 0] sub_offset    (25)    -- 32 MB per slot
+```
+
+16 PEs x 16 sub-unit slots x 32 MB = 8 GB active decode.
+
+| pe_sub_unit | Name | Budget |
+|-------------|------|--------|
+| 0 | PE_CPU_DTCM | 8 KB |
+| 1 | MATH_ENGINE_DTCM | 8 KB |
+| 2 | IPCQ | 256 KB |
+| 3 | PE_CPU_SFR | 16 KB |
+| 4 | MATH_ENGINE_SFR | 16 KB |
+| 5 | DMA_ENGINE_SFR | 192 KB |
+| 6 | PE_TCM | 2 MB |
+| 7..15 | Reserved | -- |
+
+##### 2.3.4 MCPU_LOCAL (resource_kind = 001)
+
+```text
+[33:30] MBZ            (4)
+[29:25] mcpu_sub_unit  (5)
+[24: 0] sub_offset    (25)    -- 32 MB per slot
+```
+
+1 GB active decode.
+
+| mcpu_sub_unit | Name | Budget |
+|---------------|------|--------|
+| 0 | MCPU_ITCM | 512 KB |
+| 1 | MCPU_DTCM | 512 KB |
+| 2 | IPCQ | 256 KB |
+| 3 | MCPU_SFR | 8 KB |
+| 4 | MCPU_DMA_SFR | 16 KB |
+| 5 | MCPU_SRAM | 10 MB |
+| 6..31 | Reserved | -- |
+
+##### 2.3.5 CUBE_SRAM (resource_kind = 010)
+
+```text
+[33:25] MBZ            (9)
+[24: 0] sram_offset   (25)    -- flat 32 MB
+```
+
+#### 2.4 IOCHIPLET Die Layout
+
+Only lower 1 TB of the 4 TB die-local window is assigned.
+
+```text
+[41:40] MBZ            (2)
+[39: 0] chiplet_offset (40)   -- 1 TB
+```
+
+Region split by address range:
+
+| Range | Meaning | Decode condition |
+|-------|---------|------------------|
+| [0, 2 GB) | IOCPU resource | chiplet_offset < 0x8000_0000 |
+| [2 GB, 1 TB) | UAL | chiplet_offset >= 0x8000_0000 |
+
+##### 2.4.1 IOCPU Region
+
+```text
+[30:27] iocpu_sub_unit (4)
+[26: 0] sub_offset    (27)    -- 128 MB per slot
+```
+
+16 x 128 MB slots. 2 GB active decode.
+
+| iocpu_sub_unit | Name | Budget |
+|----------------|------|--------|
+| 0 | IOCPU_ITCM | 512 KB |
+| 1 | IOCPU_DTCM | 512 KB |
+| 2 | IPCQ | 2 MB |
+| 3 | IOCPU_SFR | 8 KB |
+| 4 | IO_DMA_SFR | 16 KB |
+| 5 | IO_SRAM | 64 MB |
+| 6..15 | Reserved | -- |
+
+##### 2.4.2 UAL Region
+
+Sub-layout TBD (separate ADR).
+
+#### 2.5 Addressing Rules
+
+1. MBZ bits must be zero. An address with non-zero MBZ bits is
+   **architecturally invalid**. Implementation may raise a decode fault
+   or return an error -- behavior is not prescribed by this ADR.
+2. Fixed slot sizes are chosen for simple hardware decode; actual
+   implemented capacity may be smaller than the slot.
+3. Access beyond a sub-unit's implemented budget within a slot is
+   **architecturally invalid** (same policy as MBZ).
+
+### D3. Bitfield decoding is deterministic
+
+Given an integer address, field extraction (`sip_id`, `die_id`, `kind`,
+`sub_unit`, `offset`) is purely positional. No runtime state is required.
+Decoding deterministically maps an integer address to destination domains:
+`sip_id`, `die_id`, target kind (HBM / PE_LOCAL / MCPU_LOCAL / CUBE_SRAM /
+IOCPU / UAL).
+
+### D4. Capacity validation may depend on topology config
+
+Whether a decoded address falls within **implemented capacity** (e.g.,
+HBM 96 GB on a specific SKU) is checked against topology parameters
+provided via DI/config. Decode itself (D3) never consults topology --
+only validation does. These parameters must live in the topology/config
+layer, not in node implementations.
+
+### D5. Routing consumes decoded domains, not raw bits
+
+Routing policy uses decoded domains:
+
+- `src` location (sip / die / pe or node_id)
+- `dst` domains derived from PhysAddr decoding
+- `size_bytes` for size-aware link latency
+
+Routing must not inspect raw bit-fields directly except inside the
+decoding module.
+
+## Alternatives Considered
+
+1. **Keep `rack_id` (4 bits)**: Rejected -- never used in practice,
+   consumes 4 bits that enable die-local expansion to 42 bits
+   (IOCHIPLET 1 TB).
+
+2. **Uniform 256 GB per die**: Rejected -- IOCHIPLET UAL requires ~1 TB.
+   Freed rack_id bits enable 42-bit local_offset.
+
+3. **Variable-width die windows (AHBM 256 GB, CHIPLET 1 TB via multi-seg
+   spanning)**: Rejected -- complicates D3 (deterministic decoding).
+   Uniform 4 TB window with MBZ padding is simpler.
+
+4. **Use raw integers everywhere, decode ad-hoc in routing**: Rejected --
+   leads to duplicated logic, inconsistent routing, and hidden
+   assumptions.
+
+5. **Hardcode topology sizes (SIP/CUBE/PE counts) into decoding**:
+   Rejected -- violates SPEC R3 and breaks swappability.
+
+6. **Put decoding inside memory controllers or routers**: Rejected --
+   leaks policy into components, violates SPEC R4 / D5.
+
+## Consequences
+
+### Positive
+
+- Simple hierarchical decoder: SIP -> die -> kind -> sub-unit.
+- Clean separation of memory (HBM) vs local resource (PE/MCPU/SRAM/IOCPU).
+- Deterministic routing domains enable clear test invariants (SPEC R1, R5).
+- Expandable: 11 reserved die_id slots, reserved resource_kind / sub-unit
+  slots, reserved MBZ bits.
+- DI-first: decoder can be swapped without changing components (SPEC R4).
+
+### Tradeoffs
+
+- Sparse address holes due to power-of-2 slot alignment.
+- Large reserved/MBZ regions (intentional for future extension).
+- Requires explicit configuration for topology-derived sizes (D4).
+- Introduces a single "blessed" decoding module that must remain stable
+  and well-tested.
+
+## Supersedes
+
+- **ADR-0031 (PhysAddr PE-Resource Extension)**: stub status. The
+  PE_LOCAL / MCPU_LOCAL / CUBE_SRAM sub-unit tables in D2.3.3-D2.3.5
+  fulfill ADR-0031's stated goals.
+
+## Implementation Notes (Non-normative)
+
+- Recommended module: `src/kernbench/policy/address/phyaddr.py`
+- Tests should cover: encode/decode round-trip per kind, MBZ enforcement,
+  die_id dispatch (AHBM / IOCHIPLET / reserved), sub-unit boundary
+  values, backward compatibility of factory APIs.
+- Factory methods: `hbm_addr`, `pe_hbm_addr`, `pe_tcm_addr`,
+  `cube_sram_addr` retain signatures (minus `rack_id`); `cube_id`
+  parameter renamed to `die_id`.
+- New factories: `pe_resource_addr`, `mcpu_resource_addr`,
+  `iocpu_resource_addr`, `ual_addr`.
+
+## Appendix A. Address Examples
+
+### A.1 AHBM HBM access
+
+sip=2, die=5, HBM offset=0x1000
+
+```text
+sip_id     = 2       -> [50:47] = 0b0010
+die_id     = 5       -> [46:42] = 0b00101
+addr_space = 1       -> [37]    = 1 (HBM)
+hbm_offset = 0x1000  -> [36:0]
+
+51-bit addr = (2 << 47) | (5 << 42) | (1 << 37) | 0x1000
+```
+
+### A.2 AHBM PE_LOCAL -- PE3 PE_TCM, offset=0x400
+
+```text
+sip_id        = 0  -> [50:47] = 0
+die_id        = 0  -> [46:42] = 0
+addr_space    = 0  -> [37]    = 0
+resource_kind = 0  -> [36:34] = 000 (PE_LOCAL)
+pe_id         = 3  -> [32:29] = 0011
+pe_sub_unit   = 6  -> [28:25] = 0110 (PE_TCM)
+sub_offset    = 0x400 -> [24:0]
+
+local_offset = (0 << 34) | (3 << 29) | (6 << 25) | 0x400
+```
+
+### A.3 AHBM MCPU_LOCAL -- MCPU_SRAM, offset=0x0
+
+```text
+sip_id        = 1  -> [50:47] = 0001
+die_id        = 3  -> [46:42] = 00011
+addr_space    = 0  -> [37]    = 0
+resource_kind = 1  -> [36:34] = 001 (MCPU_LOCAL)
+mcpu_sub_unit = 5  -> [29:25] = 00101 (MCPU_SRAM)
+sub_offset    = 0  -> [24:0]  = 0
+
+local_offset = (1 << 34) | (5 << 25)
+```
+
+### A.4 IOCHIPLET -- IOCPU IPCQ, offset=0x20000
+
+```text
+sip_id         = 1   -> [50:47] = 0001
+die_id         = 17  -> [46:42] = 10001 (IOCHIPLET[1])
+iocpu_sub_unit = 2   -> [30:27] = 0010 (IPCQ)
+sub_offset     = 0x20000 -> [26:0]
+
+chiplet_offset = (2 << 27) | 0x20000
+                 (< 0x8000_0000 -> IOCPU region)
+```
+
+### A.5 IOCHIPLET -- UAL region, offset=4 GB
+
+```text
+sip_id         = 0   -> [50:47] = 0
+die_id         = 16  -> [46:42] = 10000 (IOCHIPLET[0])
+chiplet_offset = 0x1_0000_0000 (4 GB >= 2 GB -> UAL region)
+```
+
+## Links
+
+- SPEC.md: R1 (routing), R3 (configurable topology), R4 (DI-first),
+  R5 (multi-domain comm)
+- ADR-0031: Superseded

docs/adr/ADR-0001-physaddr-layout.md

@@ -1,108 +0,0 @@
-# ADR-0001: PhysAddr Layout & Address Decoding Contract
-
-## Status
-
-Accepted
-
-## Date
-
-2026-02-27
-
-## Context
-
-KernBench Graph Latency Simulator must route requests deterministically and compute end-to-end latency strictly by graph traversal.
-To model local vs remote traffic (same/different SIP, same/different CUBE, optional PE-group), requests need a stable, parsable address/location scheme that:
-
-- can be decoded into routing domains (SIP/CUBE/HBM/PE-resource, etc.)
-- remains topology-agnostic (no hardcoded counts)
-- supports swappable policy and DI-first components without leaking topology assumptions into node implementations
-
-## Decision
-
-We define a **PhysAddr value object** and an **address decoding contract** that converts an integer address into routing domains.
-
-### D1. PhysAddr is an immutable value object
-
-- PhysAddr is immutable and comparable as a pure value.
-- Any allocator returns a **fully specified PhysAddr** (not partial metadata).
-- No global state may be required to interpret a PhysAddr.
-
-### D2. PhysAddr fields (logical contract)
-
-PhysAddr must be able to represent at least:
-
-- `rack_id` (optional but reserved for scale-out)
-- `sip_id`  (device / SIP domain)
-- `sip_seg` (SIP-level segment/window selection, e.g., cube window)
-- `local_offset` (offset within the chosen segment/window)
-
-Decoded/derived fields may include (optional):
-
-- `cube_id`
-- `kind` (e.g., HBM vs PE-resource vs raw)
-- `unit_type` / `pe_id` (if PE-level addressing is modeled)
-
-**Important:** The exact bit allocation may evolve, but the *semantic fields above* must remain decodable without hidden assumptions.
-
-### D3. Decoding is deterministic and policy-compatible
-
-- Decoding must deterministically map an integer address to:
-  - destination SIP domain (`sip_id`)
-  - destination sub-domain (`cube_id` if applicable)
-  - destination target kind (HBM/PE-resource/other)
-- Decoding must not depend on runtime topology sizes; it may depend on **explicit topology parameters** provided through configuration (e.g., segment size, slice size), and those parameters must live in the topology/config layer (not in random components).
-
-### D4. Topology-derived constants live in the topology layer
-
-Constants such as segment sizes (e.g., HBM slice size / window size) are derived from topology configuration (YAML/JSON/dict) and are provided to the decoder via DI/config.
-They must not be hardcoded in node implementations.
-
-### D5. Routing consumes decoded domains, not raw bits
-
-Routing policy uses decoded domains:
-
-- `src` location (sip/cube/pe or node_id)
-- `dst` domains derived from PhysAddr decoding
-- `size_bytes` for size-aware link latency
-Routing must not inspect raw bit-fields directly except inside the decoding module.
-
-## Alternatives Considered
-
-1) **Use raw integers everywhere, decode ad-hoc in routing**
-
-- Rejected: leads to duplicated logic, inconsistent routing, and hidden assumptions embedded in multiple components.
-
-1) **Hardcode topology sizes (SIP/CUBE/PE counts) into decoding**
-
-- Rejected: violates SPEC (R3) and breaks swappability and configuration-driven topologies.
-
-1) **Put decoding inside memory controllers or routers**
-
-- Rejected: leaks policy into components and undermines DI-first, swappable implementations (SPEC R4).
-
-## Consequences
-
-### Positive
-
-- Deterministic routing domains enable clear test invariants for local vs remote paths (SPEC R1, R5).
-- Keeps topology variability (SPEC R3) while preserving consistent semantics.
-- DI-first: decoder can be swapped or extended without changing components or tests (SPEC R4).
-
-### Tradeoffs / Costs
-
-- Requires explicit configuration for any topology-derived sizes.
-- Introduces a single “blessed” decoding module that must remain stable and well-tested.
-
-## Implementation Notes (Non-normative)
-
-- Recommended module boundary:
-  - `src/kernbench/policy/address/phyaddr.py`
-
-- Tests should cover:
-  - deterministic decoding
-  - local vs remote classification from decoded fields
-  - invariants: “allocator returns full PhysAddr”, “decoding requires no global state”
-
-## Links
-
-- SPEC.md: R1 (routing), R3 (configurable topology), R4 (DI-first), R5 (multi-domain comm)

docs/adr/ADR-0002-lat-routing-distance.md

@@ -35,7 +35,7 @@ shortcuts that obscure control paths.
 
 ### D3. Bypass is explicit and graph-represented
 - All paths must be explicitly represented in the graph and subject to latency accumulation.
-- Example: PE_DMA connects to the NOC router mesh (ADR-0019). All destinations
+- Example: PE_DMA connects to the NOC router mesh (ADR-0017 D7). All destinations
   (HBM, shared SRAM, inter-cube UCIe) are reached via explicit mesh hops.
   Local HBM access has minimal hops (switching overhead only); remote access
   traverses additional routers.

docs/adr/ADR-0003-dev-target-system-hierarchy.md

@@ -35,11 +35,13 @@ We model the system hierarchy explicitly:
 
 - A CUBE contains:
   - HBM + memory controller (HBM_CTRL)
-  - NOC router mesh: 2D grid of explicit routers (from cube_mesh.yaml) with XY routing;
-    carries all intra-cube traffic including HBM data, inter-cube (UCIe),
-    command (M_CPU↔PE_CPU), and shared SRAM access.
-    HBM_CTRL is attached to PE routers (local HBM = 0 hop).
-    See ADR-0017 and ADR-0019 for full architecture.
+  - NOC (on-die fabric): carries all intra-cube traffic including HBM data,
+    inter-cube (UCIe), command (M_CPU↔PE_CPU), and shared SRAM access.
+    Must provide: full-BW PE↔local HBM path, PE↔SRAM connectivity,
+    PE↔UCIe connectivity, M_CPU↔PE command path.
+    NOC topology is an implementation choice (e.g., 2D mesh, ring, crossbar);
+    current implementation uses a 2D mesh with XY routing (see ADR-0017).
+    HBM_CTRL is attached to each PE's local NOC port (local HBM = minimal hop).
   - Shared SRAM: cube-level shared memory accessible by all PEs via NOC
   - management/control CPU (M_CPU) coordinating PE command distribution and completion aggregation
   - multiple PEs

docs/adr/ADR-0004-mem-memory-semantics-local-hbm.md

@@ -15,7 +15,7 @@ Each PE has a notion of “local HBM” that must guarantee full HBM bandwidth,
 
 - Each PE is assigned a logically defined “local HBM” region.
 - Local HBM corresponds to the pseudo-channel subset directly attached to that PE’s
-  router in the NOC mesh (ADR-0019).
+  router in the NOC mesh (ADR-0017 D4).
 - The path is: PE_DMA → local router → HBM_CTRL (switching overhead only, 0 mesh hops).
 - The mapping (HBM pseudo-channels → PE local regions) is derived from topology configuration.
 
@@ -33,12 +33,17 @@ Each PE has a notion of “local HBM” that must guarantee full HBM bandwidth,
 - This guarantee is modeled by:
   - a dedicated logical path and/or service model that enforces HBM BW at the PE-local-HBM interaction point,
   - while still incurring non-zero latency along explicitly modeled components.
+- HBM CTRL internal modeling (PC striping, cut-through, scheduling fidelity)
+  is consolidated in ADR-0033 (Latency Model: Assumptions and Known
+  Simplifications). The aggregate BW guarantee here remains the contract;
+  ADR-0033 documents how the per-PC model realizes it and which scheduler
+  effects are intentionally simplified.
 
 ### D3. Remote PE HBM semantics (intra-cube)
 
-- A PE that accesses another PE's local HBM traverses the router mesh:
-  - PE_DMA → local router → (mesh hops) → target PE's router → HBM_CTRL
-- Router mesh bandwidth and hop count may limit remote HBM access relative to local access.
+- A PE that accesses another PE's local HBM traverses the NOC:
+  - PE_DMA → NOC → (fabric hops) → target PE's NOC port → HBM_CTRL
+- NOC bandwidth and hop count may limit remote HBM access relative to local access.
 
 ### D4. Non-local HBM semantics (inter-cube / inter-SIP)
 

docs/adr/ADR-0005-dev-diagram-views-distance-layout.md

@@ -20,7 +20,9 @@ Diagrams must reflect this distance by default.
 
 ---
 
-## Global Defaults
+## Decision
+
+### D1. Global Defaults
 
 - All diagrams MUST be **distance-aware by default**.
 - All diagrams MUST render **representative views** of the architecture.
@@ -31,7 +33,7 @@ Diagrams must reflect this distance by default.
 
 ---
 
-## Representative Rendering Rule
+### D2. Representative Rendering Rule
 
 - All CUBEs share the same internal structure.
 - All PEs share the same internal structure.
@@ -47,9 +49,9 @@ unless explicitly requested.
 
 ---
 
-## Diagram Views
+### D3. Diagram Views
 
-### View A — SIP-Level Diagram
+#### View A — SIP-Level Diagram
 
 **Purpose**
 Explain system-scale structure and connectivity.
@@ -75,7 +77,7 @@ Explain system-scale structure and connectivity.
 
 ---
 
-### View B — CUBE-Level Diagram
+#### View B — CUBE-Level Diagram
 
 **Purpose**
 Explain cube-internal structure and data/control flow.
@@ -106,7 +108,7 @@ Explain cube-internal structure and data/control flow.
 
 ---
 
-### View C — PE-Level Diagram
+#### View C — PE-Level Diagram
 
 **Purpose**
 Explain internal PE behavior and execution structure.
@@ -128,14 +130,14 @@ Explain internal PE behavior and execution structure.
 
 ---
 
-## Distance-Aware Layout (Default)
+### D4. Distance-Aware Layout (Default)
 
-### Distance definition
+#### Distance definition
 
 - Distance is defined as **accumulated latency**, consistent with ADR-0002.
 - Distance is computed from a single anchor node.
 
-### Default anchor selection
+#### Default anchor selection
 
 - SIP view: IO chiplet (or Host CPU if present)
 - CUBE view: a representative PE
@@ -143,7 +145,7 @@ Explain internal PE behavior and execution structure.
 
 Anchors are **implicit defaults** and MUST NOT be required to be specified.
 
-### Layout rules
+#### Layout rules
 
 - Diagrams MUST be laid out in layers based on distance buckets.
 - Layout direction MUST be consistent within a view type
@@ -156,7 +158,7 @@ without affecting distance semantics.
 
 ---
 
-## Generation Contract (for Tools / Claude Code)
+### D5. Generation Contract (for Tools / Claude Code)
 
 When generating diagrams:
 

docs/adr/ADR-0006-dev-topology-compilation-distance-diagram.md

@@ -63,7 +63,7 @@ For each view (SIP / CUBE / PE):
 - CUBE-level projection MUST include:
   - Router mesh (from cube_mesh.yaml), HBM_CTRL, shared SRAM, M_CPU, UCIe ports,
     and PEs as opaque blocks.
-  - All paths (HBM, non-HBM, command) route through the same router mesh (ADR-0019).
+  - All paths (HBM, non-HBM, command) route through the same router mesh (ADR-0017).
 - Default anchors are implicit (ADR-0005) and MUST NOT require instance indices.
 
 ### D6. Output formats and determinism

docs/adr/ADR-0007-api-runtime-api-boundaries.md

@@ -42,21 +42,25 @@ The runtime API MUST NOT:
 
 ---
 
-### D2. Simulation engine executes and schedules requests
+### D2. Simulation engine wires components and tracks completion
 
 The simulation engine (sim_engine) MUST:
 
-- inject requests into the compiled topology graph,
+- wire components at initialization (create port stores + start wire
+  processes per the component port/wire framework — ADR-0015),
+- inject requests into the compiled topology graph at entry components
+  (e.g., PCIE_EP for memory operations, IO_CPU for kernel launch),
 - schedule and execute events using a discrete-event model,
-- manage correlation ids and completion tracking,
-- decompose operations into low-level requests when required
-  (e.g., MemoryWrite events).
+- manage correlation ids and completion tracking.
 
 The simulation engine MUST NOT:
 
 - define tensor semantics,
 - define kernel execution policies,
-- expose internal graph details to the runtime API.
+- expose internal graph details to the runtime API,
+- walk the topology path during request execution,
+- call component `run()` methods directly,
+- track per-hop latency or decompose fan-out (components own this).
 
 ---
 
@@ -87,3 +91,5 @@ component-level fan-out explicitly.
 - SPEC R4, R7, R8
 - ADR-0008 (Tensor deployment)
 - ADR-0009 (Kernel execution)
+- ADR-0015 (Component port/wire model and engine role)
+- ADR-0010 (CLI surface and execution semantics — runtime API consumer)

docs/adr/ADR-0008-api-tensor-deploy-and-allocation.md

@@ -94,7 +94,7 @@ The Phase 0 PA shard map remains a valid fast-path configuration.
 
 ## Links
 
-- ADR-0011 (PA-first)
+- ADR-0011 (Memory Addressing — PA / VA / LA)
 - ADR-0012 (Host↔IO_CPU schema)
 - ADR-0007 (runtime_api vs sim_engine boundaries)
 - ADR-0009 (Kernel execution)

docs/adr/ADR-0009-api-kernel-execution-messaging.md

@@ -0,0 +1,146 @@
+# ADR-0009: Kernel Execution Messaging and Completion Semantics
+
+## Status
+
+Accepted
+
+## Context
+
+Kernel execution is initiated by the host and proceeds through
+device control components:
+
+Host → IO_CPU → M_CPU → PE_CPU → schedulers → engines
+
+Completion propagates in reverse order.
+
+To keep benchmarks simple and topology-agnostic,
+kernel execution must be endpoint-driven with deterministic aggregation.
+
+---
+
+## Decision
+
+### D1. Kernel launch is an endpoint request
+
+A kernel launch is initiated by submitting a single KernelLaunch request
+to the IO_CPU endpoint.
+
+The runtime API MUST:
+
+- construct the kernel launch request,
+- submit it to IO_CPU,
+- await a single completion result.
+
+The runtime API MUST NOT orchestrate internal fan-out.
+
+---
+
+### D2. Tensor arguments are passed by metadata
+
+KernelLaunch requests MUST reference tensor arguments via:
+
+- host-owned tensor handles, or
+- resolved device address maps derived from those handles.
+
+Bulk tensor data MUST NOT be embedded in kernel launch messages.
+
+---
+
+### D3. Fan-out and aggregation are component responsibilities
+
+- IO_CPU fans out work to M_CPUs.
+- M_CPU fans out work to PE_CPUs.
+- PE_CPU manages kernel execution and engine dispatch.
+
+Completion semantics:
+
+- M_CPU completes when all targeted PEs complete or a failure policy triggers.
+- IO_CPU completes when all targeted CUBEs complete or a failure policy triggers.
+
+---
+
+### D4. Completion and failure propagation
+
+- All messages MUST carry correlation identifiers.
+- Completion and failure MUST propagate deterministically to the host.
+- The simulation engine provides futures/handles to observe completion.
+
+---
+
+### D5. Launch timing is endpoint-synchronized
+
+All PEs targeted by a single kernel launch MUST begin executing the kernel
+body at the same simulated time, regardless of their dispatch path length
+from the launch entry point.
+
+Rationale. The dispatch tree Host → IO_CPU → M_CPU → PE_CPU has variable
+latency at every level. PEs near their M_CPU receive the launch earlier
+than PEs farther away; cubes near an IO_CPU receive it earlier than cubes
+farther away. Without synchronization, each PE's kernel begins at a
+different `env.now`, making per-PE metrics such as `pe_exec_ns` a function
+of dispatch-path geometry rather than of the kernel's behavior —
+producing measurement artifacts in benchmarks that time kernel-internal
+waits (for example `tl.recv` on cross-cube or cross-SIP hops).
+
+Mechanism.
+
+- `KernelLaunchMsg` carries an optional `target_start_ns: float | None`.
+- **IO_CPU** is the canonical stamper. On fan-out to M_CPUs, it
+  computes `target_start_ns = env.now + max_latency` where
+  `max_latency` is the maximum, over every target (sip, cube, pe)
+  tuple, of the **two-leg dispatch chain**:
+
+  ```
+  max_latency(sip, cube, pe) =
+      compute_path_latency_ns(find_node_path(io_cpu, m_cpu(sip, cube)))
+    + compute_path_latency_ns(find_node_path(m_cpu(sip, cube), pe_cpu))
+    - io_cpu.overhead_ns
+    - m_cpu.overhead_ns
+  ```
+
+  This models the actual dispatch as **two sequential Transactions**
+  (IO_CPU → M_CPU, then M_CPU → PE_CPU). Each leg's
+  `compute_path_latency_ns` adds its endpoints' `overhead_ns`;
+  `io_cpu.overhead_ns` is subtracted because IO_CPU has already
+  paid it before this method runs, and `m_cpu.overhead_ns` is
+  subtracted once because it appears as endpoint of leg1 *and*
+  start of leg2 but is paid only once at run time. A single
+  `find_node_path(io_cpu, pe_cpu)` walk is **not** equivalent —
+  it can pick a graph path that bypasses M_CPU and silently
+  under-shoots the prediction for far cubes, breaking the D5
+  invariant.
+
+  The fanned-out sub-Transactions carry **`nbytes = 0`** for
+  `KernelLaunchMsg` (control message only). Without this,
+  large kernel-launch payloads would occupy fabric BW on the
+  shared first hop and serialize the per-cube dispatch, pushing
+  far M_CPUs past `target_start_ns` and re-introducing the
+  late-arrival violation.
+- **M_CPU** passes an already-stamped `target_start_ns` through
+  unchanged. Only when the value is absent (e.g. a direct
+  launch-to-M_CPU unit test) does M_CPU compute a per-cube barrier
+  `env.now + max(local command-path latency)`.
+- **PE_CPU** yields `env.timeout(target_start_ns - env.now)` at the top
+  of `_execute_kernel`, before recording `pe_exec_start` and invoking
+  the kernel body.
+- When `target_start_ns is None`, PE_CPU falls through to the legacy
+  unsynchronized behavior — preserving backward compatibility.
+
+IO_CPU-level stamping guarantees every PE across every targeted cube
+uses the same barrier sim-time, eliminating both the within-cube
+dispatch-offset artifact *and* the cross-cube offset artifact in
+multi-cube launches. Models a real-hardware timed-broadcast launch
+(latency-equalized dispatch tree).
+
+The synchronization is internal to the engine / IO_CPU / M_CPU / PE_CPU
+control plane — runtime API and application kernels are unchanged.
+
+---
+
+## Links
+
+- SPEC R1, R2, R7, R8
+- ADR-0007 (Runtime API boundaries)
+- ADR-0008 (Tensor deployment)
+- ADR-0013 (Verification strategy — V2 fan-out tests)
+- ADR-0015 D4 (concrete fabric path for kernel launch)

docs/adr/ADR-0009-kernel-execution-messaging.md

@@ -1,74 +0,0 @@
-# ADR-0009: Kernel Execution Messaging and Completion Semantics
-
-## Status
-
-Accepted
-
-## Context
-
-Kernel execution is initiated by the host and proceeds through
-device control components:
-
-Host → IO_CPU → M_CPU → PE_CPU → schedulers → engines
-
-Completion propagates in reverse order.
-
-To keep benchmarks simple and topology-agnostic,
-kernel execution must be endpoint-driven with deterministic aggregation.
-
----
-
-## Decision
-
-### D1. Kernel launch is an endpoint request
-
-A kernel launch is initiated by submitting a single KernelLaunch request
-to the IO_CPU endpoint.
-
-The runtime API MUST:
-
-- construct the kernel launch request,
-- submit it to IO_CPU,
-- await a single completion result.
-
-The runtime API MUST NOT orchestrate internal fan-out.
-
----
-
-### D2. Tensor arguments are passed by metadata
-
-KernelLaunch requests MUST reference tensor arguments via:
-
-- host-owned tensor handles, or
-- resolved device address maps derived from those handles.
-
-Bulk tensor data MUST NOT be embedded in kernel launch messages.
-
----
-
-### D3. Fan-out and aggregation are component responsibilities
-
-- IO_CPU fans out work to M_CPUs.
-- M_CPU fans out work to PE_CPUs.
-- PE_CPU manages kernel execution and engine dispatch.
-
-Completion semantics:
-
-- M_CPU completes when all targeted PEs complete or a failure policy triggers.
-- IO_CPU completes when all targeted CUBEs complete or a failure policy triggers.
-
----
-
-### D4. Completion and failure propagation
-
-- All messages MUST carry correlation identifiers.
-- Completion and failure MUST propagate deterministically to the host.
-- The simulation engine provides futures/handles to observe completion.
-
----
-
-## Links
-
-- SPEC R1, R2, R7, R8
-- ADR-0007 (Runtime API boundaries)
-- ADR-0008 (Tensor deployment)

docs/adr/ADR-0010-api-cli-surface-and-semantics.md

@@ -0,0 +1,152 @@
+# ADR-0010: Command Line Interface and Execution Semantics
+
+## Status
+
+Accepted
+
+## Context
+
+The `kernbench` CLI is the user-facing entry point of the simulator. It
+exposes four subcommands:
+
+- `run` — execute a benchmark against a topology.
+- `list` — enumerate registered benches.
+- `probe` — diagnostic utility for latency / BW measurement.
+- `web` — interactive topology viewer.
+
+Device enumeration is centralized in the CLI; neither the runtime API
+nor the simulation engine enumerates devices. Benchmarks remain
+single-device by design and accept a device identifier as input.
+
+## Decision
+
+### D1. Benchmark contract — single-device by design
+
+- A benchmark MUST define behavior for a single device only.
+- A benchmark MUST accept a device identifier as input.
+- Benchmarks MUST NOT enumerate or loop over multiple devices.
+
+Multi-device execution is the CLI's concern (D3), not the benchmark's.
+
+### D2. `kernbench run` — benchmark execution
+
+Required arguments:
+
+- `--topology <path>`: topology YAML file path. Loaded via
+  `resolve_topology()`.
+- `--bench <identifier>`: benchmark identifier. Resolved via
+  `kernbench.benches.registry.resolve()`, which accepts either the
+  registered kebab-case name (e.g., `gemm-single-pe`) or a numeric
+  index from `kernbench list`.
+
+Optional arguments:
+
+- `--device <selector>` (default: `all`):
+  - `all` — run once per discovered SIP (see D3).
+  - `sip:<N>` — run only on SIP N.
+  - Parsed via `resolve_device()`.
+- `--verify-data` (default: off) — enable Phase 2 data verification
+  (see ADR-0020). When set, `engine_factory` constructs the engine
+  with `enable_data=True`. After the benchmark runs, a diagnostic
+  summary of recorded ops is printed.
+
+Each invocation runs the benchmark once within a single simulation
+instance.
+
+### D3. Multi-device execution is logically parallel
+
+When `--device all` (or omitted) and the topology has multiple SIPs:
+
+- Benchmark executions are submitted to a single simulation engine
+  instance.
+- Executions are logically parallel in simulation time.
+- Inter-device contention is naturally modeled (shared fabric
+  bandwidth, cross-SIP traffic, etc.).
+
+The CLI does NOT spawn multiple OS processes or independent
+simulation runs — parallelism is internal to one simulation instance.
+
+### D4. `kernbench list` — enumerate registered benches
+
+No arguments. Prints each registered bench's auto-assigned index,
+registered name, and one-line description.
+
+Benches register themselves via the `@bench(name=..., description=...)`
+decorator (`kernbench.benches.registry`). Every non-underscore module
+under `kernbench.benches/` MUST register at least one bench; a missing
+decorator raises `RuntimeError` at package import time.
+
+Indices are assigned alphabetically by name at import time. They are a
+CLI convenience (shorthand for `--bench`), not a stable API — a new
+bench inserted alphabetically will shift later indices.
+
+### D5. `kernbench probe` — latency / BW diagnostic utility
+
+Required argument:
+
+- `--topology <path>`: topology YAML file path.
+
+Optional argument:
+
+- `--case <name>` (default: `all`) — run a predefined traffic
+  pattern, or `all` to run every defined case.
+
+Probe runs each pattern through the simulation engine and reports
+per case:
+
+- End-to-end latency (ns).
+- Effective bandwidth (nbytes / total_ns).
+- Bottleneck bandwidth (min edge BW along the chosen path).
+- Utilization (effective / bottleneck).
+
+Probe additionally validates monotonicity invariants — for example
+that local-HBM access ≤ cross-PE-within-cube ≤ cross-cube ≤
+cross-SIP — and reports violations. Probe is a developer tool for
+verifying the latency / BW model; it is not a benchmark.
+
+### D6. `kernbench web` — topology viewer
+
+Optional arguments:
+
+- `--port <N>` (default: `8765`) — HTTP port.
+- `--no-open` — do not auto-open the browser.
+
+Launches a local HTTP server that renders the compiled topology in
+the browser. Distinct from the static `docs/diagrams/` artifacts:
+
+- `docs/diagrams/` files are derived at topology-compile time
+  (ADR-0006).
+- `kernbench web` is interactive — pan/zoom, hover for component
+  attributes, switch between SIP / CUBE / PE views.
+
+### D7. Runtime API and simulation engine remain device-scoped
+
+- Runtime API calls operate on one device per invocation.
+- The simulation engine schedules all requests deterministically.
+- Neither layer enumerates devices.
+
+This invariant keeps each layer testable in isolation; device
+enumeration and multi-device fan-out live only in the CLI's `run`
+command (D3).
+
+The `probe` implementation lives under `kernbench.probes` (separate
+from `kernbench.benches`), reflecting that probes are diagnostic
+utilities, not registered benches.
+
+## Consequences
+
+- Benchmark authors write single-device logic; multi-device behavior
+  emerges from the CLI dispatching across SIPs.
+- Adding a new subcommand (e.g., trace export, replay) does not
+  require benchmark or runtime-API changes — the CLI is the
+  extension point.
+- `probe` and `web` are diagnostic / visualization tools, not
+  benchmarks; they bypass the benchmark loader path.
+
+## Links
+
+- SPEC R7, R8, R9
+- ADR-0007 (Runtime API and Simulation Engine Boundaries)
+- ADR-0020 (Two-pass data execution — `--verify-data`)
+- ADR-0006 (Topology compilation and diagram generation —
+  background for `kernbench web`)

docs/adr/ADR-0010-cli-device-selection.md

@@ -1,62 +0,0 @@
-# ADR-0010: CLI Device Selection and Multi-Device Execution Semantics
-
-## Status
-
-Accepted
-
-## Context
-
-Benchmarks represent device-agnostic workloads that operate on a single device.
-Users may want to run a benchmark:
-
-- on a specific device, or
-- across all devices in the system.
-
-Device enumeration must not leak into benchmarks or runtime APIs.
-
----
-
-## Decision
-
-### D1. Benchmarks are single-device by design
-
-- A benchmark MUST define behavior for a single device only.
-- A benchmark MUST accept a device identifier as input.
-- Benchmarks MUST NOT enumerate or loop over multiple devices.
-
----
-
-### D2. CLI controls device selection
-
-The `kernbench run` command supports an optional `--device` argument:
-
-- If `--device <id>` is specified:
-  - the benchmark executes once for the specified device.
-
-- If `--device` is omitted:
-  - the benchmark executes once using all the SIPs discovered in the topology.
-
----
-
-### D3. Multi-device execution is logically parallel
-
-When running on multiple devices:
-
-- benchmark executions are submitted to a single simulation engine instance,
-- executions are logically parallel in simulation time,
-- inter-device contention is naturally modeled.
-
----
-
-### D4. Runtime API and simulation engine remain device-scoped
-
-- Runtime API calls operate on one device per invocation.
-- The simulation engine schedules all requests deterministically.
-- Neither layer enumerates devices.
-
----
-
-## Links
-
-- SPEC R7, R8
-- ADR-0007 (Runtime API boundaries)

docs/adr/ADR-0011-mem-memory-addressing-simplification.md

@@ -0,0 +1,521 @@
+# ADR-0011: Memory Addressing — PA / VA / LA Address Models
+
+## Status
+
+Accepted.
+
+- **VA model: currently implemented (default).**
+- PA model: implemented as PageFault fallback in PE_DMA.
+- LA model: proposed, not implemented.
+
+## Context
+
+KernBench's address model evolved through three design points, each
+addressing a limitation of the previous. This ADR documents all three
+in one place because future implementation work selects among them.
+
+### PA-only baseline
+
+Phase 0 of KernBench treated all device memory operations
+(MemoryRead/MemoryWrite) as raw physical-address transfers. No
+host-side virtual addressing, no MMU/IOMMU translation. Allocators
+returned PA mappings; DMA requests carried PA directly.
+
+This was sufficient for early correctness/latency work but
+insufficient for running standard Triton kernels that use
+`base_addr + offset` patterns on sharded tensors: each PE's shard
+has a different PA, but the kernel needs a single contiguous address
+space to compute offsets.
+
+### Why VA/MMU (current default)
+
+A realistic system uses host-side virtual addressing and an
+MMU/IOMMU-style translation path for DMA: the host allocates physical
+memory at PE level, maps it into a virtual address space, installs
+mappings, and DMA requests use virtual addresses that are translated
+to physical addresses.
+
+Adopting this model lets kernels use `base_addr + offset` over a
+contiguous VA range while the device-side MMU translates each access
+to the appropriate PA.
+
+### Why LA/BAAW (proposed)
+
+VA/MMU treats HBM as a single backing space. KernBench needs to
+explore architectures where HBM is composed of multiple pseudo
+channels in parallel:
+
+- CUBE's HBM has 32 or 64 pseudo channels.
+- In a PE-Local-HBM model, each PE is assigned N pseudo channels
+  (N = `hbm_pseudo_channels / pes_per_cube`).
+- Per-channel BW (e.g. 32 GB/s) determines aggregate PE BW
+  (N × per-channel).
+
+Two channel-mapping modes need to be modelable:
+
+- **1:1 mode** — one logical access → N per-channel requests.
+  Precise per-channel BW contention modelling.
+- **n:1 mode (default)** — one logical access → one aggregated
+  request. Channels are assumed to interleave; aggregated BW model.
+
+VA's `tl.load(va_ptr)` produces a single DMA request to a single
+target. Decomposing that into per-channel requests inside PE_DMA
+requires the address layer to be aware of channels. This is the
+role of the LA (Logical Address) abstraction with BAAW
+(Logical-to-Physical Mapping Unit).
+
+Core requirements driving the LA design:
+
+- PE_DMA → HBM_CTRL effective bandwidth semantics must be identical
+  in both modes (only request shape and resource model differ).
+- Kernel programming model is unchanged — physical channel
+  information is never exposed to kernel code.
+- Mode switch is a topology-level configuration.
+
+### Design space summary
+
+| Model | Status | Key idea |
+|-------|--------|----------|
+| PA | fallback (implemented) | Direct physical addressing, no translation |
+| VA | current default (implemented) | Per-tensor contiguous VA range; MMU translates per access |
+| LA | proposed | LA + BAAW resolves to (PA, channel); supports 1:1 and n:1 channel mapping modes |
+
+---
+
+## Decision
+
+This ADR defines three address models. At any given time the system
+operates in exactly one model. Selection is topology- / configuration-
+driven; coexistence within one simulation run is not required.
+
+---
+
+### Address Model: PA (Physical Address) — fallback
+
+#### D-PA1. PA-only semantics
+
+- All device memory accesses (MemoryRead/MemoryWrite) operate on
+  device physical addresses (PA) plus size.
+- PA-only mode remains functional via the PageFault fallback path in
+  PE_DMA: if a DMA src/dst address has no MMU mapping, PE_DMA treats
+  the value as a PA directly.
+
+#### D-PA2. Allocation produces PA mappings
+
+Device allocation selects PE-local memory regions and returns PA
+mappings sufficient to execute kernels and issue DMA requests.
+
+PA model is retained primarily for backward compatibility with PA-only
+tests and as the underlying physical layer that VA / LA models resolve
+into.
+
+---
+
+### Address Model: VA (Virtual Address with MMU) — current default
+
+#### D-VA1. Virtual Address Model
+
+- Each tensor gets a single contiguous VA range (`TensorHandle.va_base`).
+- `TensorShard` does NOT carry a `va` field — shard VA is derived as
+  `va_base + offset_bytes`.
+- Kernels receive `va_base` as their pointer argument (via
+  `TensorArg.va_base`).
+- `DmaReadCmd.src_addr` and `DmaWriteCmd.dst_addr` carry VA (not PA).
+
+#### D-VA2. PE_MMU Component
+
+- Hybrid design: SimPy component (inbox for `MmuMapMsg`) + utility
+  (synchronous `translate()` called by PE_DMA).
+- Page-aligned dict lookup for O(1) VA → PA translation.
+- `tlb_overhead_ns` configurable per-access latency.
+- PageFault fallback: if VA has no mapping, PE_DMA treats it as PA
+  directly (preserves PA model for backward compatibility).
+
+#### D-VA3. Mapping Installation
+
+- `MmuMapMsg` traverses the fabric: Host → PCIE_EP → IO_CPU (cube
+  fan-out) → M_CPU (PE fan-out) → NOC → PE_MMU. Latency is measured
+  end-to-end.
+- `MmuMapMsg.target_sips` controls SIP-level routing to prevent
+  cross-SIP mapping contamination for replicated tensors.
+- Mapping strategy based on `DPPolicy.cube`:
+  - **Replicate** (`cube="replicate"`): per-(sip, cube) local mapping
+    only. Each cube's PEs see only their local PA. No cross-cube
+    mapping installed.
+  - **Sharded** (`cube="column_wise"`, etc.): broadcast all shard
+    mappings to all target cubes. Enables cross-PE and cross-cube
+    DMA.
+
+#### D-VA4. Tensor Lifecycle
+
+- `del tensor` triggers automatic cleanup via `Tensor.__del__` +
+  `weakref` to `RuntimeContext`. Sends `MmuUnmapMsg` through fabric,
+  returns VA and PA space.
+- `with RuntimeContext(...) as ctx:` provides scope-based bulk cleanup.
+- `RuntimeContext._tensors` uses `weakref.ref` to avoid preventing GC.
+- `PEMemAllocator` uses free-list with coalescing (not bump allocator).
+- `VirtualAllocator` uses free-list with coalescing for VA space.
+
+#### D-VA5. Allocators
+
+- `VirtualAllocator`: device-wide VA space, page-aligned alloc/free
+  with coalescing.
+- `PEMemAllocator`: per-PE HBM/TCM, free-list based alloc/free with
+  coalescing.
+- Page size configurable via `topology.yaml` `pe_mmu` attrs
+  (default 4096).
+
+#### Consequences (VA model)
+
+- Triton kernels use `base_addr + offset` patterns naturally on
+  sharded tensors.
+- All latency remains explicit via graph traversal, including MMU
+  mapping installation and per-access TLB overhead.
+- PA-only mode retained as fallback (PageFault → treat as PA).
+- IPCQ and other fixed-address resources bypass MMU (use PA directly).
+
+---
+
+### Address Model: LA (Logical Address with BAAW) — proposed
+
+LA replaces VA when channel-level HBM modelling is required.
+Adopting this model removes the VA/MMU infrastructure (D-LA1 lists the
+removed artifacts). Coexistence with VA in the same run is not a goal.
+
+#### D-LA1. LA introduction — replaces VA infrastructure
+
+LA is the sole address space used by kernel code (`tl.load`,
+`tl.store`, `tl.composite`). Properties:
+
+- Can map a Tensor to a contiguous logical space (like VA).
+- Expresses `(logical buffer + offset)`.
+- Does NOT contain physical channel information directly.
+- Stays as an intermediate abstraction until physical resolution.
+
+LA address space:
+
+| Item | Value |
+|------|-------|
+| LA start | `0x1_0000_0000` (4 GB, preserves former VA start) |
+| LA space size | 64 GB per PE |
+| Alignment unit | segment (see D-LA3) |
+
+LA is PE-local: different PEs may use the same LA value; BAAW segment
+tables differ → they resolve to different PAs.
+
+VA infrastructure removed when LA is adopted:
+
+| Removed | Replacement |
+|---------|-------------|
+| `policy/address/va_allocator.py` (VirtualAllocator) | LA allocator (same free-list approach, renamed) |
+| `policy/address/pe_mmu.py` (PeMMU) | BAAW segment table (inside PE_DMA) |
+| `components/builtin/pe_mmu.py` (PeMmuComponent) | Removed — BAAW is internal PE_DMA logic, not a separate component |
+| `runtime_api/kernel.py`: `MmuMapMsg`, `MmuUnmapMsg` | `BaawSegmentInstallMsg` |
+| `runtime_api/context.py`: VA alloc + MMU install | LA alloc + BAAW segment install |
+| `runtime_api/tensor.py`: `va_base` | `la_base` |
+| `topology.yaml`: `pe_mmu` component entry | Removed |
+
+#### D-LA2. Mapping mode setting
+
+Topology-level (cube) configuration:
+
+```yaml
+cube:
+  memory_map:
+    hbm_mapping_mode: n_to_one    # one_to_one | n_to_one
+    hbm_pseudo_channels: 64       # total pseudo channel count
+    hbm_channels_per_pe: 8        # per-PE local channel count
+    hbm_channel_bw_gbs: 32.0      # per-channel bandwidth
+```
+
+Consumed by the graph compiler (topology builder) and BAAW
+initialisation.
+
+#### D-LA3. Segment and BAAW
+
+Segment partitions the LA space; each segment maps to a specific HBM
+channel or channel group. Created at tensor deploy time by the runtime
+allocator. BAAW resolves LA → physical request(s) using the segment
+table.
+
+```python
+@dataclass
+class BaawSegment:
+    la_base: int          # segment start LA
+    la_size: int          # segment size (bytes)
+    mode: str             # "one_to_one" | "n_to_one"
+    # 1:1 mode fields
+    channel_count: int    # channels assigned to this segment (e.g. 8)
+    pa_bases: list[int]   # per-channel PA bases (len = channel_count)
+    channel_ids: list[int]   # per-channel logical IDs (e.g. [0..7])
+    channel_size: int     # per-channel size (la_size // channel_count)
+    # n:1 mode fields
+    agg_pa_base: int      # aggregated PA base
+    agg_node_id: str      # aggregated router node_id
+```
+
+Segment lifecycle:
+
+1. **Allocate** (tensor deploy): RuntimeContext allocates LA from LA
+   allocator. PEMemAllocator allocates per-channel PA (1:1) or
+   aggregated PA (n:1). `BaawSegmentInstallMsg` registers the segment
+   with PE_DMA.
+2. **Use** (kernel run): kernel `tl.load(la_ptr)` → `DmaReadCmd
+   (src_addr=LA)`. PE_DMA's BAAW front-end looks up the segment and
+   converts to PA(s).
+3. **Free** (tensor free): segment removed from table; LA and PA
+   returned.
+
+#### D-LA4. BAAW resolution logic
+
+BAAW is a front-end stage inside PE_DMA, not a separate SimPy
+component. Synchronous address-resolution logic executed at the start
+of PE_DMA's `handle_command()`.
+
+Input: `(LA, nbytes)`. Output:
+
+- **1:1 mode**: `list[PhysicalRequest]` — one per channel.
+- **n:1 mode**: single `PhysicalRequest`.
+
+```python
+@dataclass
+class PhysicalRequest:
+    pa: int           # 51-bit Physical Address
+    nbytes: int       # transfer size for this request
+    dst_node: str     # target node_id (channel router or aggregated router)
+
+
+def resolve(self, la: int, nbytes: int) -> list[PhysicalRequest]:
+    seg = self._find_segment(la)  # la_base <= la < la_base + la_size
+    offset = la - seg.la_base
+
+    if seg.mode == "n_to_one":
+        pa = seg.agg_pa_base + offset
+        return [PhysicalRequest(pa=pa, nbytes=nbytes, dst_node=seg.agg_node_id)]
+
+    # one_to_one
+    requests = []
+    per_ch_size = seg.channel_size
+    for i, (pa_base, ch_id) in enumerate(zip(seg.pa_bases, seg.channel_ids)):
+        ch_offset = offset % per_ch_size
+        ch_nbytes = nbytes // seg.channel_count
+        pa = pa_base + ch_offset
+        dst_node = f"{self._pe_prefix}.ch_r{ch_id}"
+        requests.append(PhysicalRequest(pa=pa, nbytes=ch_nbytes, dst_node=dst_node))
+    return requests
+```
+
+BAAW responsibilities:
+
+- Convert logical access → physical request units.
+- Apply mode-dependent fan-out (1:1) or pass-through (n:1).
+- Compute PA and target node.
+
+BAAW non-responsibilities:
+
+- Performing actual data movement.
+- Executing NOC routing.
+- Simulating bandwidth occupation (downstream components' job).
+
+BAAW output is directly usable by the simulator's routing and resource
+model without additional address decoding.
+
+#### D-LA5. PE_DMA `handle_command()` change
+
+Current (VA-based) flow:
+
+```
+DmaReadCmd.src_addr (VA)
+  → MMU.translate(VA) → PA
+  → PhysAddr.decode(PA) → PhysAddr object
+  → resolver.resolve(PhysAddr) → dst_node_id
+  → router.find_path(pe_prefix, dst_node_id) → path
+  → 1 sub-Transaction → fabric inject
+```
+
+LA-based flow:
+
+```
+DmaReadCmd.src_addr (LA)
+  → BAAW.resolve(LA, nbytes) → list[PhysicalRequest]
+  → for each PhysicalRequest:
+      → router.find_path(pe_prefix, req.dst_node) → path
+      → compute_drain_ns(path, req.nbytes) → drain
+      → sub-Transaction → fabric inject
+  → await all sub-Transactions
+  → pe_txn.done.succeed()
+```
+
+Key changes:
+
+- MMU reference removed → BAAW resolve.
+- `PhysAddr.decode()` + `resolver.resolve()` → BAAW returns `dst_node`
+  directly.
+- 1 request → N parallel requests in 1:1 mode.
+
+#### D-LA6. 1:1 mode detail
+
+- One logical access → N physical requests (N = `channels_per_pe`).
+- N = `hbm_pseudo_channels / pes_per_cube`.
+- Each request: fully-resolved 51-bit PA, targets a specific channel
+  router (`{pe_prefix}.ch_r{channel_id}`).
+- Per-channel link models BW contention.
+- PE_DMA injects N sub-transactions concurrently.
+
+Example: `hbm_pseudo_channels=64`, `pes_per_cube=8` → `channels_per_pe=8`.
+PE0 owns ch0-7.
+
+```text
+Tensor A (4 KB) → LA 0x1_0000_0000, size=4096 bytes
+BAAW segment: {
+    la_base: 0x1_0000_0000, la_size: 4096,
+    mode: "one_to_one", channel_count: 8,
+    pa_bases: [PA_ch0, PA_ch1, ..., PA_ch7],
+    channel_ids: [0, 1, 2, 3, 4, 5, 6, 7],
+    channel_size: 512,
+}
+
+BAAW resolve result (8 requests):
+  → PhysicalRequest(pa=PA_ch0, nbytes=512, dst_node="sip0.cube0.pe0.ch_r0")
+  → PhysicalRequest(pa=PA_ch1, nbytes=512, dst_node="sip0.cube0.pe0.ch_r1")
+  → ...
+  → PhysicalRequest(pa=PA_ch7, nbytes=512, dst_node="sip0.cube0.pe0.ch_r7")
+
+PE_DMA: 8 sub-transactions parallel inject
+  per-channel router → hbm_ctrl link (channel_bw_gbs) per channel
+  Total effective BW = 8 × channel_bw_gbs
+```
+
+Other N values:
+
+- `hbm_pseudo_channels=32`, `pes_per_cube=8` → `channels_per_pe=4`,
+  4 requests
+- `hbm_pseudo_channels=64`, `pes_per_cube=4` → `channels_per_pe=16`,
+  16 requests
+
+#### D-LA7. n:1 mode detail
+
+- One logical access → one aggregated request.
+- Target: aggregated router → hbm_ctrl (see ADR-0017 D8).
+- Aggregated link BW = `channels_per_pe × channel_bw_gbs`
+  (e.g. 8 × 32 = 256 GB/s).
+- Single queue / resource for modelling.
+- No per-channel PA decomposition.
+
+```text
+Tensor A (4 KB) → LA 0x1_0000_0000, size=4096 bytes
+BAAW segment: {
+    la_base: 0x1_0000_0000, la_size: 4096,
+    mode: "n_to_one",
+    agg_pa_base: PA_agg,
+    agg_node_id: "sip0.cube0.pe0.agg_router",
+}
+
+BAAW resolve result:
+  → PhysicalRequest(pa=PA_agg, nbytes=4096, dst_node="sip0.cube0.pe0.agg_router")
+
+PE_DMA: 1 sub-transaction
+  aggregated router → hbm_ctrl link (256 GB/s)
+```
+
+#### D-LA8. Kernel model preserved
+
+- Kernel still issues single memory ops (`tl.load`, `tl.store`,
+  `tl.composite`).
+- LA is the address scheme exposed to kernel code.
+- Channel decomposition / aggregation happens inside PE_DMA's BAAW.
+- Kernel code never sees physical channel information.
+
+#### Consequences (LA model, proposed)
+
+Positive:
+
+- 1:1 vs n:1 semantics live in one place (BAAW).
+- Kernel abstraction preserved — no kernel code changes.
+- Topology-based policy control (mode switch via yaml).
+- Improved simulation-model consistency and debuggability.
+- Segment-based mapping is simpler than page tables; lower overhead.
+
+Negative:
+
+- Full VA/MMU code refactor required.
+- Request-generation path more complex (N requests in 1:1 mode).
+- Reduced per-channel visibility in n:1 mode.
+- VA-related tests need rewriting.
+
+---
+
+## Migration Path
+
+- **PA → VA** was an extension. PA mode is retained as the PageFault
+  fallback inside PE_DMA. Switching does not require removing PA
+  code.
+- **VA → LA**, if adopted, is a replacement, not coexistence. See
+  D-LA1 for the VA infrastructure removal list. PA fallback inside
+  PE_DMA may be retained orthogonally for tests.
+
+## Alternatives Considered (LA model)
+
+1. **Keep VA + fan-out in MMU**: MMU returns per-channel PAs.
+   Rejected: MMU's role would grow beyond translation to request
+   decomposition; aggregation (n:1) becomes awkward to express.
+2. **Channel-aware kernel API**: kernels call per-channel load/store
+   directly. Rejected: abstraction leakage, portability loss, all
+   benchmarks need rewriting.
+3. **Always PA (no LA)**: runtime passes per-channel PA to kernel
+   directly. Rejected: incompatible with aggregation; conversion
+   timing unclear; channel info leaks to kernel.
+
+## Test Requirements
+
+### VA model (current, regression)
+
+- Cross-PE / cross-cube DMA paths over installed mappings.
+- `MmuMapMsg` / `MmuUnmapMsg` fabric traversal with measured latency.
+- TLB-overhead-per-access timing.
+- PageFault fallback path preserves PA-only behaviour.
+
+### LA model (when implemented)
+
+- 1:1 mode: same logical access → N per-channel requests.
+- n:1 mode: same logical access → 1 aggregated request.
+- Bandwidth equivalence between modes for identical workload.
+- 1:1 mode: per-channel contention modelled correctly.
+- n:1 mode: aggregated bandwidth correctly reflected.
+- Kernel code unchanged across mode switch.
+- BAAW segment install / uninstall correctness.
+- Multiple tensors in distinct segments do not collide.
+
+## Implementation Order (LA, when scheduled)
+
+1. LA type (`policy/address/la_allocator.py`).
+2. BAAW segment table (`policy/address/baaw.py`).
+3. `BaawSegmentInstallMsg` (`runtime_api/kernel.py`).
+4. PE_DMA BAAW integration (`components/builtin/pe_dma.py`
+   `handle_command()`).
+5. RuntimeContext: LA alloc + segment install
+   (`runtime_api/context.py`).
+6. `Tensor.va_base` → `Tensor.la_base` (`runtime_api/tensor.py`).
+7. Remove VA/MMU code.
+8. Remove `pe_mmu` from `topology.yaml`; add mapping mode settings.
+9. Test migration:
+
+| Test file | Action |
+|-----------|--------|
+| `tests/test_mmu_component.py` | Remove → BAAW segment install tests |
+| `tests/test_mmu_fabric.py` | Remove → BAAW + fabric integration tests |
+| `tests/test_pe_mmu.py` | Remove |
+| `tests/test_va_allocator.py` | Replace with LA allocator tests |
+| `tests/test_va_integration.py` | Replace with LA + BAAW integration tests |
+| `tests/test_va_offset.py` | Replace with LA offset tests |
+
+## Links
+
+- ADR-0007 (runtime_api vs sim_engine boundaries)
+- ADR-0008 (tensor deployment)
+- ADR-0009 (kernel execution)
+- ADR-0014 (PE-internal execution model)
+- ADR-0015 (component port/wire model)
+- ADR-0017 (Cube NOC and HBM connectivity — LA model topology consumer)
+- ADR-0013 (Verification strategy — V1 PA tagging)
+- SPEC R2 (latency by traversal), R10 (memory addressing)

docs/adr/ADR-0011-memory-addressing-simplification.md

@@ -1,100 +0,0 @@
-# ADR-0011: Memory Addressing — PA-first with VA/MMU Extension
-
-## Status
-
-Accepted (Phase 1 VA/MMU implemented)
-
-## Context
-
-A realistic system uses host-side virtual addressing and an MMU/IOMMU-style
-translation path for DMA: host allocates physical memory at PE level, maps it
-into a virtual address space, installs mappings, and DMA requests use virtual
-addresses that are translated to physical addresses.
-
-The PA-only model (Phase 0) was insufficient for running standard Triton kernels
-that use `base_addr + offset` patterns on sharded tensors — each PE's shard has
-a different PA, but the kernel needs a single contiguous address space.
-
----
-
-## Decision
-
-### D1. Phase 0 model is PA-only (original, retained as fallback)
-
-- All device memory accesses (MemoryRead/MemoryWrite) operate on device physical
-  addresses (PA) plus size.
-- PA-only mode remains functional via PageFault fallback in PE_DMA.
-
-### D2. Allocation produces PA mappings
-
-Device allocation selects PE-local memory regions and returns PA mappings
-sufficient to execute kernels and issue DMA requests.
-
-### D3. Phase 1: VA/MMU layer (implemented)
-
-#### D3.1 Virtual Address Model
-
-- Each tensor gets a single contiguous VA range (`TensorHandle.va_base`).
-- `TensorShard` does NOT carry a `va` field — shard VA is derived as
-  `va_base + offset_bytes`.
-- Kernels receive `va_base` as their pointer argument (via `TensorArg.va_base`).
-- `DmaReadCmd.src_addr` and `DmaWriteCmd.dst_addr` carry VA (not PA).
-
-#### D3.2 PE_MMU Component
-
-- Hybrid design: SimPy component (inbox for MmuMapMsg) + utility (synchronous
-  `translate()` called by PE_DMA).
-- Page-aligned dict lookup for O(1) VA→PA translation.
-- `tlb_overhead_ns` configurable per-access latency.
-- PageFault fallback: if VA has no mapping, PE_DMA treats it as PA directly
-  (backward compatibility with PA-only tests).
-
-#### D3.3 Mapping Installation
-
-- `MmuMapMsg` traverses the fabric: Host → PCIE_EP → IO_CPU (cube fan-out) →
-  M_CPU (PE fan-out) → NOC → PE_MMU. Latency is measured end-to-end.
-- `MmuMapMsg.target_sips` controls SIP-level routing to prevent cross-SIP
-  mapping contamination for replicated tensors.
-- Mapping strategy based on `DPPolicy.cube`:
-  - **Replicate** (`cube="replicate"`): per-(sip, cube) local mapping only.
-    Each cube's PEs see only their local PA. No cross-cube mapping installed.
-  - **Sharded** (`cube="column_wise"`, etc.): broadcast all shard mappings to all
-    target cubes. Enables cross-PE and cross-cube DMA.
-
-#### D3.4 Tensor Lifecycle
-
-- `del tensor` triggers automatic cleanup via `Tensor.__del__` + `weakref` to
-  RuntimeContext. Sends `MmuUnmapMsg` through fabric, returns VA and PA space.
-- `with RuntimeContext(...) as ctx:` provides scope-based bulk cleanup.
-- `RuntimeContext._tensors` uses `weakref.ref` to avoid preventing GC.
-- `PEMemAllocator` uses free-list with coalescing (not bump allocator).
-- `VirtualAllocator` uses free-list with coalescing for VA space.
-
-#### D3.5 Allocators
-
-- `VirtualAllocator`: device-wide VA space, page-aligned alloc/free with
-  coalescing.
-- `PEMemAllocator`: per-PE HBM/TCM, free-list based alloc/free with coalescing.
-- Page size configurable via `topology.yaml` pe_mmu attrs (default 4096).
-
----
-
-## Consequences
-
-- Triton kernels use `base_addr + offset` patterns naturally on sharded tensors.
-- All latency remains explicit via graph traversal, including MMU mapping
-  installation and per-access TLB overhead.
-- PA-only mode retained as fallback (PageFault → treat as PA).
-- Benchmark parameter renamed `ctx` → `torch` for PyTorch code compatibility.
-- IPCQ and other fixed-address resources bypass MMU (use PA directly).
-
----
-
-## Links
-
-- ADR-0007 (runtime_api vs sim_engine boundaries)
-- ADR-0008 (tensor deployment)
-- ADR-0009 (kernel execution)
-- ADR-0014 (PE-internal execution model)
-- ADR-0015 (component port/wire model)
-- SPEC R2 (latency by traversal)

docs/adr/ADR-0012-api-host-io-message-schema.md

@@ -18,7 +18,7 @@ We define stable, minimal message schemas for Host ↔ IO_CPU so that:
 - IO_CPU-internal fan-out/aggregation can evolve independently,
 - completion and failure propagation is deterministic.
 
-We also require PE-tagging (A 방식): each shard explicitly carries (sip,cube,pe)
+We also require PE-tagging (Scheme A): each shard explicitly carries (sip,cube,pe)
 so IO_CPU can deterministically route/fan-out without relying on PA decoding.
 
 ---
@@ -93,7 +93,7 @@ Rules:
 Mandatory fields:
 
 - common envelope fields (D3)
-- destination placement tags (A 방식):
+- destination placement tags (Scheme A):
   - `dst_sip: int`
   - `dst_cube: int`
   - `dst_pe: int`
@@ -130,7 +130,7 @@ Notes:
 Mandatory fields:
 
 - common envelope fields (D3)
-- source placement tags (A 방식):
+- source placement tags (Scheme A):
   - `src_sip: int`
   - `src_cube: int`
   - `src_pe: int`
@@ -183,7 +183,7 @@ Tensor arg (mandatory):
 
 - `shards: list[TensorShard]`
 
-`TensorShard` MUST have (A 방식 강제):
+`TensorShard` MUST have (Scheme A enforced):
 
 - `sip: int`
 - `cube: int`
@@ -226,7 +226,8 @@ Tests SHOULD validate:
 
 ## Links
 
-- ADR-0011 (PA-first memory addressing)
+- ADR-0011 (Memory Addressing — PA / VA / LA)
 - ADR-0007 (runtime_api vs sim_engine boundaries)
 - ADR-0009 (kernel execution fan-out/aggregation)
+- ADR-0013 (Verification strategy — V1 message schema validation)
 - SPEC R2, R7, R8

docs/adr/ADR-0013-ver-verification-strategy.md

@@ -134,6 +134,6 @@ Phase 2 (Apply) MUST:
 ## Links
 
 - SPEC 0.1, R2, R6
-- ADR-0011 (PA-first memory addressing)
+- ADR-0011 (Memory Addressing — PA / VA / LA)
 - ADR-0012 (Host ↔ IO_CPU message schema)
 - ADR-0009 (Kernel execution semantics)

docs/adr/ADR-0014-dev-pe-pipeline-execution-model.md

@@ -0,0 +1,451 @@
+# ADR-0014: PE Pipeline Execution Model
+
+## Status
+
+Accepted
+
+## Context
+
+This ADR defines the PE-internal kernel execution model:
+
+- Role decomposition of PE-internal components
+- Command dispatch paths (simple / composite / multi-op composite with epilogue)
+- TileToken-based self-routing pipeline (scheduler does dispatch + completion only)
+- TCM-centric dataflow with a register-file intermediary
+- Engine resource model
+- Observability and trace contract
+- Topology representation
+
+PE-internal structure (7 components in scope; 2 cross-referenced):
+
+- `pe_cpu`, `pe_scheduler`, `pe_dma`, `pe_fetch_store`, `pe_gemm`, `pe_math`,
+  `pe_tcm` — defined here
+- `pe_mmu` — VA model, defined in ADR-0011 D-VA
+- `pe_ipcq` — collective communication, defined in ADR-0023
+
+The goal is a deterministic, trace-friendly execution contract that keeps
+each block independently swappable.
+
+## Decision
+
+### D1. PE-internal component roles
+
+**PE_CPU**
+
+- Executes kernel instruction stream / control logic.
+- Generates PE commands and submits them to `PE_SCHEDULER` (via
+  `PeInternalTxn`).
+- Does NOT enqueue work directly into engine queues.
+
+**PE_SCHEDULER**
+
+- Sole dispatcher inside a PE.
+- Receives commands from `PE_CPU`. Dispatch by command type:
+  - Simple command (`DmaReadCmd`, `DmaWriteCmd`, `GemmCmd`, `MathCmd`)
+    → forward directly to the target engine.
+  - `CompositeCmd` → generate a `TilePlan`, feed tiles into the pipeline
+    via a single `_feed_loop` (D6).
+- Does not participate in stage-to-stage chaining within a composite;
+  that is handled by token self-routing (D6).
+
+**PE_DMA**
+
+- Handles memory transfers between TCM and external memory domains
+  (HBM, shared SRAM, cross-cube UCIe) through the cube NOC.
+- Two execution channels:
+  - `DMA_READ` (capacity = 1) and `DMA_WRITE` (capacity = 1) — see D4.
+- Additional virtual channels:
+  - `vc_compute` — load/store/writeback traffic for GEMM/MATH tiles.
+  - `vc_comm` — IPCQ collective send data (defined in ADR-0023 D8).
+
+**PE_FETCH_STORE**
+
+- TCM ↔ Register File transfer unit.
+- Isolates register-file access semantics from compute engines so that
+  GEMM/MATH stay pure compute components.
+- BW-based latency model; TCM access contention naturally serializes
+  through `PE_TCM`'s BW resource.
+
+**PE_GEMM**
+
+- MAC array. Reads operands from the register file; writes results to
+  the register file. Does not touch `PE_TCM` directly.
+
+**PE_MATH**
+
+- Element-wise / reduction / SIMD unit. Reads / writes the register file.
+
+**PE_TCM**
+
+- Tightly-coupled scratchpad with BW-serialized access. Two logical
+  regions partitioned by ownership (see D5).
+
+**Cross-referenced components** (defined elsewhere):
+
+- `pe_mmu` — VA→PA translation per access (ADR-0011 D-VA).
+- `pe_ipcq` — collective ring buffers and peer endpoint metadata
+  (ADR-0023).
+
+### D2. Command lifecycle and queues
+
+`PE_SCHEDULER` maintains three logical structures:
+
+**SubmissionQueue** — written by `PE_CPU`; consumed by the scheduler.
+
+**InflightTable** — owned and mutated only by `PE_SCHEDULER`; tracks
+expanded sub-commands, dependency state, engine assignment, and
+completion status.
+
+**CompletionQueue** — written by `PE_SCHEDULER`; holds final completion
+records.
+
+**Single-writer rule**: only `PE_SCHEDULER` mutates command completion
+state. Engines report completion via explicit events / messages
+consumed by the scheduler.
+
+**Command completion**: when all sub-commands complete, `PE_SCHEDULER`
+publishes a completion record.
+
+### D3. Dispatch modes
+
+#### D3.1 Simple command
+
+A simple command expands to exactly one engine sub-command:
+
+- `DmaReadCmd` / `DmaWriteCmd` → `PE_DMA`
+- `GemmCmd` → `PE_GEMM`
+- `MathCmd` → `PE_MATH`
+
+Flow:
+
+```text
+PE_CPU → SubmissionQueue → PE_SCHEDULER → engine queue → engine execution
+       → completion → PE_SCHEDULER → CompletionQueue
+```
+
+#### D3.2 Composite command (single-op tiled pipeline)
+
+The default `CompositeCmd` runs a single compute op as a tile-pipelined
+sequence:
+
+```text
+DMA_READ → FETCH (TCM → RF) → COMPUTE (GEMM | MATH) → STORE (RF → TCM) → DMA_WRITE
+```
+
+`PE_SCHEDULER` splits the DMA payload into hardware tiles and emits one
+`TileToken` per tile with a monotonically increasing `tile_id`.
+
+Tile dependency (within one tile `t`):
+
+```text
+DMA_READ(t) → FETCH(t) → COMPUTE(t) → STORE(t) → DMA_WRITE(t)
+```
+
+Inter-tile overlap is allowed wherever engine resources permit
+(D4 governs the constraints):
+
+```text
+DMA_READ(t+1) ∥ COMPUTE(t)
+DMA_WRITE(t-1) ∥ COMPUTE(t)
+```
+
+#### D3.3 Multi-op composite (head + epilogue with scope)
+
+A `CompositeCmd` MAY carry `ops: tuple[OpSpec, ...]` to express a
+multi-op pipeline:
+
+```python
+@dataclass(frozen=True)
+class OpSpec:
+    kind: str         # "gemm" | "math.exp" | "math.bias_add" | ...
+    scope: Scope      # "per_k_tile" | "per_output_tile" | "once"
+    ...
+```
+
+- `ops[0]` (head) defines tile geometry (e.g., the head GEMM determines
+  M/K/N partition).
+- `ops[1:]` (epilogue) are subsequent stages whose `scope` decides how
+  often they fire:
+  - `per_k_tile` — every K-reduction step.
+  - `per_output_tile` — once per output tile.
+  - `once` — once per kernel.
+
+Cross-engine chains (e.g., GEMM head → MATH epilogue) are natural —
+each stage is dispatched via token self-routing (D6), so GEMM and MATH
+participate serially within the same composite even though they share
+the compute slot (D4).
+
+The empty-`ops` form is the legacy single-op path.
+
+### D4. Engine resource model
+
+**DMA engine**:
+
+- `DMA_READ`: `simpy.Resource(capacity=1)`.
+- `DMA_WRITE`: `simpy.Resource(capacity=1)`.
+- Both channels run concurrently (READ ∥ WRITE allowed).
+- Within a channel, requests serialize (READ ∥ READ disallowed; same
+  for WRITE).
+- `vc_comm` is an orthogonal channel for IPCQ traffic defined in
+  ADR-0023 D8 — out of scope for this ADR.
+
+**Compute engine**:
+
+- `accel_slot`: `simpy.Resource(capacity=1)` shared by `PE_GEMM` and
+  `PE_MATH`.
+- At most one compute op runs at a time within a PE.
+- Multi-op composite chains (D3.3) execute their compute stages serially
+  through this slot; token self-routing (D6) ensures the next stage
+  starts only after the previous compute releases the slot.
+
+**Engine completion**: each engine emits a completion event consumed by
+the scheduler / `PipelineContext` (D6).
+
+### D5. Dataflow
+
+**Input path (HBM source)**:
+
+```text
+HBM → cube NOC → PE_DMA (DMA_READ) → PE_TCM
+PE_TCM → PE_FETCH_STORE → Register File
+Register File → PE_GEMM | PE_MATH
+```
+
+**Input path (shared SRAM source)**:
+
+```text
+Shared SRAM → cube NOC → PE_DMA (DMA_READ) → PE_TCM
+PE_TCM → PE_FETCH_STORE → Register File
+```
+
+**Output path (HBM destination)**:
+
+```text
+Register File → PE_FETCH_STORE → PE_TCM
+PE_TCM → PE_DMA (DMA_WRITE) → cube NOC → HBM
+```
+
+GEMM/MATH never touch `PE_TCM` directly — `PE_FETCH_STORE` is the
+single TCM↔register-file gateway. This makes TCM BW contention
+explicit and lets fetch unit policies (e.g., prefetch) be replaced
+independently of compute engines.
+
+#### D5.1 PE_TCM partitioning
+
+`PE_TCM` is split into two logical regions:
+
+**SchedulerReservedTCM**
+
+- Owned exclusively by `PE_SCHEDULER`.
+- Holds composite-command tile buffers.
+- `PE_SCHEDULER` partitions this region, assigns buffers per DMA_READ /
+  COMPUTE / DMA_WRITE stage, guarantees input/output separation, and
+  manages tile-buffer lifetimes.
+
+**AllocatableTCM**
+
+- General-purpose region managed by `PEMemAllocator`.
+- Used for host / DP-visible allocations.
+
+**Visibility rule (hard isolation)**: `PEMemAllocator` MUST NOT see or
+allocate inside `SchedulerReservedTCM`. The reserved region is excluded
+from allocator-managed ranges by construction.
+
+**Tile buffer rules**:
+
+- Input and output buffers within `SchedulerReservedTCM` MUST NOT
+  overlap during a tile's active lifetime.
+- A tile buffer remains valid until the corresponding `DMA_WRITE`
+  completes.
+- Buffer reuse is permitted only after the consuming tile's lifetime
+  ends.
+
+### D6. TileToken self-routing pipeline
+
+A composite's stage-to-stage progression happens **without** routing
+through the scheduler. Each component forwards the token directly to
+the next stage's component using the token's `plan`:
+
+```text
+Scheduler → DMA → Fetch → GEMM → Math (epi) → Store → DMA_WB → (complete)
+              ↑ chaining: no scheduler hop                          ↑
+                                                  PipelineContext.complete_tile()
+```
+
+This mirrors real-HW done-wire chains. The scheduler handles only
+**initial dispatch + completion aggregation**.
+
+#### TilePlan / Stage
+
+```python
+class StageType(Enum):
+    DMA_READ = 0
+    FETCH = 1
+    GEMM = 2
+    MATH = 3
+    STORE = 4
+    DMA_WRITE = 5
+
+@dataclass(frozen=True)
+class Stage:
+    stage_type: StageType
+    component: str         # topology node id (e.g., "sip0.cube0.pe0.pe_dma")
+    params: dict           # stage-specific parameters
+
+@dataclass(frozen=True)
+class TilePlan:
+    tile_id: int
+    stages: tuple[Stage, ...]
+```
+
+#### TileToken
+
+```python
+@dataclass
+class TileToken:
+    tile_id: int
+    pipeline_ctx: PipelineContext
+    plan: TilePlan
+    stage_idx: int
+    params: dict             # cached current stage params
+    data_op: bool = True     # op_log opt-in (ADR-0020 D4)
+```
+
+Single-owner invariant: a token is owned by exactly one component at a
+time. Lifecycle: scheduler creates with `stage_idx=0` → component
+`_process()` → increment `stage_idx` → put to next stage's `in_port` →
+last stage calls `pipeline_ctx.complete_tile()`.
+
+#### PipelineContext (exactly-once completion)
+
+```python
+@dataclass
+class PipelineContext:
+    id: str
+    total_tiles: int
+    completed_tiles: int = 0
+    done_event: simpy.Event = None
+
+    def complete_tile(self) -> None:
+        self.completed_tiles += 1
+        if self.completed_tiles == self.total_tiles:
+            self.done_event.succeed()
+```
+
+Each tile's last stage MUST call `complete_tile()` exactly once.
+Duplicate calls are bugs (SimPy `Event` can succeed at most once).
+
+#### Feed ordering
+
+`PE_SCHEDULER` has exactly one `_feed_loop` process consuming a
+`_pending_feeds` FIFO. Composite commands are enqueued in submission
+order; tile feed for a command runs to completion before the next
+command's feed begins. **Tile-feed interleaving between commands is
+disallowed.**
+
+Within a single command's tiles, downstream pipeline overlap arises
+naturally — earlier tiles progress through later stages while the feeder
+keeps pushing remaining tiles into the first stage queue (SimPy Store
+backpressure governs flow control). If the first-stage queue is full,
+only the feeder blocks; the scheduler worker's inbox processing
+continues.
+
+#### Token routing pattern (base class)
+
+```python
+def _pipeline_worker(self, env):
+    while True:
+        token = yield self._inbox.get()
+        yield from self._process(env, token)       # stage-specific logic
+        next_idx = token.stage_idx + 1
+        if next_idx < len(token.plan.stages):
+            next_stage = token.plan.stages[next_idx]
+            token.stage_idx = next_idx
+            token.params = next_stage.params
+            yield self.out_ports[next_stage.component].put(token)
+        else:
+            token.pipeline_ctx.complete_tile()
+```
+
+Each component implements only `_process()`; chaining lives in the
+base class.
+
+### D7. Observability and trace contract
+
+The simulator emits deterministic trace events:
+
+- `command_submitted`
+- `sub_command_dispatched`
+- `engine_start`
+- `engine_complete`
+- `tile_ready`
+- `command_complete`
+
+For identical inputs, trace ordering MUST be deterministic.
+
+### D8. Topology representation
+
+PE-internal components are declared in `cube.pe_template`:
+
+```yaml
+pe_template:
+  components:
+    pe_cpu:         { kind: pe_cpu,         impl: builtin.pe_cpu,         attrs: { overhead_ns: ... } }
+    pe_scheduler:   { kind: pe_scheduler,   impl: builtin.pe_scheduler,   attrs: { overhead_ns: ... } }
+    pe_dma:         { kind: pe_dma,         impl: builtin.pe_dma,         attrs: { rd_engines: 1, wr_engines: 1 } }
+    pe_fetch_store: { kind: pe_fetch_store, impl: builtin.pe_fetch_store, attrs: { ... } }
+    pe_gemm:        { kind: pe_gemm,        impl: builtin.pe_gemm,        attrs: { shared_resource: accel_slot, ... } }
+    pe_math:        { kind: pe_math,        impl: builtin.pe_math,        attrs: { shared_resource: accel_slot, ... } }
+    pe_tcm:         { kind: pe_tcm,         impl: builtin.pe_tcm,         attrs: { size_mb: ..., read_bw_gbs: ..., write_bw_gbs: ... } }
+    pe_mmu:         { kind: pe_mmu,         impl: builtin.pe_mmu,         attrs: { ... } }   # ADR-0011 D-VA
+    pe_ipcq:        { kind: pe_ipcq,        impl: builtin.pe_ipcq,        attrs: { ... } }   # ADR-0023
+  links:
+    # Scheduler dispatch edges (initial)
+    scheduler_to_dma_mm:         0.0
+    scheduler_to_fetch_store_mm: 0.0
+    scheduler_to_gemm_mm:        0.0
+    scheduler_to_math_mm:        0.0
+    # Pipeline chaining edges (token self-routing per D6)
+    dma_to_fetch_store_mm:       0.0
+    fetch_store_to_gemm_mm:      0.0
+    fetch_store_to_math_mm:      0.0
+    gemm_to_fetch_store_mm:      0.0
+    gemm_to_math_mm:             0.0
+    math_to_fetch_store_mm:      0.0
+    fetch_store_to_dma_mm:       0.0
+    fetch_store_to_tcm_bw_gbs:   ...
+```
+
+Template is instantiated once per PE. PE instances are derived from
+`cube.pe_layout` (corner placement). External connectivity (PE_DMA ↔
+cube NOC ↔ HBM, etc.) is modeled at the cube level (ADR-0017 D4).
+
+## Consequences
+
+### Positive
+
+- Each block is an independent topology node — individually swappable
+  via DI (ADR-0015).
+- PE-internal structure is visible in the topology graph.
+- Components do not know their downstream — plan-based routing gives
+  flexibility (e.g., epilogue chains require no scheduler change).
+- DMA and compute overlap naturally via SimPy Store backpressure.
+- Multi-op composite expresses fused operations (e.g., GEMM + bias_add)
+  without engine-level coupling.
+- TCM access contention is realistic — `PE_FETCH_STORE` is the single
+  TCM↔RF gateway.
+
+### Negative
+
+- Intra-PE component count is higher than a coarser model (7 base + 2
+  cross-referenced) — more topology nodes/edges.
+- Intra-PE token forwarding is explicit in traces (acceptable trade for
+  HW fidelity).
+
+## Links
+
+- ADR-0011 D-VA (PE_MMU component, VA translation)
+- ADR-0015 D4 (component port/wire model)
+- ADR-0020 (greenlet kernel execution / two-pass)
+- ADR-0023 (PE_IPCQ + PE_DMA virtual channels)
+- SPEC R3, R4

docs/adr/ADR-0014-pe-internal-execution-model.md

@@ -1,365 +0,0 @@
-# ADR-0014: PE Internal Execution Model (PE_CPU, PE_SCHEDULER, and Composite Commands)
-
-## Status
-
-Accepted
-
-## Context
-
-ADR-0003 (system hierarchy) and ADR-0009 (kernel execution semantics) reference PE internals but do not define:
-
-- the dispatch model inside a PE,
-- the responsibilities of PE_SCHEDULER,
-- the PE_TCM-centric dataflow contract used by accelerator engines.
-
-We need a deterministic and debuggable PE-internal execution contract that supports:
-
-- simple single-engine commands
-- composite commands that build a tiled pipeline across DMA and accelerator engines
-
-The simulator must produce deterministic traces and allow modeling of PE-internal pipelining without introducing nondeterministic engine scheduling.
-
-## Decision
-
-### D1. PE internal component roles
-
-Each PE contains the following logical components.
-
-**PE_CPU**
-
-- Executes kernel instruction stream or kernel control logic.
-- Generates PE commands.
-- Submits commands to PE_SCHEDULER.
-- PE_CPU does NOT enqueue work directly into engine queues.
-
-**PE_SCHEDULER**
-
-- The sole dispatcher inside a PE.
-- Receives commands from PE_CPU.
-- Expands composite commands into sub-commands.
-- Tracks dependencies and command state.
-- Dispatches work to engine queues.
-- Manages tile scheduling for composite commands.
-
-**PE_DMA**
-
-- Handles memory transfers between PE_TCM and external memory domains.
-- PE_DMA connects to the NOC router mesh at the CUBE level (ADR-0019):
-  - All destinations (HBM, shared SRAM, inter-cube UCIe) are reached via the router mesh
-  - Local HBM access: PE_DMA → local router → hbm_ctrl (switching overhead only)
-  - Remote/shared: PE_DMA → local router → (mesh hops) → destination
-- Supported directions include:
-  - HBM → PE_TCM (via router mesh)
-  - PE_TCM → HBM (via router mesh)
-  - PE_TCM → shared SRAM (via router mesh)
-  - PE_TCM → other memory domains (via router mesh, if supported by topology)
-
-**PE_GEMM**
-
-- Matrix multiplication engine.
-- Reads activations from PE_TCM.
-- May stream weights directly from HBM.
-
-**PE_MATH**
-
-- Element-wise computation engine.
-- Reads and writes PE_TCM.
-
-**PE_TCM**
-
-- Local SRAM used as the staging memory for accelerator operations.
-
----
-
-### D2. Command lifecycle and queues
-
-PE_SCHEDULER maintains three logical structures.
-
-**SubmissionQueue**
-
-- Written by PE_CPU.
-- Contains incoming PE commands waiting to be processed.
-
-**InflightTable**
-
-- Owned and mutated only by PE_SCHEDULER.
-- Tracks:
-  - expanded sub-commands
-  - dependency state
-  - engine assignment
-  - completion status
-
-**CompletionQueue**
-
-- Written by PE_SCHEDULER.
-- Contains final completion records for commands.
-
-**Single-writer rule**
-
-- Only PE_SCHEDULER is allowed to mutate command completion state.
-- Engine components must report completion via explicit completion events/messages.
-
-**Command completion**
-
-A command becomes DONE when:
-
-- all sub-commands complete
-- PE_SCHEDULER publishes a completion record to CompletionQueue.
-
----
-
-### D3. Dispatch modes
-
-PE commands are divided into two categories.
-
-#### D3.1 Simple command
-
-A simple command expands to exactly one engine sub-command.
-
-Examples include:
-
-- DMA transfer
-- GEMM compute
-- MATH compute
-
-Execution flow:
-
-```text
-PE_CPU → SubmissionQueue → PE_SCHEDULER → engine queue → engine execution → completion event → PE_SCHEDULER → CompletionQueue
-```
-
-#### D3.2 Composite command (tiled pipeline)
-
-Composite commands implement tiled pipelined execution across engines.
-
-Each tile executes the following pipeline:
-
-```text
-Input DMA (READ)
-→ Compute (GEMM or MATH)
-→ Output DMA (WRITE)
-```
-
-**Tiling rule**
-
-If the DMA payload exceeds hardware tile size, PE_SCHEDULER splits the transfer into tiles.
-Each tile is assigned a monotonically increasing `tile_id`.
-
-**Tile dependency rules**
-
-For tile `t`:
-
-- Compute must wait for input DMA: `DMA_READ(t) → COMPUTE(t)`
-- Output DMA must wait for compute: `COMPUTE(t) → DMA_WRITE(t)`
-- All dependencies are enforced by PE_SCHEDULER.
-
-**Overlap policy (Phase 0 default)**
-
-Operations for different tiles may overlap when engine resources permit.
-
-Allowed overlaps:
-
-```text
-DMA_READ(t+1) ∥ COMPUTE(t)
-DMA_WRITE(t−1) ∥ COMPUTE(t)
-DMA_READ(t) ∥ DMA_WRITE(t)
-```
-
-Disallowed overlaps:
-
-```text
-GEMM(t) ∥ GEMM(t′)
-MATH(t) ∥ MATH(t′)
-GEMM(t) ∥ MATH(t′)
-```
-
----
-
-### D4. Engine execution model (Phase 0 default)
-
-Each engine behaves as a deterministic service resource.
-
-**DMA engine**
-
-PE_DMA contains two independent channels.
-
-```text
-DMA_READ capacity  = 1
-DMA_WRITE capacity = 1
-```
-
-Rules:
-
-- DMA_READ and DMA_WRITE may execute concurrently.
-- Multiple READs cannot overlap.
-- Multiple WRITEs cannot overlap.
-
-Example allowed:
-
-```text
-DMA_READ(t+1) ∥ DMA_WRITE(t)
-```
-
-Example not allowed:
-
-```text
-DMA_READ(t) ∥ DMA_READ(t+1)
-DMA_WRITE(t) ∥ DMA_WRITE(t+1)
-```
-
-**Compute engine**
-
-Compute operations share a single compute resource.
-
-```text
-PE_ACCEL capacity = 1
-```
-
-Both GEMM and MATH require this shared compute slot.
-
-Consequences:
-
-- GEMM ∥ GEMM not allowed
-- MATH ∥ MATH not allowed
-- GEMM ∥ MATH not allowed
-
-Only one compute operation can run in a PE at a time.
-
-**Compute opcode restriction**
-
-Composite commands contain one compute opcode only.
-
-Examples:
-
-```text
-COMPOSITE_GEMM
-COMPOSITE_MATH
-```
-
-Mixed compute pipelines such as `GEMM → MATH` are not supported in Phase 0.
-
-**Engine completion signaling**
-
-Every engine emits a completion event when a sub-command finishes.
-Completion events are delivered to PE_SCHEDULER.
-
----
-
-### D5. Dataflow model
-
-Compute operations use a TCM-centric dataflow model.
-
-**Input path (HBM)**
-
-```text
-HBM → router mesh → PE_DMA (DMA_READ) → PE_TCM
-```
-
-**Input path (shared SRAM)**
-
-```text
-Shared SRAM → NOC → PE_DMA (DMA_READ) → PE_TCM
-```
-
-**Compute stage**
-
-Compute engines read input tensors from PE_TCM.
-
-```text
-PE_TCM → GEMM / MATH
-```
-
-Weights for GEMM may optionally stream directly from HBM (via router mesh).
-
-**Output path (HBM)**
-
-Compute results are written to PE_TCM, then DMA writes to HBM.
-
-```text
-PE_TCM → PE_DMA (DMA_WRITE) → router mesh → HBM
-```
-
-**Output path (shared SRAM)**
-
-```text
-PE_TCM → PE_DMA (DMA_WRITE) → NOC → Shared SRAM
-```
-
-#### D5.1 PE_TCM partitioning and ownership boundary
-
-The PE_TCM address space is partitioned into two logical regions.
-
-**SchedulerReservedTCM**
-
-- A staging region owned exclusively by PE_SCHEDULER.
-- This region is used for composite command tile buffers.
-- PE_SCHEDULER:
-  - partitions this region into tile buffers
-  - assigns buffers for DMA_READ, COMPUTE, and DMA_WRITE stages
-  - guarantees input/output buffer separation
-  - manages tile buffer lifetime
-
-**AllocatableTCM**
-
-- General-purpose region managed by PEMemAllocator.
-- Used by host or DP-visible allocations.
-
-**Visibility rule (hard isolation)**
-
-- PEMemAllocator must not see or allocate memory inside SchedulerReservedTCM.
-- SchedulerReservedTCM is excluded from allocator-managed ranges by construction.
-- This prevents DP or host allocations from interfering with scheduler staging buffers.
-
-**Tile buffer rules**
-
-Within SchedulerReservedTCM:
-
-- input buffers and output buffers must not overlap
-- PE_SCHEDULER assigns tile buffers for DMA and compute stages
-- tile buffers remain valid until the corresponding DMA_WRITE completes
-- Buffer reuse is allowed only after the tile lifetime finishes.
-
----
-
-### D6. Observability and trace contract
-
-The simulator must emit deterministic trace events.
-
-Required events include:
-
-- `command_submitted`
-- `sub_command_dispatched`
-- `engine_start`
-- `engine_complete`
-- `tile_ready`
-- `command_complete`
-
-Trace ordering must be deterministic for identical inputs.
-
----
-
-### D7. Topology representation
-
-PE internal components are declared in `cube.pe_template`.
-
-The template is instantiated once per PE.
-
-PE instances are derived from `cube.pe_layout`.
-
-External connectivity such as:
-
-- PE_DMA → router mesh → HBM (data path, ADR-0019)
-- PE_DMA → router mesh → shared SRAM, inter-cube UCIe (non-HBM data path)
-- router mesh → PE_CPU (command path from M_CPU)
-
-is modeled at the CUBE level (see ADR-0003 D3).
-
----
-
-## Links
-
-- SPEC R3, R4
-- ADR-0003 D4 (PE-level system hierarchy)
-- ADR-0005 View C (PE-level diagram)
-- ADR-0008 D2 (PA-level allocation at PE scope; PEMemAllocator is the per-PE allocator instance)
-- ADR-0009 D3 (kernel execution fan-out and PE_CPU dispatch)

docs/adr/ADR-0015-dev-component-port-wire-model.md

@@ -6,20 +6,19 @@ Accepted
 
 ## Context
 
-ADR-0007 D2 assigns path-walking and low-level request decomposition to the simulation engine.
-In practice, the engine iterates the topology path and calls `run()` on each component
-sequentially — conflating routing policy with component behavior and preventing realistic
-hardware modeling (queues, contention, fan-out).
-
-ADR-0007 D3 already states that components own fan-out and aggregation, but the current
-implementation does not enforce this for fabric traversal.
+Realistic hardware modeling — queues, contention, fan-out — requires
+that components own fabric traversal while the simulation engine
+handles only initialization and completion observation. Direct method
+calls between components, or path-walking inside the engine, defeat
+queueing and contention semantics.
 
 This ADR defines:
 
 - how components communicate via typed port queues,
 - how propagation delay is modeled (wire processes with BW occupancy),
-- the fabric paths for Memory R/W (M_CPU bypass) and Kernel Launch (via M_CPU),
-- the reduced role of the simulation engine,
+- the fabric paths for Memory R/W (M_CPU bypass) and Kernel Launch
+  (via M_CPU),
+- the engine's reduced role (wire init + completion observation only),
 - M_CPU.DMA as an internal subcomponent of M_CPU.
 
 ---
@@ -88,9 +87,6 @@ The simulation engine MUST NOT:
 - call component `run()` methods directly,
 - track per-hop latency or decompose fan-out.
 
-This supersedes ADR-0007 D2's "decompose operations into low-level requests" clause.
-ADR-0007 D2 must be amended accordingly.
-
 ---
 
 ### D4. Fabric paths for Memory R/W and Kernel Launch
@@ -192,16 +188,15 @@ It is used for shard comparison in `_route_kernel` and as a regression guard.
 - Propagation delay is modeled accurately per edge.
 - Engine is decoupled from routing policy.
 - Component implementations remain swappable via DI (ADR-0007 D3).
-- ADR-0007 D2 must be amended to remove path-walking from engine responsibilities.
-- ADR-0009 D3 should be updated to reference the unified fabric path (D4 above).
 
 ---
 
 ## Links
 
-- ADR-0007 D2 (to be amended: engine path-walking clause)
-- ADR-0009 D3 (kernel execution fan-out; fabric path to be referenced)
+- ADR-0007 D2 (engine role boundary)
+- ADR-0009 D3 (kernel execution fan-out hierarchy)
 - ADR-0014 D4 (DMA engine capacity=1)
 - ADR-0012 D1 (host ↔ IO_CPU message schema; M_CPU.DMA is component-internal)
 - ADR-0016 (IOChiplet NOC and memory data path)
 - ADR-0017 (cube NOC 2D mesh architecture)
+- ADR-0033 (Latency model assumptions built on these mechanisms)

docs/adr/ADR-0017-cube-noc-2d-mesh.md

@@ -1,189 +0,0 @@
-# ADR-0017: Cube NOC 2D Mesh Architecture
-
-## Status
-
-Accepted
-
-## Context
-
-ADR-0003 D3 defines the cube-level NOC as a "distributed on-die fabric" but
-does not specify the internal routing model, contention semantics, or
-attachment topology. The implementation uses a 2D mesh router grid with
-XY routing and per-segment contention modeling. This ADR formalizes that
-architecture.
-
-## Decision
-
-### D1. NOC node and router grid
-
-Each cube contains a 2D router mesh generated by `mesh_gen.py`.
-Each router is a separate topology node (`sip{S}.cube{C}.r{row}c{col}`)
-implemented as `forwarding_v1`. (Supersedes the original single-node
-`noc_2d_mesh_v1` design — see ADR-0019.)
-
-Grid properties:
-
-- Default dimensions: 6x6 routers (derived from PE layout + UCIe connections)
-- Router naming: `r{row}c{col}` (e.g., `r0c0`, `r5c5`)
-- HBM exclusion zone: center rows/columns are excluded where HBM physically
-  occupies space (e.g., r2c2, r2c3, r3c2, r3c3)
-- Router positions are derived from physical PE corner placement and cube
-  geometry
-
-The NOC overhead_ns is 0.0. Latency is modeled by Manhattan distance
-traversal within the mesh (distance_mm x ns_per_mm).
-
-### D2. XY routing algorithm
-
-The NOC uses deterministic XY routing:
-
-1. Horizontal segment: route from source X to destination X at source Y
-2. Vertical segment: route from destination X at source Y to destination Y
-
-Each directed segment is identified by a unique link key:
-
-- Horizontal: `("H", y_band, x_min, x_max, direction)`
-- Vertical: `("V", x_band, y_min, y_max, direction)`
-
-Grid positions are snapped to the router grid, excluding the HBM zone.
-
-### D3. Contention model
-
-Each directed XY segment is a `simpy.Resource(capacity=1)`. Transactions
-sharing a segment (same row or column band, same direction) contend for the
-resource. This models link-level serialization in a wormhole-routed mesh.
-
-With no contention, NOC traversal latency equals the Manhattan distance
-multiplied by `ns_per_mm`. Under contention, additional queueing delay
-is added by SimPy's resource scheduling.
-
-### D4. NOC attachment points
-
-The NOC connects to all major cube-level components:
-
-```text
-                    UCIe-N (conn x4)
-                         |
-           +---------+---+---+---------+
-           |         |       |         |
-PE0.dma ---+  r0c0   |  ...  |  r0c5  +--- PE2.dma
-PE0.cpu <--+         |       |         +--< PE2.cpu
-           |         |       |         |
-UCIe-W ----+  ...    | [HBM] |  ...   +---- UCIe-E
-(conn x4)  |         | zone  |         |  (conn x4)
-           |  r2c0   |       |         |
-M_CPU <--->+         |       |         |
-           |  r3c0   |       |         |
-SRAM <---->+         |       |         |
-           |         |       |         |
-PE4.dma ---+  r4c0   |  ...  |  r4c5  +--- PE6.dma
-PE4.cpu <--+         |       |         +--< PE6.cpu
-           |         |       |         |
-           +---------+---+---+---------+
-                         |
-                    UCIe-S (conn x4)
-
-HBM attach: PE가 있는 라우터에 hbm_ctrl도 연결 (ADR-0019 D1)
-(xbar_top/xbar_bot은 ADR-0019에 의해 제거됨)
-```
-
-### D5. NOC edge bandwidths and distances
-
-| Connection | BW (GB/s) | Distance | Notes |
-| --- | --- | --- | --- |
-| PE_DMA -> NOC | 256.0 | Physical (PE pos) | Matches HBM slice BW |
-| NOC -> PE_CPU | - | 0.0 mm | Command path only |
-| Router <-> HBM_CTRL | 256.0 | 0.0 mm | Per PE router (ADR-0019) |
-| NOC <-> M_CPU | - | 0.0 mm | Command path |
-| NOC <-> SRAM | 128.0 x4 | 0.0 mm | 512 GB/s aggregate |
-| NOC <-> UCIe conn | 128.0 | 0.0 mm | Per connection, 4 per port |
-
-Distance 0.0 mm for most connections reflects the distributed nature of
-the NOC; the actual traversal distance is computed internally via Manhattan
-distance within the router grid.
-
-### D6. UCIe decomposition and inter-cube traffic
-
-Each cube has 4 UCIe ports (N, S, E, W). Each port is decomposed into:
-
-- 1 `ucie-{PORT}` node: UCIe protocol endpoint (overhead = 8.0 ns)
-- 4 `ucie-{PORT}.conn{0-3}` nodes: connection bridges between NOC and UCIe
-
-This decomposition enables N=4 independent NOC-to-UCIe connections per port,
-each with 128 GB/s bandwidth. Total aggregate per port: 512 GB/s.
-
-Inter-cube traffic path:
-
-```text
-Source: PE_DMA -> NOC -> conn{i} -> ucie-{PORT}
-                    [UCIe link: 512 GB/s, 1.0mm seam distance]
-Target: ucie-{PORT} -> conn{i} -> r{x}c{y} -> (mesh hops) -> hbm_ctrl
-```
-
-UCIe overhead (8.0 ns) is applied at each ucie-{PORT} node, so a
-full crossing incurs 16 ns (TX port + RX port).
-
-### D7. Data paths through the NOC
-
-**PE DMA to local HBM (same half):**
-
-```text
-PE_DMA -> r{x}c{y} -> hbm_ctrl  (local: 0 mesh hops, switching overhead only)
-```
-
-**PE DMA to remote PE's HBM:**
-
-```text
-PE_DMA -> r{x}c{y} -> (mesh hops) -> r{x'}c{y'} -> hbm_ctrl
-```
-
-**PE DMA to remote cube HBM:**
-
-```text
-PE_DMA -> r{x}c{y} -> conn -> ucie-E -> [seam] -> ucie-W -> conn -> r{x'}c{y'} -> hbm_ctrl
-```
-
-**Kernel Launch command to PE:**
-
-```text
-[from io_noc] -> ucie -> conn -> r{x}c{y} -> (mesh hops) -> M_CPU -> (mesh hops) -> PE_CPU
-```
-
-**Shared SRAM access:**
-
-```text
-PE_DMA -> r{x}c{y} -> (mesh hops) -> SRAM
-```
-
-### D8. Mesh generation
-
-The router grid is generated by `mesh_gen.py` based on:
-
-- `cube.pe_layout`: corner placement (NW, NE, SW, SE) and PEs per corner
-- `cube.geometry`: cube physical dimensions and HBM zone
-- `cube.ucie.n_connections`: determines router count for UCIe attachment
-
-The generator produces a `mesh_data` dictionary containing:
-
-- Router grid with positions and HBM exclusion zones
-- PE-to-router attachments (pe_dma, pe_cpu per PE)
-- UCIe-to-router attachments (N/S/E/W, distributed across edge routers)
-- M_CPU and SRAM router attachments
-- HBM attachment per PE router (ADR-0019)
-
-## Consequences
-
-- NOC provides position-aware routing with deterministic latency
-- Contention is captured per directed segment (not per-node)
-- All cube-internal traffic is explicitly routed through the NOC
-- HBM exclusion zone reflects physical die layout constraints
-- The mesh generation is fully parameterized by `topology.yaml`
-
-## Links
-
-- ADR-0003 D3 (cube-level NOC definition — extended by this ADR)
-- ADR-0004 D1 (PE DMA to local HBM path via router mesh)
-- ADR-0014 D1 (PE_DMA egress via router mesh)
-- ADR-0019 (NOC-Local HBM — xbar/bridge 제거, 명시적 라우터 mesh)
-- ADR-0015 D4 (fabric paths for Memory R/W and Kernel Launch)
-- ADR-0016 D1 (IOChiplet io_noc — analogous pattern at IO chiplet level)