From 168b0c89f0ced5b4b539a5e8f37bc3207edb3ce4 Mon Sep 17 00:00:00 2001 From: Yangwook Date: Wed, 20 May 2026 08:17:56 -0700 Subject: [PATCH] ADR: translate adr-ko/ to Korean, fix ADR-0013 slug, refine Status check Follow-up to the bilingual-structure commit: docs/adr-ko/ now holds only Korean versions (24 files translated from English placeholders), ADR-0013 slug uses kebab-case in both folders, and the verify tool allows translated parenthetical commentary in the Status block. - Translate 24 English files in docs/adr-ko/ to Korean. The previous bilingual-structure commit had left these as English copies because their source content was already English; this commit fulfills the policy that docs/adr-ko/ contains only Korean. - Rename ADR-0013 in both adr/ and adr-ko/ from ver-verification_strategy.md to ver-verification-strategy.md (kebab-case consistency with other ADRs). - CLAUDE.md (ADR Translation Discipline): clarify that only the Status lifecycle keyword (Accepted / Proposed / Stub / Draft / Superseded by ADR-NNNN / Merged into ADR-NNNN) must match across EN and KO; parenthetical commentary and trailing list items may be translated. - tools/verify_adr_lang_pairs.py: replace byte-equal Status check with normalize_status_keyword() which strips parenthetical commentary and takes only the first non-empty line. - tests/test_verify_adr_lang_pairs.py: update existing test names, add coverage for translated parenthetical, translated trailing list, and Superseded-by-NNNN keyword equality. Co-Authored-By: Claude Opus 4.7 (1M context) --- CLAUDE.md | 7 +- docs/adr-ko/ADR-0001-mem-physaddr-layout.md | 246 ++++---- docs/adr-ko/ADR-0002-lat-routing-distance.md | 138 ++-- .../ADR-0003-dev-target-system-hierarchy.md | 78 +-- ...ADR-0004-mem-memory-semantics-local-hbm.md | 100 +-- ...-0005-dev-diagram-views-distance-layout.md | 206 +++--- ...v-topology-compilation-distance-diagram.md | 146 ++--- .../ADR-0007-api-runtime-api-boundaries.md | 102 +-- ...R-0008-api-tensor-deploy-and-allocation.md | 94 +-- ...ADR-0009-api-kernel-execution-messaging.md | 164 +++-- .../ADR-0010-api-cli-surface-and-semantics.md | 162 +++-- ...11-mem-memory-addressing-simplification.md | 592 +++++++++--------- .../ADR-0012-api-host-io-message-schema.md | 188 +++--- .../ADR-0013-ver-verification-strategy.md | 145 +++++ ...DR-0014-dev-pe-pipeline-execution-model.md | 360 ++++++----- .../ADR-0015-dev-component-port-wire-model.md | 170 ++--- ...-0016-dev-iochiplet-noc-and-memory-path.md | 99 +-- ...-0017-dev-cube-noc-and-hbm-connectivity.md | 277 ++++---- .../ADR-0022-prog-program-id-2d-grid.md | 58 +- .../ADR-0032-algo-intercube-allreduce.md | 265 ++++---- .../ADR-0033-lat-latency-model-assumptions.md | 246 ++++---- ...0034-dev-hbm-controller-internal-design.md | 344 +++++----- ...dev-m-cpu-and-m-cpu-dma-component-model.md | 339 +++++----- .../ADR-0036-dev-io-cpu-component-model.md | 265 ++++---- .../ADR-0037-dev-forwarding-component.md | 261 ++++---- .../ADR-0013-ver-verification-strategy.md} | 0 .../adr/ADR-0013-ver-verification_strategy.md | 139 ---- tests/test_verify_adr_lang_pairs.py | 52 +- tools/verify_adr_lang_pairs.py | 39 +- 29 files changed, 2631 insertions(+), 2651 deletions(-) create mode 100644 docs/adr-ko/ADR-0013-ver-verification-strategy.md rename docs/{adr-ko/ADR-0013-ver-verification_strategy.md => adr/ADR-0013-ver-verification-strategy.md} (100%) delete mode 100644 docs/adr/ADR-0013-ver-verification_strategy.md diff --git a/CLAUDE.md b/CLAUDE.md index f48675c..1b2fa6f 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -402,8 +402,11 @@ Mechanics: - 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` block content must remain byte-identical between - the EN and KO files (e.g., both say `Accepted`). +- 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. diff --git a/docs/adr-ko/ADR-0001-mem-physaddr-layout.md b/docs/adr-ko/ADR-0001-mem-physaddr-layout.md index 5f33108..83b1735 100644 --- a/docs/adr-ko/ADR-0001-mem-physaddr-layout.md +++ b/docs/adr-ko/ADR-0001-mem-physaddr-layout.md @@ -1,10 +1,10 @@ -# ADR-0001: 51-bit Physical Address Layout & Decoding Contract +# ADR-0001: 51비트 물리 주소 레이아웃 및 디코딩 계약 ## 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.) +Accepted (Revision 2 — 2026-04-27: 구체적인 비트 레이아웃, rack_id 제거, +Tray->SIP / SIP->DIE 명칭 변경, PE/MCPU/IOCPU 서브 유닛 표. +ADR-0031을 대체함.) ## Date @@ -12,40 +12,39 @@ Supersedes ADR-0031.) ## Context -KernBench requires a stable, parsable physical address scheme that: +KernBench에는 다음과 같은 요건을 만족하는 안정적이고 파싱 가능한 물리 주소 체계가 필요하다. -- 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 +- 라우팅 도메인(SIP / die / HBM / PE-resource / IOCPU)으로 디코딩 가능 +- 토폴로지에 비의존적(개수를 하드코딩하지 않음) +- 교체 가능한 정책과 DI-first 컴포넌트를 지원 +- 다수의 SIP, AHBM die, IO chiplet die를 통합된 공간에서 다룸 -### 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. +- 최초 ADR-0001은 `rack_id(4) + sip_id(4) + sip_seg(5) + local_offset(38)` + 로 구성된 51비트 레이아웃을 정의했다. `rack_id`는 실제로 사용된 적이 없다. +- ADR-0031(스텁)은 PE-resource 범위 분할을 요청했으나 구현되지 않았다. -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. +Revision 2에서는 `rack_id`를 제거하고 `sip_seg`를 `die_id`로 개명하며, +PE, MCPU, CUBE_SRAM, IOCPU 리소스에 대한 구체적인 서브 유닛 표를 제공한다. +ADR-0031은 본 ADR로 대체된다. ## Decision -We define a **PhysAddr value object** and an **address decoding contract** -that converts an integer address into routing domains. +**PhysAddr 값 객체**와, 정수 주소를 라우팅 도메인으로 변환하는 +**주소 디코딩 계약**을 정의한다. -### D1. PhysAddr is an immutable value object +### D1. PhysAddr는 불변 값 객체이다 -- 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. +- PhysAddr는 불변이며 순수한 값으로 비교 가능하다. +- 모든 할당자는 **완전히 명세된 PhysAddr**(부분적인 메타데이터가 아님)를 반환한다. +- PhysAddr를 해석하기 위해 전역 상태를 필요로 해서는 안 된다. -### D2. 51-bit Physical Address Layout +### D2. 51비트 물리 주소 레이아웃 -A 51-bit physical address is adopted. +51비트 물리 주소를 채택한다. -#### 2.1 Top-Level Address Map +#### 2.1 최상위 주소 맵 ```text [50:47] sip_id (4) -- 16 SIPs @@ -60,17 +59,17 @@ A 51-bit physical address is adopted. +---------+----------+-------------------------+ ``` -#### 2.2 die_id Allocation +#### 2.2 die_id 할당 -| die_id | Meaning | +| die_id | 의미 | |--------|---------| | 0..15 | AHBM dies | | 16..20 | IOCHIPLET dies | | 21..31 | Reserved | -#### 2.3 AHBM Die Layout +#### 2.3 AHBM Die 레이아웃 -Only lower 256 GB of the 4 TB die-local window is assigned. +4 TB die-local 윈도우 중 하위 256 GB만 할당된다. ```text [41:38] MBZ (4) @@ -78,35 +77,35 @@ Only lower 256 GB of the 4 TB die-local window is assigned. [36: 0] sub-address (37) ``` -| addr_space | Meaning | +| addr_space | 의미 | |------------|---------| | 0 | Local resource | | 1 | HBM memory | -##### 2.3.1 HBM Window (addr_space = 1) +##### 2.3.1 HBM 윈도우 (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). +아키텍처상의 디코드 윈도우는 128 GB로 고정된다. 실제 구현 용량은 +SKU/토폴로지에 따라 더 작을 수 있다(D4 참조). -##### 2.3.2 Resource Window (addr_space = 0) +##### 2.3.2 Resource 윈도우 (addr_space = 0) ```text [36:34] resource_kind (3) [33: 0] kind_local (34) -- 16 GB per kind ``` -| resource_kind | Meaning | +| resource_kind | 의미 | |---------------|---------| | 000 | PE_LOCAL | | 001 | MCPU_LOCAL | | 010 | CUBE_SRAM | | 011..111 | Reserved | -Each kind gets a 16 GB decode region. +각 kind는 16 GB 디코드 영역을 갖는다. ##### 2.3.3 PE_LOCAL (resource_kind = 000) @@ -117,9 +116,9 @@ Each kind gets a 16 GB decode region. [24: 0] sub_offset (25) -- 32 MB per slot ``` -16 PEs x 16 sub-unit slots x 32 MB = 8 GB active decode. +16 PE x 16 서브 유닛 슬롯 x 32 MB = 8 GB 활성 디코드. -| pe_sub_unit | Name | Budget | +| pe_sub_unit | 이름 | 예산 | |-------------|------|--------| | 0 | PE_CPU_DTCM | 8 KB | | 1 | MATH_ENGINE_DTCM | 8 KB | @@ -138,9 +137,9 @@ Each kind gets a 16 GB decode region. [24: 0] sub_offset (25) -- 32 MB per slot ``` -1 GB active decode. +1 GB 활성 디코드. -| mcpu_sub_unit | Name | Budget | +| mcpu_sub_unit | 이름 | 예산 | |---------------|------|--------| | 0 | MCPU_ITCM | 512 KB | | 1 | MCPU_DTCM | 512 KB | @@ -157,32 +156,32 @@ Each kind gets a 16 GB decode region. [24: 0] sram_offset (25) -- flat 32 MB ``` -#### 2.4 IOCHIPLET Die Layout +#### 2.4 IOCHIPLET Die 레이아웃 -Only lower 1 TB of the 4 TB die-local window is assigned. +4 TB die-local 윈도우 중 하위 1 TB만 할당된다. ```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 +##### 2.4.1 IOCPU 영역 ```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. +16 x 128 MB 슬롯. 2 GB 활성 디코드. -| iocpu_sub_unit | Name | Budget | +| iocpu_sub_unit | 이름 | 예산 | |----------------|------|--------| | 0 | IOCPU_ITCM | 512 KB | | 1 | IOCPU_DTCM | 512 KB | @@ -192,110 +191,107 @@ Region split by address range: | 5 | IO_SRAM | 64 MB | | 6..15 | Reserved | -- | -##### 2.4.2 UAL Region +##### 2.4.2 UAL 영역 -Sub-layout TBD (separate ADR). +서브 레이아웃은 별도 ADR에서 정의한다(TBD). -#### 2.5 Addressing Rules +#### 2.5 주소 지정 규칙 -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). +1. MBZ 비트는 반드시 0이어야 한다. MBZ 비트가 0이 아닌 주소는 + **아키텍처적으로 유효하지 않다**. 구현체는 디코드 폴트를 발생시키거나 + 오류를 반환할 수 있다 — 본 ADR은 동작을 규정하지 않는다. +2. 단순한 하드웨어 디코드를 위해 고정된 슬롯 크기를 채택한다. 실제 구현 + 용량은 슬롯보다 작을 수 있다. +3. 슬롯 내에서 서브 유닛의 구현 예산을 초과하는 접근은 **아키텍처적으로 + 유효하지 않다**(MBZ와 동일한 정책). -### D3. Bitfield decoding is deterministic +### D3. 비트필드 디코딩은 결정론적이다 -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). +정수 주소가 주어지면 필드 추출(`sip_id`, `die_id`, `kind`, `sub_unit`, +`offset`)은 순수하게 위치 기반이다. 런타임 상태가 필요하지 않다. +디코딩은 정수 주소를 결정론적으로 목적지 도메인(`sip_id`, `die_id`, +타깃 종류 HBM / PE_LOCAL / MCPU_LOCAL / CUBE_SRAM / IOCPU / UAL)으로 매핑한다. -### D4. Capacity validation may depend on topology config +### D4. 용량 검증은 토폴로지 설정에 의존할 수 있다 -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. +디코딩된 주소가 **구현된 용량** 안에 들어가는지(예: 특정 SKU의 HBM 96 GB)는 +DI/설정을 통해 제공된 토폴로지 파라미터로 검증한다. 디코딩 자체(D3)는 +토폴로지를 참조하지 않으며 — 검증 단계에서만 참조한다. 이러한 파라미터는 +컴포넌트 구현이 아니라 토폴로지/설정 레이어에 존재해야 한다. -### D5. Routing consumes decoded domains, not raw bits +### D5. 라우팅은 원시 비트가 아닌 디코딩된 도메인을 소비한다 -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 +- `src` 위치 (sip / die / pe 또는 node_id) +- PhysAddr 디코딩에서 도출된 `dst` 도메인 +- 크기 인지 링크 레이턴시를 위한 `size_bytes` -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). +1. **`rack_id`(4비트) 유지**: 기각 — 실제로 사용된 적이 없으며, 4비트를 + 소비함으로써 die-local 확장을 42비트(IOCHIPLET 1 TB)까지 가능하게 하는 + 기회를 막는다. -2. **Uniform 256 GB per die**: Rejected -- IOCHIPLET UAL requires ~1 TB. - Freed rack_id bits enable 42-bit local_offset. +2. **die당 256 GB로 균일화**: 기각 — IOCHIPLET UAL은 약 1 TB가 필요하다. + 해제된 rack_id 비트를 활용하여 42비트 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. +3. **가변 폭 die 윈도우(AHBM 256 GB, CHIPLET 1 TB를 다중 seg 스패닝으로 구현)**: + 기각 — D3(결정론적 디코딩)를 복잡하게 만든다. MBZ 패딩을 갖는 균일한 + 4 TB 윈도우가 더 단순하다. -4. **Use raw integers everywhere, decode ad-hoc in routing**: Rejected -- - leads to duplicated logic, inconsistent routing, and hidden - assumptions. +4. **모든 곳에서 원시 정수를 사용하고, 라우팅에서 임시로 디코딩**: 기각 — + 로직이 중복되고 라우팅이 일관성을 잃으며 가정이 숨겨진다. -5. **Hardcode topology sizes (SIP/CUBE/PE counts) into decoding**: - Rejected -- violates SPEC R3 and breaks swappability. +5. **토폴로지 크기(SIP/CUBE/PE 개수)를 디코딩에 하드코딩**: 기각 — + SPEC R3를 위반하고 교체 가능성을 깬다. -6. **Put decoding inside memory controllers or routers**: Rejected -- - leaks policy into components, violates SPEC R4 / D5. +6. **디코딩을 메모리 컨트롤러나 라우터 내부에 둠**: 기각 — 정책이 컴포넌트로 + 누출되며 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). +- 단순한 계층적 디코더: SIP -> die -> kind -> 서브 유닛. +- 메모리(HBM)와 로컬 리소스(PE/MCPU/SRAM/IOCPU)의 깔끔한 분리. +- 결정론적 라우팅 도메인은 명확한 테스트 불변식을 가능하게 한다(SPEC R1, R5). +- 확장 가능: 11개의 예약된 die_id 슬롯, 예약된 resource_kind / 서브 유닛 + 슬롯, 예약된 MBZ 비트. +- DI-first: 컴포넌트를 변경하지 않고도 디코더를 교체할 수 있다(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. +- power-of-2 슬롯 정렬로 인한 희소한 주소 공백. +- 큰 예약/MBZ 영역(향후 확장을 위해 의도된 것). +- 토폴로지에서 유도된 크기에 대해 명시적인 설정이 필요하다(D4). +- 안정적이고 잘 테스트된 상태로 유지되어야 하는 단일 "정통" 디코딩 모듈이 + 도입된다. -## 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. +- **ADR-0031 (PhysAddr PE-Resource Extension)**: 스텁 상태였음. D2.3.3-D2.3.5의 + PE_LOCAL / MCPU_LOCAL / CUBE_SRAM 서브 유닛 표가 ADR-0031에서 제시한 + 목표를 충족한다. -## 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`. +- 권장 모듈: `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`. -## Appendix A. Address Examples +## 부록 A. 주소 예시 -### A.1 AHBM HBM access +### A.1 AHBM HBM 접근 sip=2, die=5, HBM offset=0x1000 @@ -347,7 +343,7 @@ chiplet_offset = (2 << 27) | 0x20000 (< 0x8000_0000 -> IOCPU region) ``` -### A.5 IOCHIPLET -- UAL region, offset=4 GB +### A.5 IOCHIPLET -- UAL 영역, offset=4 GB ```text sip_id = 0 -> [50:47] = 0 @@ -355,7 +351,7 @@ 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) diff --git a/docs/adr-ko/ADR-0002-lat-routing-distance.md b/docs/adr-ko/ADR-0002-lat-routing-distance.md index 19849f9..c3e922a 100644 --- a/docs/adr-ko/ADR-0002-lat-routing-distance.md +++ b/docs/adr-ko/ADR-0002-lat-routing-distance.md @@ -1,4 +1,4 @@ -# ADR-0002: Routing Distance, Ordering & Bypass Rules +# ADR-0002: 라우팅 거리, 순서 및 우회 규칙 ## Status Accepted @@ -7,96 +7,94 @@ Accepted 2026-02-27 ## Context -The KernBench Graph Latency Simulator must compare kernel execution time -across different architectures and topologies by computing end-to-end -latency from graph traversal. +KernBench Graph Latency Simulator는 서로 다른 아키텍처·토폴로지에 대한 +커널 실행 시간을 비교해야 하며, 그래프 순회로부터 end-to-end 레이턴시를 +계산하여 이를 달성한다. -To support meaningful comparison: -- routing must be deterministic -- latency must reflect actual interconnect structure -- local vs remote traffic must be distinguishable -- “bypass” optimizations must not undermine debuggability or correctness +의미 있는 비교를 지원하려면: +- 라우팅이 결정론적이어야 한다 +- 레이턴시가 실제 인터커넥트 구조를 반영해야 한다 +- 로컬과 리모트 트래픽이 구분 가능해야 한다 +- "우회(bypass)" 최적화가 디버깅 가능성이나 정확성을 훼손해서는 안 된다 -The simulator also aims to avoid software-managed metadata and hidden -shortcuts that obscure control paths. +또한 시뮬레이터는 소프트웨어가 관리하는 메타데이터 및 제어 경로를 +가리는 숨겨진 지름길을 피하는 것을 목표로 한다. ## Decision -### D1. Distance is accumulated latency, not hop count -- Routing “distance” is defined as the **sum of per-node and per-link latency**. -- Hop count alone must not be used for ordering or path selection. -- Size-aware serialization latency (bytes / BW) contributes to distance. +### D1. 거리(distance)는 hop 수가 아니라 누적 레이턴시이다 +- 라우팅 "거리"는 **노드별·링크별 레이턴시의 합**으로 정의된다. +- 순서 결정이나 경로 선택에 hop 수만을 사용해서는 안 된다. +- 크기 인지(size-aware) 직렬화 레이턴시(bytes / BW)가 거리에 기여한다. -### D2. Routing order is derived from graph traversal -- The chosen route is the path with minimum accumulated latency - given the constructed graph and routing policy. -- Deterministic ordering must be guaranteed for identical inputs - (topology + policy + request). +### D2. 라우팅 순서는 그래프 순회에서 유도된다 +- 선택된 경로는 구성된 그래프와 라우팅 정책 하에서 + 누적 레이턴시가 최소인 경로이다. +- 동일 입력(토폴로지 + 정책 + 요청)에 대해 결정론적 순서가 보장되어야 한다. -### 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-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. -- Implicit or “magic” bypass paths are disallowed. +### D3. 우회는 명시적이며 그래프로 표현된다 +- 모든 경로는 그래프에 명시적으로 표현되며 레이턴시 누적의 대상이 되어야 한다. +- 예: PE_DMA는 NOC 라우터 메시(ADR-0017 D7)에 연결된다. 모든 목적지 + (HBM, 공유 SRAM, 큐브 간 UCIe)는 명시적 메시 hop을 통해 도달한다. + 로컬 HBM 접근은 hop 수가 최소(스위칭 오버헤드만)이며, 리모트 접근은 + 추가 라우터를 거친다. +- 암묵적이거나 "마법 같은" 우회 경로는 금지된다. -### D4. No zero-latency end-to-end paths +### D4. end-to-end 레이턴시가 0인 경로는 없다 -- Every routed request must incur **end-to-end** latency > 0. -- Individual fabric segments (e.g., NOC hops) MAY have distance_mm = 0 - when the fabric is distributed and distance is not meaningful at that granularity. - This is allowed because other components on the same path (e.g., PE_DMA, SRAM, - UCIe endpoints) contribute non-zero latency, ensuring the end-to-end invariant holds. -- Fully zero-latency end-to-end paths are disallowed, except for explicit - test-only stubs clearly marked as such. +- 모든 라우팅 요청은 **end-to-end** 레이턴시가 > 0이어야 한다. +- 개별 패브릭 세그먼트(예: NOC hop)는 패브릭이 분산되어 있고 해당 granularity에서 + 거리가 의미가 없을 때 distance_mm = 0을 가질 수 있다. + 이는 같은 경로상의 다른 컴포넌트(예: PE_DMA, SRAM, UCIe 엔드포인트)가 + 0이 아닌 레이턴시에 기여하여 end-to-end 불변성을 유지하므로 허용된다. +- end-to-end가 완전히 0 레이턴시인 경로는 금지된다. 단, 명시적으로 + 표시된 테스트 전용 stub만 예외이다. -### D5. Policy vs topology responsibility split -- Topology builder: - - defines nodes and links and their latency/BW parameters -- Routing policy: - - selects among available graph paths based on decoded domains -- Routing policy must not assume missing links; missing connectivity - is a topology construction error. +### D5. 정책과 토폴로지의 책임 분리 +- 토폴로지 빌더: + - 노드와 링크 및 그들의 레이턴시/BW 파라미터를 정의한다 +- 라우팅 정책: + - 디코딩된 도메인을 바탕으로 사용 가능한 그래프 경로 중에서 선택한다 +- 라우팅 정책은 누락된 링크를 가정해서는 안 된다. 누락된 연결성은 + 토폴로지 구성 오류이다. -### D6. No software-managed routing metadata -- Routing decisions must not rely on per-request software-managed metadata - that tracks distance, hop count, or ordering outside the graph model. -- All distance/order computation is derived from traversal itself. +### D6. 소프트웨어 관리 라우팅 메타데이터 금지 +- 라우팅 결정은 그래프 모델 외부에서 거리·hop 수·순서를 추적하는 + 요청별 소프트웨어 관리 메타데이터에 의존해서는 안 된다. +- 모든 거리·순서 계산은 순회 자체에서 유도된다. ## Alternatives Considered -1) **Hop-count based routing** -- Rejected: ignores heterogeneous latency/BW and misrepresents - architectural differences. +1) **Hop 수 기반 라우팅** +- 기각: 이질적인 레이턴시·BW를 무시하고 아키텍처 차이를 잘못 표현한다. -2) **Implicit local shortcuts** -- Rejected: breaks debuggability and violates traversal-based latency. +2) **암묵적 로컬 지름길** +- 기각: 디버깅 가능성을 해치고 순회 기반 레이턴시 원칙을 위반한다. -3) **Software-managed distance metadata** -- Rejected: increases control overhead and obscures routing semantics. +3) **소프트웨어 관리 거리 메타데이터** +- 기각: 제어 오버헤드를 증가시키고 라우팅 시맨틱을 모호하게 만든다. ## Consequences -### Positive -- Clear, debuggable hop-by-hop traces (SPEC R2, R4). -- Architecture comparisons reflect real interconnect structure. -- Routing behavior is reproducible and deterministic. +### 긍정적 +- 명확하고 디버깅 가능한 hop-by-hop 트레이스 (SPEC R2, R4). +- 아키텍처 비교가 실제 인터커넥트 구조를 반영한다. +- 라우팅 동작이 재현 가능하고 결정론적이다. -### Tradeoffs / Costs -- Graph construction must be correct and complete. -- Bypass modeling requires explicit graph representation, - which slightly increases topology description complexity. +### 트레이드오프 / 비용 +- 그래프 구성이 정확하고 완전해야 한다. +- 우회 모델링이 명시적 그래프 표현을 요구하므로 토폴로지 기술이 + 약간 더 복잡해진다. ## Implementation Notes (Non-normative) -- Recommended responsibilities: - - Graph builder: ensure all required paths exist. - - Router: select next hop based on decoded domains and policy. -- Tests should assert: - - non-zero end-to-end latency - - deterministic routing for identical inputs - - bypass paths appear explicitly in emitted traces +- 권장 책임 분담: + - 그래프 빌더: 필요한 모든 경로가 존재함을 보장. + - 라우터: 디코딩된 도메인과 정책을 바탕으로 다음 hop 선택. +- 테스트가 검증해야 할 항목: + - end-to-end 레이턴시 > 0 + - 동일 입력에 대한 결정론적 라우팅 + - 우회 경로가 출력 트레이스에 명시적으로 나타남 ## Links -- SPEC.md: R1 (routing), R2 (latency), R3 (topology), R5 (multi-domain comm) -- ADR-0001: PhysAddr layout & decoding contract +- SPEC.md: R1 (라우팅), R2 (레이턴시), R3 (토폴로지), R5 (다중 도메인 통신) +- ADR-0001: PhysAddr 레이아웃 및 디코딩 계약 diff --git a/docs/adr-ko/ADR-0003-dev-target-system-hierarchy.md b/docs/adr-ko/ADR-0003-dev-target-system-hierarchy.md index e5acc7d..d5cd5d1 100644 --- a/docs/adr-ko/ADR-0003-dev-target-system-hierarchy.md +++ b/docs/adr-ko/ADR-0003-dev-target-system-hierarchy.md @@ -1,4 +1,4 @@ -# ADR-0003: Target System Hierarchy & Modeling Scope +# ADR-0003: 타겟 시스템 계층 및 모델링 범위 ## Status @@ -6,63 +6,63 @@ Accepted ## Context -We need a system-level simulator to evaluate LLM kernel performance on our AI Accelerator platform. -The platform is organized as a compute tray containing multiple identical SIPs connected via PCIe or UAL -through switching fabrics, with a host CPU issuing commands/kernels. +자사 AI Accelerator 플랫폼에서 LLM 커널 성능을 평가하기 위해 시스템 수준의 시뮬레이터가 필요하다. +해당 플랫폼은 PCIe 또는 UAL을 통해 스위칭 패브릭으로 연결된 다수의 동일한 SIP를 포함하는 컴퓨트 트레이로 구성되며, +호스트 CPU가 명령/커널을 발급한다. ## Decision -We model the system hierarchy explicitly: +시스템 계층을 다음과 같이 명시적으로 모델링한다. ### D1. Tray-level -- A compute tray contains: - - Host CPU (issues requests / coordinates runtime & data placement) - - Multiple identical SIPs (accelerators) - - Interconnect fabric between SIPs (PCIe and/or UAL via switches) +- 하나의 컴퓨트 트레이는 다음을 포함한다: + - 호스트 CPU (요청 발급 / 런타임 및 데이터 배치 조정) + - 다수의 동일한 SIP (가속기) + - SIP 간 인터커넥트 패브릭 (스위치를 통한 PCIe 및/또는 UAL) ### D2. SIP-level -- A SIP is a multi-die package composed of: - - Multiple CUBEs (HBM die + compute PEs + UCIe) - - One or more IO chiplets (host/SIP interfaces) -- IO chiplets: - - provide interfaces: PCIe-EP, IO_CPU, optionally UAL-EP - - can be multiple per SIP - - placement constrained to SIP shoreline (top/bottom/left/right); each shoreline may host 1–2 IO chiplets +- SIP는 다음으로 구성된 멀티 다이 패키지이다: + - 다수의 CUBE (HBM 다이 + 컴퓨트 PE + UCIe) + - 하나 이상의 IO 칩렛 (호스트/SIP 인터페이스) +- IO 칩렛: + - 다음 인터페이스를 제공한다: PCIe-EP, IO_CPU, 선택적으로 UAL-EP + - SIP 당 다수가 존재할 수 있다 + - 배치는 SIP shoreline(상/하/좌/우)으로 제약되며, 각 shoreline에는 1~2개의 IO 칩렛이 위치할 수 있다 ### D3. CUBE-level -- A CUBE contains: - - HBM + memory controller (HBM_CTRL) - - 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 - - up to 4 UCIe endpoints (N/E/W/S) for CUBE↔CUBE and CUBE↔IO connectivity +- 하나의 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 -- A PE can execute one kernel instance -- PE contains internal control + accelerators (modeled at PE view granularity): - - PE_CPU, command handler, PE_TCM, DMA/GEMM/MATH engines, internal queues +- 하나의 PE는 하나의 커널 인스턴스를 실행할 수 있다 +- PE는 내부 제어 + 가속기를 포함한다 (PE 뷰 단위로 모델링): + - PE_CPU, 명령 핸들러, PE_TCM, DMA/GEMM/MATH 엔진, 내부 큐 ## Consequences -- The simulator supports abstraction by “views”: - - SIP view hides PE internals - - CUBE view treats each PE as a single block - - PE view expands PE internals -- Topology remains parameterized; sizes/counts/links come from configuration. +- 시뮬레이터는 "뷰" 단위의 추상화를 지원한다: + - SIP 뷰는 PE 내부를 숨긴다 + - CUBE 뷰는 각 PE를 단일 블록으로 다룬다 + - PE 뷰는 PE 내부를 전개한다 +- 토폴로지는 매개변수화된 상태로 유지되며, 크기/개수/링크는 설정으로부터 주어진다. ## Links - SPEC R3/R5 -- ADR-0005 (diagram views) -- ADR-0017 (cube NOC 2D mesh architecture) +- ADR-0005 (다이어그램 뷰) +- ADR-0017 (큐브 NoC 2D 메시 아키텍처) diff --git a/docs/adr-ko/ADR-0004-mem-memory-semantics-local-hbm.md b/docs/adr-ko/ADR-0004-mem-memory-semantics-local-hbm.md index d9144b0..855e84d 100644 --- a/docs/adr-ko/ADR-0004-mem-memory-semantics-local-hbm.md +++ b/docs/adr-ko/ADR-0004-mem-memory-semantics-local-hbm.md @@ -1,4 +1,4 @@ -# ADR-0004: Memory Semantics & Local-HBM Bandwidth Guarantee +# ADR-0004: 메모리 시맨틱 및 로컬 HBM 대역폭 보장 ## Status @@ -6,71 +6,73 @@ Accepted ## Context -Accurately modeling PE↔HBM behavior is essential for kernel latency estimation. -Each PE has a notion of “local HBM” that must guarantee full HBM bandwidth, independent of intervening on-die fabric bandwidth. +PE↔HBM 동작을 정확하게 모델링하는 것은 커널 레이턴시 추정에 필수적이다. +각 PE는 "로컬 HBM"이라는 개념을 가지며, 이는 중간 온칩 패브릭 대역폭과 +무관하게 HBM 전체 대역폭을 보장해야 한다. ## Decision -### D1. Local HBM definition +### D1. 로컬 HBM의 정의 -- 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-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. +- 각 PE에는 논리적으로 정의된 "로컬 HBM" 영역이 할당된다. +- 로컬 HBM은 NOC 메시(ADR-0017 D4) 내에서 해당 PE의 라우터에 직접 연결된 + pseudo-channel 부분집합에 대응한다. +- 경로는: PE_DMA → 로컬 라우터 → HBM_CTRL (스위칭 오버헤드만, 메시 hop 0개). +- 매핑(HBM pseudo-channel → PE 로컬 영역)은 토폴로지 구성에서 유도된다. -### D2. Local HBM bandwidth guarantee contract +### D2. 로컬 HBM 대역폭 보장 계약 -- Accesses from a PE to its local HBM MUST guarantee full effective HBM - read/write bandwidth independent of intervening fabric bandwidth limits. -- Effective HBM bandwidth = spec bandwidth x efficiency factor. - The efficiency factor (configured via `hbm_ctrl.attrs.efficiency`, default 0.8) - models real-world DRAM inefficiencies (refresh cycles, bank conflicts, page - misses). For example: 256 GB/s spec x 0.8 = 204.8 GB/s effective. -- The topology builder applies the efficiency factor to router-to-hbm edge - bandwidth at graph construction time, so all downstream routing and latency - computation uses the effective value. -- 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. +- 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. Remote PE HBM semantics (intra-cube) +### D3. 리모트 PE HBM 시맨틱 (큐브 내) -- 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. +- 한 PE가 다른 PE의 로컬 HBM에 접근할 때는 NOC를 거친다: + - PE_DMA → NOC → (패브릭 hop) → 대상 PE의 NOC 포트 → HBM_CTRL +- NOC의 대역폭과 hop 수에 의해 리모트 HBM 접근이 로컬 접근 대비 제한될 수 있다. -### D4. Non-local HBM semantics (inter-cube / inter-SIP) +### D4. 비로컬 HBM 시맨틱 (큐브 간 / SIP 간) -- Accesses from a PE to HBM in a different cube or SIP MAY be limited by: - - NOC bandwidth within the cube, - - inter-cube UCIe links, - - inter-SIP fabric (PCIe/UAL). -- These paths MUST be explicit and traceable. +- PE에서 다른 큐브나 SIP에 있는 HBM으로의 접근은 다음에 의해 제한될 수 있다: + - 큐브 내 NOC 대역폭, + - 큐브 간 UCIe 링크, + - SIP 간 패브릭 (PCIe/UAL). +- 이 경로들은 명시적이고 추적 가능해야 한다. -### D5. Shared SRAM semantics +### D5. 공유 SRAM 시맨틱 -- Each CUBE contains a shared SRAM accessible by all PEs in that CUBE. -- Access path: PE_DMA → NOC → shared SRAM. -- Shared SRAM bandwidth is limited by the NOC↔SRAM link bandwidth. -- Shared SRAM is not part of the HBM address space; it is a separate memory domain. +- 각 CUBE는 해당 CUBE의 모든 PE가 접근 가능한 공유 SRAM을 포함한다. +- 접근 경로: PE_DMA → NOC → 공유 SRAM. +- 공유 SRAM의 대역폭은 NOC↔SRAM 링크 대역폭으로 제한된다. +- 공유 SRAM은 HBM 주소 공간의 일부가 아니라 별도의 메모리 도메인이다. ## Verification Notes -Tests should cover: +테스트가 다뤄야 할 케이스: -- local-HBM case: BW matches HBM BW regardless of fabric BW parameter -- remote PE HBM case: latency includes mesh hop traversal -- non-local cases (inter-cube/inter-SIP): BW/latency respond to fabric/link parameters -- shared SRAM case: access via NOC with correct BW +- 로컬 HBM 케이스: 패브릭 BW 파라미터와 무관하게 대역폭이 HBM 대역폭과 일치 +- 리모트 PE HBM 케이스: 레이턴시가 메시 hop 순회를 포함 +- 비로컬 케이스(큐브 간/SIP 간): 패브릭/링크 파라미터에 대역폭·레이턴시가 반응 +- 공유 SRAM 케이스: NOC 경유 접근이 올바른 대역폭으로 수행됨 ## Links - SPEC R2/R5 -- ADR-0002 (distance/order & explicit bypass) -- ADR-0017 D7 (PE DMA data paths through NOC to HBM) +- ADR-0002 (거리/순서 및 명시적 우회) +- ADR-0017 D7 (NOC를 통한 PE DMA → HBM 데이터 경로) diff --git a/docs/adr-ko/ADR-0005-dev-diagram-views-distance-layout.md b/docs/adr-ko/ADR-0005-dev-diagram-views-distance-layout.md index 6391f19..986c41c 100644 --- a/docs/adr-ko/ADR-0005-dev-diagram-views-distance-layout.md +++ b/docs/adr-ko/ADR-0005-dev-diagram-views-distance-layout.md @@ -1,4 +1,4 @@ -# ADR-0005: Diagram Views & Distance-Aware Layout Rules +# ADR-0005: 다이어그램 뷰 및 거리 기반 레이아웃 규칙 ## Status @@ -6,17 +6,17 @@ Accepted ## Context -We require verifiable and inspectable system modeling for a large-scale, -parameterized AI Accelerator system. +대규모, 매개변수화된 AI Accelerator 시스템에 대해 검증 가능하고 점검 가능한 +시스템 모델링이 필요하다. -Humans must be able to: +사람이 다음을 할 수 있어야 한다: -- visually inspect the modeled topology, -- reason about communication structure and relative distance, -- do so at multiple abstraction levels without being overwhelmed by detail. +- 모델링된 토폴로지를 시각적으로 점검하고, +- 통신 구조와 상대적 거리에 대해 추론하고, +- 세부 사항에 압도되지 않으면서 여러 추상화 수준에서 이를 수행한다. -The simulator models distance (accumulated latency) as a first-class concept. -Diagrams must reflect this distance by default. +시뮬레이터는 거리(누적 레이턴시)를 1급 개념(first-class concept)으로 모델링한다. +다이어그램은 기본적으로 이 거리를 반영해야 한다. --- @@ -24,163 +24,163 @@ Diagrams must reflect this distance by default. ### D1. Global Defaults -- All diagrams MUST be **distance-aware by default**. -- All diagrams MUST render **representative views** of the architecture. -- Instance indices (e.g., sip0, cube2, pe3) MUST NOT be required for diagram generation. -- Instance indices MAY be used ONLY: - - to define a distance anchor in asymmetric or debugging scenarios, or - - when explicitly requested. +- 모든 다이어그램은 기본적으로 **거리 인식(distance-aware)** 이어야 한다. +- 모든 다이어그램은 아키텍처의 **대표 뷰(representative view)** 를 렌더링해야 한다. +- 인스턴스 인덱스(예: sip0, cube2, pe3)는 다이어그램 생성에 필수가 아니어야 한다. +- 인스턴스 인덱스는 다음의 경우에만 사용될 수 있다: + - 비대칭 또는 디버깅 시나리오에서 거리 앵커를 정의하기 위한 경우, 또는 + - 명시적으로 요청된 경우. --- ### D2. Representative Rendering Rule -- All CUBEs share the same internal structure. -- All PEs share the same internal structure. +- 모든 CUBE는 동일한 내부 구조를 공유한다. +- 모든 PE는 동일한 내부 구조를 공유한다. -Therefore: +따라서: -- SIP-level diagrams render representative CUBEs and IO chiplets. -- CUBE-level diagrams render representative PEs as opaque blocks. -- PE-level diagrams render a representative PE with fully expanded internals. +- SIP 수준 다이어그램은 대표 CUBE와 IO 칩렛을 렌더링한다. +- CUBE 수준 다이어그램은 대표 PE를 불투명 블록으로 렌더링한다. +- PE 수준 다이어그램은 내부가 완전히 전개된 대표 PE를 렌더링한다. -Diagrams MUST NOT depend on specific SIP, CUBE, or PE indices -unless explicitly requested. +다이어그램은 명시적으로 요청되지 않는 한 +특정 SIP, CUBE, 또는 PE 인덱스에 의존해서는 안 된다. --- ### D3. Diagram Views -#### View A — SIP-Level Diagram +#### View A — SIP 수준 다이어그램 -**Purpose** -Explain system-scale structure and connectivity. +**목적** +시스템 규모의 구조와 연결성을 설명한다. -**Visible elements** +**가시 요소** -- SIP boundaries (optional) -- CUBEs (opaque blocks) -- IO chiplets (opaque blocks) -- Optional UCIe stubs only if needed to clarify connectivity +- SIP 경계 (선택사항) +- CUBE (불투명 블록) +- IO 칩렛 (불투명 블록) +- 연결성 명확화에 필요한 경우에만 선택적 UCIe 스텁 -**Hidden elements** +**비가시 요소** -- PE internals -- CUBE internal fabric -- IO chiplet internals +- PE 내부 +- CUBE 내부 패브릭 +- IO 칩렛 내부 -**Visible links** +**가시 링크** -- Host ↔ IO chiplets (PCIe) -- SIP ↔ SIP (PCIe / UAL via switches) -- IO ↔ CUBE (on-package links) +- 호스트 ↔ IO 칩렛 (PCIe) +- SIP ↔ SIP (스위치를 통한 PCIe / UAL) +- IO ↔ CUBE (온패키지 링크) --- -#### View B — CUBE-Level Diagram +#### View B — CUBE 수준 다이어그램 -**Purpose** -Explain cube-internal structure and data/control flow. +**목적** +큐브 내부 구조와 데이터/제어 흐름을 설명한다. -**Visible elements** +**가시 요소** -- Router mesh: 2D grid of NOC routers (from cube_mesh.yaml), all traffic routes through mesh -- HBM_CTRL attached to PE routers (local HBM = 0 hop) -- HBM subsystem (HBM_CTRL) -- Shared SRAM: cube-level shared memory -- Management CPU (M_CPU) -- PEs as opaque blocks (PE[0..N−1]) -- UCIe endpoints (N/E/W/S) as ports +- 라우터 메시: 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) -**Hidden elements** +**비가시 요소** -- PE internals +- PE 내부 -**Visible links** +**가시 링크** -- PE → router (HBM + non-HBM data path via mesh) -- Router ↔ HBM_CTRL (local HBM access) -- Router ↔ Router (mesh hops for remote access) -- Router ↔ UCIe endpoints -- Router ↔ shared SRAM -- M_CPU ↔ router (command path) -- Router → PE_CPU (command delivery, collapsed into PE block) +- PE → 라우터 (메시를 통한 HBM + 비-HBM 데이터 경로) +- 라우터 ↔ HBM_CTRL (로컬 HBM 액세스) +- 라우터 ↔ 라우터 (원격 액세스를 위한 메시 홉) +- 라우터 ↔ UCIe 엔드포인트 +- 라우터 ↔ 공유 SRAM +- M_CPU ↔ 라우터 (명령 경로) +- 라우터 → PE_CPU (명령 전달, PE 블록 내부로 축약됨) --- -#### View C — PE-Level Diagram +#### View C — PE 수준 다이어그램 -**Purpose** -Explain internal PE behavior and execution structure. +**목적** +PE 내부 동작과 실행 구조를 설명한다. -**Visible elements** +**가시 요소** - PE_CPU -- Command handler / scheduler -- PE_TCM (local SRAM) -- HW accelerators (DMA, GEMM, MATH, etc.) -- Local HBM interface -- Optional IPCQ / messaging endpoints +- 명령 핸들러 / 스케줄러 +- PE_TCM (로컬 SRAM) +- HW 가속기 (DMA, GEMM, MATH 등) +- 로컬 HBM 인터페이스 +- 선택적 IPCQ / 메시징 엔드포인트 -**Visible links** +**가시 링크** -- Control paths (CPU → scheduler → engines) -- Data paths (engines ↔ TCM, DMA ↔ local HBM) -- External fabric ports as abstract ports only +- 제어 경로 (CPU → 스케줄러 → 엔진) +- 데이터 경로 (엔진 ↔ TCM, DMA ↔ 로컬 HBM) +- 외부 패브릭 포트는 추상 포트로만 표현 --- -### D4. Distance-Aware Layout (Default) +### D4. 거리 기반 레이아웃 (기본) -#### Distance definition +#### 거리 정의 -- Distance is defined as **accumulated latency**, consistent with ADR-0002. -- Distance is computed from a single anchor node. +- 거리는 ADR-0002와 정합되도록 **누적 레이턴시(accumulated latency)** 로 정의된다. +- 거리는 단일 앵커 노드로부터 계산된다. -#### Default anchor selection +#### 기본 앵커 선택 -- SIP view: IO chiplet (or Host CPU if present) -- CUBE view: a representative PE -- PE view: PE_CPU or Command Handler +- SIP 뷰: IO 칩렛 (또는 존재한다면 호스트 CPU) +- CUBE 뷰: 대표 PE +- PE 뷰: PE_CPU 또는 명령 핸들러 -Anchors are **implicit defaults** and MUST NOT be required to be specified. +앵커는 **암묵적 기본값**이며, 지정이 강제되어서는 안 된다. -#### Layout rules +#### 레이아웃 규칙 -- Diagrams MUST be laid out in layers based on distance buckets. -- Layout direction MUST be consistent within a view type - (preferred: left-to-right). -- Nodes with equal distance MUST have stable ordering - (by role or identifier, deterministically). +- 다이어그램은 거리 버킷에 기반한 레이어로 배치되어야 한다. +- 레이아웃 방향은 뷰 유형 내에서 일관되어야 한다 + (선호: 좌→우). +- 동일 거리의 노드는 결정론적으로 안정된 순서를 가져야 한다 + (역할 또는 식별자 기준). -Cycles MAY be rendered using dashed or curved edges for readability, -without affecting distance semantics. +가독성을 위해 사이클은 점선 또는 곡선 엣지로 렌더링될 수 있으며, +이는 거리 의미에 영향을 주지 않는다. --- -### D5. Generation Contract (for Tools / Claude Code) +### D5. 생성 컨트랙트 (도구 / Claude Code용) -When generating diagrams: +다이어그램 생성 시: -- Assume distance-aware layout by default. -- Assume representative rendering by default. -- Do NOT ask for SIP/CUBE/PE indices unless required. -- Do NOT expand hidden abstraction levels. -- Prefer architectural clarity over micro-hop fidelity. +- 기본적으로 거리 기반 레이아웃을 가정한다. +- 기본적으로 대표 렌더링을 가정한다. +- 필요한 경우가 아니면 SIP/CUBE/PE 인덱스를 묻지 않는다. +- 숨겨진 추상화 수준을 전개하지 않는다. +- 마이크로 홉의 정밀도보다 아키텍처적 명확성을 우선한다. --- ## Consequences -- Diagrams are stable across topology scaling. -- Changes in distance or routing policy are reflected visually. -- Diagrams serve as verifiable artifacts derived from the simulator model, - not as hand-maintained documentation. +- 다이어그램은 토폴로지 스케일링에 걸쳐 안정적으로 유지된다. +- 거리 또는 라우팅 정책의 변경이 시각적으로 반영된다. +- 다이어그램은 수작업으로 유지되는 문서가 아닌, 시뮬레이터 모델로부터 + 파생된 검증 가능한 산출물의 역할을 한다. --- ## Links - SPEC Section 4 (Output, Debuggability, and Diagrams) -- ADR-0002 (Routing distance semantics) -- ADR-0006 (Topology compilation & automatic diagram generation) +- ADR-0002 (라우팅 거리 의미) +- ADR-0006 (토폴로지 컴파일 및 자동 다이어그램 생성) diff --git a/docs/adr-ko/ADR-0006-dev-topology-compilation-distance-diagram.md b/docs/adr-ko/ADR-0006-dev-topology-compilation-distance-diagram.md index 4b3767c..55e25a0 100644 --- a/docs/adr-ko/ADR-0006-dev-topology-compilation-distance-diagram.md +++ b/docs/adr-ko/ADR-0006-dev-topology-compilation-distance-diagram.md @@ -1,4 +1,4 @@ -# ADR-0006: Topology Compilation, Distance Extraction, and Automatic Diagram Generation +# ADR-0006: 토폴로지 컴파일, 거리 추출, 그리고 자동 다이어그램 생성 ## Status @@ -6,125 +6,125 @@ Accepted ## Context -The simulator compiles topology configuration (e.g., topology.yaml) into an explicit model graph, -and computes routing and accumulated latency (distance). -Diagrams should be generated from these authoritative artifacts to ensure consistency and avoid -hand-maintained topology drawings. +시뮬레이터는 토폴로지 설정(예: topology.yaml)을 명시적인 모델 그래프로 컴파일하고, +라우팅 및 누적 레이턴시(거리)를 계산한다. +정합성을 보장하고 수작업으로 유지되는 토폴로지 도면을 피하기 위해, +다이어그램은 이 권위 있는 산출물로부터 생성되어야 한다. -Additionally, for usability, diagrams should be emitted automatically into a stable location -so that developers can preview them immediately in the repository. +또한 사용성을 위해, 다이어그램은 안정적인 위치로 자동 방출되어 +개발자가 저장소 내에서 즉시 미리볼 수 있어야 한다. --- ## Decision -### D1. Topology compilation is the single source of truth +### D1. 토폴로지 컴파일은 유일한 진실 공급원이다 -- topology.yaml (or equivalent config) is compiled into: - - an explicit system graph, - - node/link attributes, - - routing policies. -This compiled graph is the authoritative representation of the system. +- topology.yaml(또는 동등한 설정)은 다음으로 컴파일된다: + - 명시적인 시스템 그래프, + - 노드/링크 속성, + - 라우팅 정책. +이 컴파일된 그래프가 시스템의 권위 있는 표현이다. -### D2. Distance extraction during compilation +### D2. 컴파일 중 거리 추출 -- During or immediately after topology compilation, the simulator MUST compute distance metadata - (accumulated latency) consistent with ADR-0002. -- Distance metadata MUST be sufficient to support distance-aware diagram layout as defined in ADR-0005. -- Distributed fabric segments (e.g., NOC) MAY have distance_mm = 0 per ADR-0002 D4; - layout placement for such nodes uses explicit position metadata rather than distance buckets. +- 토폴로지 컴파일 중 또는 그 직후, 시뮬레이터는 ADR-0002와 정합되는 + 거리 메타데이터(누적 레이턴시)를 계산해야 한다. +- 거리 메타데이터는 ADR-0005에서 정의한 거리 기반 다이어그램 레이아웃을 지원하기에 충분해야 한다. +- 분산된 패브릭 세그먼트(예: NoC)는 ADR-0002 D4에 따라 distance_mm = 0을 가질 수 있다. + 이러한 노드의 레이아웃 배치는 거리 버킷이 아닌 명시적 위치 메타데이터를 사용한다. -### D3. Diagram generation is a derived artifact +### D3. 다이어그램 생성은 파생 산출물이다 -- Diagrams MUST be generated from: - - the compiled topology graph, - - extracted distance metadata, - - view/layout rules defined in ADR-0005. -- Diagram generation MUST NOT require additional hand-written topology descriptions. +- 다이어그램은 다음으로부터 생성되어야 한다: + - 컴파일된 토폴로지 그래프, + - 추출된 거리 메타데이터, + - ADR-0005에 정의된 뷰/레이아웃 규칙. +- 다이어그램 생성은 추가적인 수작업 토폴로지 기술을 요구해서는 안 된다. -### D4. Automatic diagram emission to the repository +### D4. 저장소로의 자동 다이어그램 방출 -- As part of topology compilation, the implementation MUST produce the following diagrams by default: - - SIP-level diagram (representative, distance-aware) - - CUBE-level diagram (representative, distance-aware) - - PE-level diagram (representative, distance-aware) -- The default output directory is: +- 토폴로지 컴파일의 일부로서, 구현은 기본적으로 다음 다이어그램을 생성해야 한다: + - SIP 수준 다이어그램 (대표, 거리 인식) + - CUBE 수준 다이어그램 (대표, 거리 인식) + - PE 수준 다이어그램 (대표, 거리 인식) +- 기본 출력 디렉터리는 다음과 같다: - `docs/diagrams/` -- The generator MUST overwrite/update only when the compiled topology (or diagram rules) changes. +- 생성기는 컴파일된 토폴로지(또는 다이어그램 규칙)가 변경되었을 때에만 덮어쓰기/업데이트해야 한다. -### D5. View-specific projection and layout +### D5. 뷰별 투영 및 레이아웃 -For each view (SIP / CUBE / PE): +각 뷰(SIP / CUBE / PE)에 대해: -- The generator MUST project the compiled graph into a reduced view graph: - - hide/collapse nodes according to ADR-0005, - - preserve connectivity semantics relevant to that view, - - compute distance buckets and assign layout layers deterministically. -- 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-0017). -- Default anchors are implicit (ADR-0005) and MUST NOT require instance indices. +- 생성기는 컴파일된 그래프를 축소된 뷰 그래프로 투영해야 한다: + - ADR-0005에 따라 노드를 숨기거나 축약하고, + - 해당 뷰와 관련된 연결성 의미를 보존하고, + - 거리 버킷을 계산하여 레이아웃 레이어를 결정론적으로 할당한다. +- CUBE 수준 투영은 다음을 포함해야 한다: + - 라우터 메시 (cube_mesh.yaml로부터), HBM_CTRL, 공유 SRAM, M_CPU, UCIe 포트, + 그리고 불투명 블록으로 표현된 PE. + - 모든 경로(HBM, 비-HBM, 명령)는 동일한 라우터 메시를 통해 라우팅된다 (ADR-0017). +- 기본 앵커는 암묵적이며 (ADR-0005) 인스턴스 인덱스를 요구해서는 안 된다. -### D6. Output formats and determinism +### D6. 출력 포맷과 결정론 -- The generator MUST output at least one of: - - Mermaid (Markdown-native) - - Graphviz DOT (rank-based control) - - SVG (mm-accurate layout, no external dependencies) -- SVG is preferred when mm-accurate position metadata is available from the compiled topology. -- Output MUST be deterministic: - - same topology + same rules → identical diagram text -- File naming MUST be deterministic and stable (see "Output Conventions"). +- 생성기는 다음 중 최소 하나를 출력해야 한다: + - Mermaid (Markdown 네이티브) + - Graphviz DOT (rank 기반 제어) + - SVG (mm 단위 정확도 레이아웃, 외부 의존성 없음) +- 컴파일된 토폴로지로부터 mm 단위 정확도의 위치 메타데이터가 가용한 경우 SVG가 선호된다. +- 출력은 결정론적이어야 한다: + - 동일한 토폴로지 + 동일한 규칙 → 동일한 다이어그램 텍스트 +- 파일 이름은 결정론적이고 안정적이어야 한다 (아래의 "출력 컨벤션" 참조). -### D7. Performance and caching +### D7. 성능 및 캐싱 -- Diagram generation MAY be lazy and/or cached, as long as the outputs in `docs/diagrams/` - remain consistent with the compiled topology. -- The implementation SHOULD use a cache key based on: - - topology content hash, - - routing policy version, - - diagram rules version, - - view type (SIP/CUBE/PE). +- 다이어그램 생성은 지연(lazy) 및/또는 캐시될 수 있으며, `docs/diagrams/`의 출력이 + 컴파일된 토폴로지와 정합을 유지하는 한 그렇다. +- 구현은 다음을 기반으로 한 캐시 키를 사용해야 한다(SHOULD): + - 토폴로지 콘텐츠 해시, + - 라우팅 정책 버전, + - 다이어그램 규칙 버전, + - 뷰 유형 (SIP/CUBE/PE). --- -## Output Conventions +## 출력 컨벤션 -### Directory +### 디렉터리 -- `docs/diagrams/` is the canonical output directory for generated diagrams. +- `docs/diagrams/`는 생성된 다이어그램의 표준 출력 디렉터리이다. -### File names (recommended, deterministic) +### 파일 이름 (권장, 결정론적) - `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` -Optionally, for multi-topology workflows: +선택적으로, 멀티 토폴로지 워크플로우용: - `sip_view__{topology_id}.svg` - `cube_view__{topology_id}.svg` - `pe_view__{topology_id}.svg` -### Repository policy +### 저장소 정책 -- Generated diagram files MAY be committed to the repository to enable diff-based review. -- If committed, they MUST be reproducible from topology compilation. +- 생성된 다이어그램 파일은 diff 기반 리뷰가 가능하도록 저장소에 커밋될 수 있다. +- 커밋된 경우, 이는 토폴로지 컴파일로부터 재현 가능해야 한다. --- ## Consequences -- Diagrams are always consistent with simulator behavior. -- Architectural changes automatically propagate to visualizations. -- Diagram diffs become meaningful indicators of architectural change. +- 다이어그램은 항상 시뮬레이터 동작과 정합한다. +- 아키텍처 변경이 시각화에 자동으로 전파된다. +- 다이어그램 diff는 아키텍처 변경의 의미 있는 지표가 된다. --- ## Links - SPEC Section 4 (Output, Debuggability, and Diagrams) -- ADR-0002 (Distance semantics) -- ADR-0005 (Diagram views and layout rules) +- ADR-0002 (거리 의미) +- ADR-0005 (다이어그램 뷰 및 레이아웃 규칙) diff --git a/docs/adr-ko/ADR-0007-api-runtime-api-boundaries.md b/docs/adr-ko/ADR-0007-api-runtime-api-boundaries.md index 9522b9d..5d27e1e 100644 --- a/docs/adr-ko/ADR-0007-api-runtime-api-boundaries.md +++ b/docs/adr-ko/ADR-0007-api-runtime-api-boundaries.md @@ -1,4 +1,4 @@ -# ADR-0007: Runtime API and Simulation Engine Boundaries +# ADR-0007: 런타임 API 및 시뮬레이션 엔진 경계 ## Status @@ -6,90 +6,90 @@ Accepted ## Context -The simulator consists of multiple layers with distinct responsibilities: +시뮬레이터는 책임이 명확히 다른 여러 계층으로 구성된다: -- a host-facing API layer used by benchmarks and user code, -- a discrete-event simulation engine that executes requests, -- device components that model hardware behavior. +- 벤치마크와 사용자 코드가 사용하는 호스트 대상 API 계층, +- 요청을 실행하는 이산 이벤트 시뮬레이션 엔진, +- 하드웨어 동작을 모델링하는 디바이스 컴포넌트. -Without strict boundaries, orchestration logic can leak into components, -or simulation internals can become entangled with user-facing APIs. +엄격한 경계가 없으면 오케스트레이션 로직이 컴포넌트로 누출되거나 +시뮬레이션 내부가 사용자 대상 API와 얽힐 수 있다. -This ADR defines clear responsibility boundaries between: +본 ADR은 다음 사이의 명확한 책임 경계를 정의한다: -- runtime API, -- simulation engine (sim_engine), -- hardware components. +- 런타임 API, +- 시뮬레이션 엔진 (sip_engine), +- 하드웨어 컴포넌트. --- ## Decision -### D1. Runtime API is host-facing orchestration only +### D1. 런타임 API는 호스트 대상 오케스트레이션만 담당 -The runtime API represents host/driver-level behavior and MUST: +런타임 API는 호스트/드라이버 수준의 동작을 표현하며 다음을 해야 한다: -- expose high-level operations (tensor deployment, kernel launch), -- submit requests only to endpoint components (e.g., IO_CPU), -- await completion via futures/handles, -- own and persist host-side metadata (tensor allocation maps, kernel bindings). +- 고수준 동작 노출 (텐서 배포, 커널 launch), +- 엔드포인트 컴포넌트(예: IO_CPU)에만 요청 제출, +- futures/handles로 완료 대기, +- 호스트측 메타데이터(텐서 할당 맵, 커널 바인딩)의 소유와 영속화. -The runtime API MUST NOT: +런타임 API가 해서는 안 되는 것: -- hardcode hop-by-hop routing or fan-out, -- directly invoke internal components (M_CPU, PE_CPU, engines), -- embed topology- or routing-specific assumptions. +- hop-by-hop 라우팅 또는 fan-out 하드코딩, +- 내부 컴포넌트(M_CPU, PE_CPU, 엔진) 직접 호출, +- 토폴로지나 라우팅 관련 가정 내장. --- -### D2. Simulation engine wires components and tracks completion +### D2. 시뮬레이션 엔진은 컴포넌트를 연결하고 완료를 추적 -The simulation engine (sim_engine) MUST: +시뮬레이션 엔진(sim_engine)은 다음을 해야 한다: -- 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. +- 초기화 시점에 컴포넌트 연결 (컴포넌트 포트/와이어 프레임워크에 따라 + 포트 store 생성 + 와이어 프로세스 시작 — ADR-0015), +- 컴파일된 토폴로지 그래프의 진입 컴포넌트(예: 메모리 동작은 PCIE_EP, + 커널 launch는 IO_CPU)에 요청 주입, +- 이산 이벤트 모델로 이벤트 스케줄링과 실행, +- correlation ID와 완료 추적 관리. -The simulation engine MUST NOT: +시뮬레이션 엔진이 해서는 안 되는 것: -- define tensor semantics, -- define kernel execution policies, -- 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). +- 텐서 시맨틱 정의, +- 커널 실행 정책 정의, +- 런타임 API에 내부 그래프 세부사항 노출, +- 요청 실행 중에 토폴로지 경로를 따라 걷기, +- 컴포넌트의 `run()` 메서드 직접 호출, +- hop별 레이턴시 추적 또는 fan-out 분해 (컴포넌트의 책임). --- -### D3. Components own fan-out and aggregation +### D3. 컴포넌트가 fan-out과 집계를 담당 -Device-side components MUST: +디바이스측 컴포넌트는 다음을 해야 한다: -- fan-out requests to downstream domains - (IO_CPU → M_CPU → PE_CPU → schedulers/engines), -- aggregate completion and failure signals, -- propagate results deterministically upstream. +- 요청을 하위 도메인으로 fan-out + (IO_CPU → M_CPU → PE_CPU → 스케줄러/엔진), +- 완료·실패 신호 집계, +- 결정론적으로 상위로 결과 전파. -Neither the runtime API nor the simulation engine may orchestrate -component-level fan-out explicitly. +런타임 API와 시뮬레이션 엔진 모두 컴포넌트 수준의 fan-out을 명시적으로 +오케스트레이션해서는 안 된다. --- ## Consequences -- Runtime APIs remain stable as topology and routing evolve. -- Simulation internals can change without affecting user-facing code. -- Component implementations remain swappable via DI. +- 토폴로지와 라우팅이 변해도 런타임 API는 안정적이다. +- 시뮬레이션 내부는 사용자 대상 코드에 영향을 주지 않고 변경 가능하다. +- 컴포넌트 구현은 DI로 교체 가능한 상태가 유지된다. --- ## Links - 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) +- ADR-0008 (텐서 배포) +- ADR-0009 (커널 실행) +- ADR-0015 (컴포넌트 포트/와이어 모델과 엔진 역할) +- ADR-0010 (CLI 표면과 실행 시맨틱 — 런타임 API 소비자) diff --git a/docs/adr-ko/ADR-0008-api-tensor-deploy-and-allocation.md b/docs/adr-ko/ADR-0008-api-tensor-deploy-and-allocation.md index 36ca4da..1b95499 100644 --- a/docs/adr-ko/ADR-0008-api-tensor-deploy-and-allocation.md +++ b/docs/adr-ko/ADR-0008-api-tensor-deploy-and-allocation.md @@ -1,4 +1,4 @@ -# ADR-0008: Tensor Deployment and Allocation (Host Allocator, PA-first) +# ADR-0008: 텐서 배포 및 할당 (호스트 할당기, PA 우선) ## Status @@ -6,95 +6,95 @@ Accepted ## Context -Benchmarks require PyTorch-like tensor semantics: +벤치마크는 PyTorch와 유사한 텐서 시맨틱을 요구한다: -- tensor creation (empty, fill), -- deployment to accelerator devices (tensor.to()). +- 텐서 생성 (empty, fill), +- 가속기 디바이스로의 배포 (tensor.to()). -In the realistic system, host software manages allocation/mapping and installs -mappings for DMA/MMU. For Phase 0 we simplify (ADR-0011): +현실적인 시스템에서는 호스트 소프트웨어가 할당·매핑을 관리하고 DMA/MMU +매핑을 설치한다. Phase 0에서는 (ADR-0011) 다음으로 단순화한다: -- device memory operations use PA only, -- VA/MMU/IOMMU is not modeled. +- 디바이스 메모리 동작은 PA만 사용, +- VA/MMU/IOMMU는 모델링하지 않는다. -To keep the host↔device interface minimal, we avoid a separate -AllocateTensorMeta message. Instead, host allocation produces a PA shard map -that is used directly by MemoryWrite/Read and KernelLaunch. +호스트↔디바이스 인터페이스를 최소로 유지하기 위해 별도의 +AllocateTensorMeta 메시지는 피한다. 대신 호스트 할당은 PA 샤드 맵을 +생성하여 MemoryWrite/Read와 KernelLaunch가 직접 사용한다. --- ## Decision -### D1. Tensor is a host-owned handle with PA shard mapping +### D1. Tensor는 PA 샤드 매핑을 가진 호스트 소유 핸들 -A Tensor object is a host-owned handle that encapsulates: +Tensor 객체는 다음을 캡슐화하는 호스트 소유 핸들이다: -- shape and dtype, -- initialization intent, -- device placement and allocation metadata as a PA shard map. +- shape과 dtype, +- 초기화 의도, +- PA 샤드 맵 형태의 디바이스 배치 및 할당 메타데이터. -After deployment, the Tensor handle MUST contain: +배포 이후 Tensor 핸들은 다음을 포함해야 한다: -- a list of shards, each with (sip,cube,pe,pa,nbytes,offset_bytes). +- 각각 (sip, cube, pe, pa, nbytes, offset_bytes)를 가진 샤드 리스트. -This PA shard mapping is the single source of truth for kernel argument binding. +이 PA 샤드 매핑이 커널 인수 바인딩의 단일 진실 원천이다. --- -### D2. Deployment uses a host allocator (Phase 0) +### D2. 배포는 호스트 할당기를 사용한다 (Phase 0) -In Phase 0, tensor deployment produces PA shard mappings via a host allocator: +Phase 0에서 텐서 배포는 호스트 할당기를 통해 PA 샤드 매핑을 생성한다: -- placement (split/replicate/hybrid) is decided by a DP policy, -- allocation assigns PA ranges at the PE level and returns shard mappings, -- the Tensor handle stores the resulting shard list deterministically. +- 배치(split/replicate/hybrid)는 DP 정책에 의해 결정, +- 할당은 PE 수준에서 PA 범위를 부여하고 샤드 매핑을 반환, +- Tensor 핸들은 결정론적으로 결과 샤드 리스트를 저장. -No separate host-visible device allocation RPC is required in Phase 0. +Phase 0에서는 호스트가 보는 별도의 디바이스 할당 RPC는 필요하지 않다. --- -### D3. Data initialization and transfer uses MemoryWrite/Read only +### D3. 데이터 초기화와 전송은 MemoryWrite/Read만 사용 -Any data initialization or transfer implied by a tensor (e.g., fill, copy) -MUST be represented using Host ↔ IO_CPU messages only: +텐서가 함의하는 모든 데이터 초기화나 전송(예: fill, copy)은 +Host ↔ IO_CPU 메시지만으로 표현되어야 한다: - MemoryWrite - MemoryRead -Rules: +규칙: -- MemoryWrite/Read MUST reference PA + (sip,cube,pe) tags (ADR-0012). -- Allocation metadata MUST NOT be embedded as a separate allocation message. -- Bulk tensor data MUST NOT be embedded in Phase 0 messages. +- MemoryWrite/Read는 PA + (sip, cube, pe) 태그를 참조해야 한다 (ADR-0012). +- 할당 메타데이터는 별도의 할당 메시지로 임베드되어서는 안 된다. +- 대량 텐서 데이터는 Phase 0 메시지에 임베드되어서는 안 된다. -The simulation engine schedules MemoryWrite/Read through the graph so that -latency is computed by explicit traversal. +시뮬레이션 엔진은 MemoryWrite/Read를 그래프를 통해 스케줄하므로 레이턴시는 +명시적 순회로 계산된다. --- -### D4. Extension path (non-breaking) +### D4. 확장 경로 (호환성 유지) -Future ADRs MAY introduce optional VA/MMU/IOMMU modeling by adding: +향후 ADR이 다음을 추가하여 선택적인 VA/MMU/IOMMU 모델링을 도입할 수 있다: -- virtual addressing in tensor handles, -- mapping install steps, -- translation latency/page granularity. +- 텐서 핸들에 가상 주소, +- 매핑 설치 단계, +- 변환 레이턴시·페이지 granularity. -The Phase 0 PA shard map remains a valid fast-path configuration. +Phase 0의 PA 샤드 맵은 유효한 fast-path 구성으로 유지된다. --- ## Consequences -- Host↔IO_CPU contract remains minimal (MemoryRead/Write + KernelLaunch). -- KernelLaunch can pass per-PE data placement explicitly via shard tags. -- Early implementation stays simple and testable. +- Host↔IO_CPU 계약이 최소(MemoryRead/Write + KernelLaunch)로 유지된다. +- KernelLaunch가 샤드 태그를 통해 PE별 데이터 배치를 명시적으로 전달할 수 있다. +- 초기 구현이 단순하고 테스트 가능하게 유지된다. --- ## Links -- 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) +- ADR-0011 (메모리 주소 지정 — PA / VA / LA) +- ADR-0012 (Host↔IO_CPU 스키마) +- ADR-0007 (runtime_api vs sim_engine 경계) +- ADR-0009 (커널 실행) diff --git a/docs/adr-ko/ADR-0009-api-kernel-execution-messaging.md b/docs/adr-ko/ADR-0009-api-kernel-execution-messaging.md index a94be07..1204d5e 100644 --- a/docs/adr-ko/ADR-0009-api-kernel-execution-messaging.md +++ b/docs/adr-ko/ADR-0009-api-kernel-execution-messaging.md @@ -1,4 +1,4 @@ -# ADR-0009: Kernel Execution Messaging and Completion Semantics +# ADR-0009: 커널 실행 메시징 및 완료 시맨틱 ## Status @@ -6,89 +6,86 @@ Accepted ## Context -Kernel execution is initiated by the host and proceeds through -device control components: +커널 실행은 호스트에서 시작되어 디바이스 측 제어 컴포넌트를 통해 진행된다: -Host → IO_CPU → M_CPU → PE_CPU → schedulers → engines +Host → IO_CPU → M_CPU → PE_CPU → 스케줄러 → 엔진 -Completion propagates in reverse order. +완료는 역방향으로 전파된다. -To keep benchmarks simple and topology-agnostic, -kernel execution must be endpoint-driven with deterministic aggregation. +벤치마크를 단순하고 토폴로지에 비의존적으로 유지하기 위해, 커널 실행은 +엔드포인트 기반(endpoint-driven)이어야 하며 완료 집계는 결정론적이어야 한다. --- ## Decision -### D1. Kernel launch is an endpoint request +### D1. 커널 런치는 엔드포인트 요청이다 -A kernel launch is initiated by submitting a single KernelLaunch request -to the IO_CPU endpoint. +커널 런치는 IO_CPU 엔드포인트에 단일 KernelLaunch 요청을 제출함으로써 +시작된다. -The runtime API MUST: +runtime API는 반드시: -- construct the kernel launch request, -- submit it to IO_CPU, -- await a single completion result. +- 커널 런치 요청을 구성하고, +- 이를 IO_CPU로 제출하며, +- 단일 완료 결과를 대기해야 한다. -The runtime API MUST NOT orchestrate internal fan-out. +runtime API는 내부 팬아웃(fan-out)을 직접 조율해서는 안 된다. --- -### D2. Tensor arguments are passed by metadata +### D2. 텐서 인자는 메타데이터로 전달된다 -KernelLaunch requests MUST reference tensor arguments via: +KernelLaunch 요청은 텐서 인자를 다음을 통해 참조해야 한다: -- 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 +### D3. 팬아웃과 집계는 컴포넌트의 책임이다 -- IO_CPU fans out work to M_CPUs. -- M_CPU fans out work to PE_CPUs. -- PE_CPU manages kernel execution and engine dispatch. +- IO_CPU는 작업을 M_CPU들에게 팬아웃한다. +- M_CPU는 작업을 PE_CPU들에게 팬아웃한다. +- PE_CPU는 커널 실행과 엔진 디스패치를 관리한다. -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. +- M_CPU는 대상 PE들이 모두 완료되거나 실패 정책이 트리거되면 완료된다. +- IO_CPU는 대상 큐브들이 모두 완료되거나 실패 정책이 트리거되면 완료된다. --- -### D4. Completion and failure propagation +### D4. 완료 및 실패 전파 -- All messages MUST carry correlation identifiers. -- Completion and failure MUST propagate deterministically to the host. -- The simulation engine provides futures/handles to observe completion. +- 모든 메시지는 correlation ID를 포함해야 한다. +- 완료와 실패는 호스트로 결정론적으로 전파되어야 한다. +- 시뮬레이션 엔진은 완료를 관찰할 수 있는 future/handle을 제공한다. --- -### D5. Launch timing is endpoint-synchronized +### D5. 런치 타이밍은 엔드포인트 동기화된다 -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. +단일 커널 런치가 지정한 모든 PE는 런치 진입점으로부터의 디스패치 경로 길이와 +무관하게, 동일한 시뮬레이션 시각에 커널 본문 실행을 시작해야 한다. -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). +근거. 디스패치 트리 Host → IO_CPU → M_CPU → PE_CPU는 모든 레벨에서 가변 +레이턴시를 가진다. M_CPU에 가까운 PE는 멀리 있는 PE보다 런치를 더 일찍 +수신하고, IO_CPU에 가까운 큐브는 먼 큐브보다 더 일찍 수신한다. 동기화가 +없으면 각 PE의 커널은 서로 다른 `env.now`에서 시작되어, `pe_exec_ns`와 같은 +PE별 메트릭이 커널 자체의 동작이 아니라 디스패치 경로 기하 구조의 함수가 +된다 — 그 결과 커널 내부 대기(예: 큐브 간 또는 SIP 간 홉에서의 `tl.recv`)를 +타이밍하는 벤치마크에서 측정 아티팩트가 발생한다. -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**: +- `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) = @@ -98,49 +95,44 @@ Mechanism. - 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. + 이는 실제 디스패치를 **두 개의 순차적 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 불변식을 위반하게 된다. - 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. + 팬아웃된 하위 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-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). +IO_CPU 레벨 스탬핑은 모든 대상 큐브의 모든 PE가 동일한 배리어 시뮬레이션 +시각을 사용하도록 보장하여, 큐브 내 디스패치 오프셋 아티팩트와 다중 큐브 +런치에서의 큐브 간 오프셋 아티팩트를 모두 제거한다. 실제 하드웨어의 +타이밍 브로드캐스트 런치(레이턴시 등화 디스패치 트리)를 모델링한다. -The synchronization is internal to the engine / IO_CPU / M_CPU / PE_CPU -control plane — runtime API and application kernels are unchanged. +이 동기화는 엔진 / IO_CPU / M_CPU / PE_CPU 제어 평면 내부에서 수행된다 — +runtime API와 애플리케이션 커널은 변경되지 않는다. --- ## 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) +- ADR-0007 (Runtime API 경계) +- ADR-0008 (텐서 배치) +- ADR-0013 (검증 전략 — V2 팬아웃 테스트) +- ADR-0015 D4 (커널 런치의 구체적 패브릭 경로) diff --git a/docs/adr-ko/ADR-0010-api-cli-surface-and-semantics.md b/docs/adr-ko/ADR-0010-api-cli-surface-and-semantics.md index 4925a64..0e29028 100644 --- a/docs/adr-ko/ADR-0010-api-cli-surface-and-semantics.md +++ b/docs/adr-ko/ADR-0010-api-cli-surface-and-semantics.md @@ -1,4 +1,4 @@ -# ADR-0010: Command Line Interface and Execution Semantics +# ADR-0010: 명령줄 인터페이스 및 실행 시맨틱 ## Status @@ -6,126 +6,120 @@ Accepted ## Context -The `kernbench` CLI is the user-facing entry point of the simulator. It -exposes three subcommands: +`kernbench` CLI는 시뮬레이터의 사용자 대면 진입점이다. 세 개의 서브명령을 +노출한다: -- `run` — execute a benchmark against a topology. -- `probe` — diagnostic utility for latency / BW measurement. -- `web` — interactive topology viewer. +- `run` — 토폴로지에 대해 벤치마크를 실행한다. +- `probe` — 레이턴시 / 대역폭 측정을 위한 진단 유틸리티. +- `web` — 인터랙티브 토폴로지 뷰어. -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. +디바이스 열거는 CLI에 중앙 집중화되어 있다. runtime API와 시뮬레이션 엔진 +모두 디바이스를 열거하지 않는다. 벤치마크는 설계상 단일 디바이스를 +유지하며 입력으로 디바이스 식별자를 받는다. ## Decision -### D1. Benchmark contract — single-device by design +### D1. 벤치마크 계약 — 설계상 단일 디바이스 -- 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. +다중 디바이스 실행은 벤치마크의 관심사가 아니라 CLI의 관심사이다(D3). -### D2. `kernbench run` — benchmark execution +### D2. `kernbench run` — 벤치마크 실행 -Required arguments: +필수 인자: -- `--topology `: topology YAML file path. Loaded via - `resolve_topology()`. -- `--bench `: benchmark name. Resolved via - `benches.loader.resolve_bench()`. +- `--topology `: 토폴로지 YAML 파일 경로. `resolve_topology()`를 + 통해 로드된다. +- `--bench `: 벤치마크 이름. `benches.loader.resolve_bench()`를 + 통해 해석된다. -Optional arguments: +선택 인자: -- `--device ` (default: `all`): - - `all` — run once per discovered SIP (see D3). - - `sip:` — 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. +- `--device ` (기본값: `all`): + - `all` — 발견된 SIP마다 한 번씩 실행한다(D3 참고). + - `sip:` — SIP N에서만 실행한다. + - `resolve_device()`를 통해 파싱된다. +- `--verify-data` (기본값: off) — Phase 2 데이터 검증을 활성화한다 + (ADR-0020 참고). 설정되면 `engine_factory`가 엔진을 + `enable_data=True`로 구성한다. 벤치마크 실행 후, 기록된 op들의 진단 + 요약이 출력된다. -Each invocation runs the benchmark once within a single simulation -instance. +각 호출은 단일 시뮬레이션 인스턴스 내에서 벤치마크를 한 번 실행한다. -### D3. Multi-device execution is logically parallel +### D3. 다중 디바이스 실행은 논리적으로 병렬이다 -When `--device all` (or omitted) and the topology has multiple SIPs: +`--device all`(또는 생략) 상태이며 토폴로지에 SIP가 여러 개일 때: -- 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.). +- 벤치마크 실행은 단일 시뮬레이션 엔진 인스턴스에 제출된다. +- 시뮬레이션 시간 상에서 실행은 논리적으로 병렬이다. +- 디바이스 간 경합(공유 패브릭 대역폭, SIP 간 트래픽 등)이 자연스럽게 + 모델링된다. -The CLI does NOT spawn multiple OS processes or independent -simulation runs — parallelism is internal to one simulation instance. +CLI는 여러 OS 프로세스나 독립된 시뮬레이션 실행을 생성하지 **않는다** — +병렬성은 단일 시뮬레이션 인스턴스 내부에서 일어난다. -### D4. `kernbench probe` — latency / BW diagnostic utility +### D4. `kernbench probe` — 레이턴시 / 대역폭 진단 유틸리티 -Required argument: +필수 인자: -- `--topology `: topology YAML file path. +- `--topology `: 토폴로지 YAML 파일 경로. -Optional argument: +선택 인자: -- `--case ` (default: `all`) — run a predefined traffic - pattern, or `all` to run every defined case. +- `--case ` (기본값: `all`) — 미리 정의된 트래픽 패턴을 실행하거나, + `all`로 정의된 모든 케이스를 실행한다. -Probe runs each pattern through the simulation engine and reports -per case: +Probe는 시뮬레이션 엔진을 통해 각 패턴을 실행하고 케이스별로 다음을 +보고한다: -- End-to-end latency (ns). -- Effective bandwidth (nbytes / total_ns). -- Bottleneck bandwidth (min edge BW along the chosen path). -- Utilization (effective / bottleneck). +- 종단 간 레이턴시(ns). +- 유효 대역폭(nbytes / total_ns). +- 병목 대역폭(선택된 경로상의 최소 엣지 BW). +- 활용률(유효 / 병목). -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. +Probe는 추가로 단조성 불변식을 검증한다 — 예를 들어 local-HBM 접근 ≤ +큐브 내 PE 간 ≤ 큐브 간 ≤ SIP 간 — 그리고 위반을 보고한다. Probe는 +레이턴시 / 대역폭 모델을 검증하기 위한 개발자 도구이다; 벤치마크가 +아니다. -### D5. `kernbench web` — topology viewer +### D5. `kernbench web` — 토폴로지 뷰어 -Optional arguments: +선택 인자: -- `--port ` (default: `8765`) — HTTP port. -- `--no-open` — do not auto-open the browser. +- `--port ` (기본값: `8765`) — HTTP 포트. +- `--no-open` — 브라우저를 자동으로 열지 않는다. -Launches a local HTTP server that renders the compiled topology in -the browser. Distinct from the static `docs/diagrams/` artifacts: +컴파일된 토폴로지를 브라우저에서 렌더링하는 로컬 HTTP 서버를 띄운다. +정적인 `docs/diagrams/` 산출물과는 구별된다: -- `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. +- `docs/diagrams/` 파일은 토폴로지 컴파일 시점에 파생된다(ADR-0006). +- `kernbench web`은 인터랙티브이다 — 팬/줌, 컴포넌트 속성 호버, + SIP / CUBE / PE 뷰 간 전환. -### D6. Runtime API and simulation engine remain device-scoped +### D6. runtime API와 시뮬레이션 엔진은 디바이스 스코프를 유지한다 -- Runtime API calls operate on one device per invocation. -- The simulation engine schedules all requests deterministically. -- Neither layer enumerates devices. +- runtime API 호출은 호출당 하나의 디바이스에서 동작한다. +- 시뮬레이션 엔진은 모든 요청을 결정론적으로 스케줄링한다. +- 어느 레이어도 디바이스를 열거하지 않는다. -This invariant keeps each layer testable in isolation; device -enumeration and multi-device fan-out live only in the CLI's `run` -command (D3). +이 불변식은 각 레이어를 독립적으로 테스트 가능하게 유지한다; 디바이스 +열거와 다중 디바이스 팬아웃은 오직 CLI의 `run` 명령에만 존재한다(D3). ## 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. +- 벤치마크 작성자는 단일 디바이스 로직을 작성한다; 다중 디바이스 동작은 + CLI가 SIP들에 걸쳐 디스패치함으로써 자연스럽게 도출된다. +- 새로운 서브명령(예: 트레이스 내보내기, 리플레이) 추가는 벤치마크나 + runtime API 변경을 요구하지 않는다 — CLI가 확장 포인트이다. +- `probe`와 `web`은 진단/시각화 도구이며 벤치마크가 아니다; 벤치마크 로더 + 경로를 우회한다. ## 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`) +- ADR-0007 (Runtime API와 시뮬레이션 엔진 경계) +- ADR-0020 (Two-pass 데이터 실행 — `--verify-data`) +- ADR-0006 (토폴로지 컴파일과 다이어그램 생성 — `kernbench web`의 배경) diff --git a/docs/adr-ko/ADR-0011-mem-memory-addressing-simplification.md b/docs/adr-ko/ADR-0011-mem-memory-addressing-simplification.md index 064c365..59ed154 100644 --- a/docs/adr-ko/ADR-0011-mem-memory-addressing-simplification.md +++ b/docs/adr-ko/ADR-0011-mem-memory-addressing-simplification.md @@ -1,288 +1,273 @@ -# ADR-0011: Memory Addressing — PA / VA / LA Address Models +# ADR-0011: 메모리 주소 지정 — PA / VA / LA 주소 모델 ## Status Accepted. -- **VA model: currently implemented (default).** -- PA model: implemented as PageFault fallback in PE_DMA. -- LA model: proposed, not implemented. +- **VA 모델: 현재 구현됨 (기본값).** +- PA 모델: PE_DMA의 PageFault fallback으로 구현됨. +- LA 모델: 제안됨, 미구현. ## 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. +KernBench의 주소 모델은 각 단계마다 이전 단계의 한계를 해결하면서 +세 단계의 설계 지점을 거쳐 발전해 왔다. 본 ADR은 미래의 구현 작업이 +이 셋 중 하나를 선택해야 하므로 셋 모두를 한 곳에 기록한다. -### PA-only baseline +### PA 단독 베이스라인 -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. +KernBench Phase 0는 모든 디바이스 메모리 동작(MemoryRead/MemoryWrite)을 +순수 물리 주소 전송으로 다뤘다. 호스트측 가상 주소 지정 없음, MMU/IOMMU +변환 없음. 할당기는 PA 매핑을 반환하고, DMA 요청은 PA를 직접 운반했다. -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. +이는 초기 정확성·레이턴시 작업에는 충분했지만, 샤딩된 텐서에 대해 +`base_addr + offset` 패턴을 사용하는 표준 Triton 커널을 실행하기에는 +부족했다. 각 PE의 샤드는 서로 다른 PA를 갖지만, 커널은 offset을 계산하기 +위해 연속된 단일 주소 공간이 필요하기 때문이다. -### Why VA/MMU (current default) +### VA/MMU를 채택한 이유 (현재 기본값) -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. +현실적인 시스템은 호스트측 가상 주소 지정과 DMA를 위한 MMU/IOMMU 스타일 +변환 경로를 사용한다. 호스트는 PE 수준에서 물리 메모리를 할당하고, +그것을 가상 주소 공간에 매핑하여 매핑을 설치한 뒤, DMA 요청은 가상 +주소를 사용하며 그것이 물리 주소로 변환된다. -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. +이 모델을 채택하면 커널이 연속된 VA 범위에 대해 `base_addr + offset`을 +사용할 수 있고, 디바이스측 MMU가 각 접근을 적절한 PA로 변환한다. -### Why LA/BAAW (proposed) +### LA/BAAW를 제안한 이유 -VA/MMU treats HBM as a single backing space. KernBench needs to -explore architectures where HBM is composed of multiple pseudo -channels in parallel: +VA/MMU는 HBM을 단일 backing 공간으로 다룬다. KernBench는 HBM이 병렬로 +여러 pseudo channel로 구성된 아키텍처를 탐색해야 한다: -- CUBE's HBM has 32 or 64 pseudo channels. -- In a PE-Local-HBM model, each PE is assigned N pseudo channels +- CUBE의 HBM은 32 또는 64개의 pseudo channel을 갖는다. +- PE-Local-HBM 모델에서 각 PE에는 N개의 pseudo channel이 할당된다 (N = `hbm_pseudo_channels / pes_per_cube`). -- Per-channel BW (e.g. 32 GB/s) determines aggregate PE BW - (N × per-channel). +- 채널당 대역폭(예: 32 GB/s)이 PE의 총 대역폭을 결정한다 + (N × 채널당). -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. +- **1:1 모드** — 하나의 논리 접근 → N개의 채널별 요청. + 채널별 대역폭 경쟁을 정밀하게 모델링. +- **n:1 모드 (기본값)** — 하나의 논리 접근 → 하나의 집계 요청. + 채널들이 interleave된다고 가정; 집계된 대역폭 모델. -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). +VA의 `tl.load(va_ptr)`은 하나의 목표에 대한 하나의 DMA 요청을 생성한다. +이를 PE_DMA 내부에서 채널별 요청으로 분해하려면 주소 계층이 채널을 +인지해야 한다. 이것이 BAAW(Logical-to-Physical Mapping Unit)를 가진 +LA(Logical Address) 추상화의 역할이다. -Core requirements driving the LA design: +LA 설계를 이끄는 핵심 요구사항: -- 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. +- PE_DMA → HBM_CTRL 유효 대역폭 시맨틱이 두 모드에서 동일해야 한다 + (요청 형태와 자원 모델만 다름). +- 커널 프로그래밍 모델은 변경되지 않는다 — 물리 채널 정보는 커널 코드에 + 절대 노출되지 않는다. +- 모드 전환은 토폴로지 수준의 설정이다. -### 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 | +| 모델 | 상태 | 핵심 아이디어 | +|------|------|--------------| +| PA | fallback (구현됨) | 직접 물리 주소 지정, 변환 없음 | +| VA | 현재 기본값 (구현됨) | 텐서별 연속 VA 범위; MMU가 접근별로 변환 | +| LA | 제안됨 | LA + BAAW가 (PA, 채널)로 해석; 1:1 및 n:1 채널 매핑 모드 지원 | --- ## 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. +본 ADR은 세 개의 주소 모델을 정의한다. 어느 시점에도 시스템은 정확히 +한 모델로 동작한다. 선택은 토폴로지·설정 주도이며, 단일 시뮬레이션 실행 +내에서의 공존은 요구되지 않는다. --- -### Address Model: PA (Physical Address) — fallback +### 주소 모델: PA (물리 주소) — fallback -#### D-PA1. PA-only semantics +#### D-PA1. PA 단독 시맨틱 -- 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. +- 모든 디바이스 메모리 접근(MemoryRead/MemoryWrite)은 디바이스 물리 주소(PA)와 + 크기에 대해 동작한다. +- PA 단독 모드는 PE_DMA의 PageFault fallback 경로를 통해 여전히 동작한다. + DMA src/dst 주소에 MMU 매핑이 없으면 PE_DMA는 그 값을 PA로 직접 다룬다. -#### D-PA2. Allocation produces PA mappings +#### D-PA2. 할당은 PA 매핑을 생성한다 -Device allocation selects PE-local memory regions and returns PA -mappings sufficient to execute kernels and issue DMA requests. +디바이스 할당은 PE 로컬 메모리 영역을 선택하고 커널 실행 및 DMA 요청 +발행에 충분한 PA 매핑을 반환한다. -PA model is retained primarily for backward compatibility with PA-only -tests and as the underlying physical layer that VA / LA models resolve -into. +PA 모델은 주로 PA 단독 테스트와의 하위 호환성을 위해, 그리고 VA / LA +모델이 해석되어 들어가는 기저 물리 계층으로 유지된다. --- -### Address Model: VA (Virtual Address with MMU) — current default +### 주소 모델: VA (MMU를 동반한 가상 주소) — 현재 기본값 -#### D-VA1. Virtual Address Model +#### D-VA1. 가상 주소 모델 -- 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). +- 각 텐서는 하나의 연속된 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 Component +#### D-VA2. PE_MMU 컴포넌트 -- 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). +- 하이브리드 설계: SimPy 컴포넌트(`MmuMapMsg`용 inbox) + 유틸리티 + (PE_DMA가 호출하는 동기식 `translate()`). +- 페이지 정렬 dict 조회로 O(1) VA → PA 변환. +- `tlb_overhead_ns`로 접근당 레이턴시 설정 가능. +- PageFault fallback: VA에 매핑이 없으면 PE_DMA가 그것을 PA로 직접 + 다룬다 (PA 모델과의 하위 호환성 유지). -#### D-VA3. Mapping Installation +#### D-VA3. 매핑 설치 -- `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. +- `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. Tensor Lifecycle +#### D-VA4. 텐서 라이프사이클 -- `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. +- `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. Allocators +#### D-VA5. 할당기 -- `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). +- `VirtualAllocator`: 디바이스 전체의 VA 공간, coalescing을 동반한 + 페이지 정렬 alloc/free. +- `PEMemAllocator`: PE별 HBM/TCM, coalescing을 동반한 free-list 기반 + alloc/free. +- 페이지 크기는 `topology.yaml`의 `pe_mmu` attrs로 설정 가능 + (기본 4096). -#### Consequences (VA model) +#### Consequences (VA 모델) -- 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). +- Triton 커널은 샤딩된 텐서에 대해 `base_addr + offset` 패턴을 자연스럽게 + 사용한다. +- 모든 레이턴시는 MMU 매핑 설치와 접근당 TLB 오버헤드를 포함하여 + 그래프 순회를 통해 명시적이다. +- PA 단독 모드는 fallback으로 유지된다 (PageFault → PA로 처리). +- IPCQ와 그 외 고정 주소 자원은 MMU를 우회한다 (PA 직접 사용). --- -### Address Model: LA (Logical Address with BAAW) — proposed +### 주소 모델: LA (BAAW를 동반한 논리 주소) — 제안됨 -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. +LA는 채널 수준 HBM 모델링이 필요할 때 VA를 대체한다. +이 모델을 채택하면 VA/MMU 인프라가 제거된다 (D-LA1이 제거되는 산출물을 +나열한다). 동일 실행 내에서 VA와의 공존은 목표가 아니다. -#### D-LA1. LA introduction — replaces VA infrastructure +#### D-LA1. LA 도입 — VA 인프라 대체 -LA is the sole address space used by kernel code (`tl.load`, -`tl.store`, `tl.composite`). Properties: +LA는 커널 코드(`tl.load`, `tl.store`, `tl.composite`)가 사용하는 +유일한 주소 공간이다. 속성: -- 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. +- Tensor를 연속된 논리 공간에 매핑할 수 있다 (VA처럼). +- `(논리 버퍼 + offset)`을 표현한다. +- 물리 채널 정보를 직접 포함하지 **않는다**. +- 물리적 해석이 일어나기 전까지는 중간 추상화로 유지된다. -LA address space: +LA 주소 공간: -| 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 시작 | `0x1_0000_0000` (4 GB, 이전 VA 시작과 동일) | +| LA 공간 크기 | PE당 64 GB | +| 정렬 단위 | segment (D-LA3 참조) | -LA is PE-local: different PEs may use the same LA value; BAAW segment -tables differ → they resolve to different PAs. +LA는 PE 로컬이다: 서로 다른 PE가 동일한 LA 값을 사용할 수 있지만, +BAAW segment 테이블이 다르므로 서로 다른 PA로 해석된다. -VA infrastructure removed when LA is adopted: +LA가 채택되면 제거되는 VA 인프라: -| 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 | +| `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` component entry | Removed | +| `topology.yaml`: `pe_mmu` 컴포넌트 entry | 제거 | -#### D-LA2. Mapping mode setting +#### D-LA2. 매핑 모드 설정 -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 + hbm_pseudo_channels: 64 # 전체 pseudo channel 수 + hbm_channels_per_pe: 8 # PE당 로컬 채널 수 + hbm_channel_bw_gbs: 32.0 # 채널당 대역폭 ``` -Consumed by the graph compiler (topology builder) and BAAW -initialisation. +그래프 컴파일러(토폴로지 빌더)와 BAAW 초기화가 이 값을 소비한다. -#### D-LA3. Segment and BAAW +#### D-LA3. Segment와 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. +Segment는 LA 공간을 분할한다. 각 segment는 특정 HBM 채널 또는 채널 +그룹에 매핑된다. 텐서 deploy 시점에 런타임 할당기가 생성한다. BAAW는 +segment 테이블을 사용하여 LA → 물리 요청(들)로 해석한다. ```python @dataclass class BaawSegment: - la_base: int # segment start LA - la_size: int # segment size (bytes) + la_base: int # segment 시작 LA + la_size: int # segment 크기 (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 + # 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 lifecycle: +Segment 라이프사이클: -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. +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 resolution logic +#### D-LA4. BAAW 해석 로직 -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()`. +BAAW는 PE_DMA 내부의 프론트엔드 단계이며, 별도의 SimPy 컴포넌트가 아니다. +PE_DMA의 `handle_command()` 시작 시점에 실행되는 동기식 주소 해석 로직. -Input: `(LA, nbytes)`. Output: +입력: `(LA, nbytes)`. 출력: -- **1:1 mode**: `list[PhysicalRequest]` — one per channel. -- **n:1 mode**: single `PhysicalRequest`. +- **1:1 모드**: `list[PhysicalRequest]` — 채널당 하나. +- **n:1 모드**: 단일 `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) + pa: int # 51-bit 물리 주소 + nbytes: int # 이 요청의 전송 크기 + dst_node: str # 대상 node_id (채널 라우터 또는 집계 라우터) def resolve(self, la: int, nbytes: int) -> list[PhysicalRequest]: @@ -305,65 +290,65 @@ def resolve(self, la: int, nbytes: int) -> list[PhysicalRequest]: return requests ``` -BAAW responsibilities: +BAAW의 책임: -- Convert logical access → physical request units. -- Apply mode-dependent fan-out (1:1) or pass-through (n:1). -- Compute PA and target node. +- 논리 접근 → 물리 요청 단위로 변환. +- 모드에 따라 fan-out(1:1) 또는 pass-through(n:1) 적용. +- PA와 대상 노드 계산. -BAAW non-responsibilities: +BAAW가 하지 않는 것: -- Performing actual data movement. -- Executing NOC routing. -- Simulating bandwidth occupation (downstream components' job). +- 실제 데이터 이동 수행. +- NOC 라우팅 실행. +- 대역폭 점유 시뮬레이션 (하위 컴포넌트의 역할). -BAAW output is directly usable by the simulator's routing and resource -model without additional address decoding. +BAAW의 출력은 추가적인 주소 디코딩 없이 시뮬레이터의 라우팅·자원 +모델에서 바로 사용 가능하다. -#### D-LA5. PE_DMA `handle_command()` change +#### D-LA5. PE_DMA `handle_command()` 변경 -Current (VA-based) flow: +현재(VA 기반) 흐름: ``` DmaReadCmd.src_addr (VA) → MMU.translate(VA) → PA - → PhysAddr.decode(PA) → PhysAddr object + → PhysAddr.decode(PA) → PhysAddr 객체 → resolver.resolve(PhysAddr) → dst_node_id → router.find_path(pe_prefix, dst_node_id) → path - → 1 sub-Transaction → fabric inject + → 1 sub-Transaction → 패브릭 주입 ``` -LA-based flow: +LA 기반 흐름: ``` DmaReadCmd.src_addr (LA) → BAAW.resolve(LA, nbytes) → list[PhysicalRequest] - → for each PhysicalRequest: + → 각 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 + → sub-Transaction → 패브릭 주입 + → 모든 sub-Transaction 대기 → 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. +- MMU 참조 제거 → BAAW resolve. +- `PhysAddr.decode()` + `resolver.resolve()` → BAAW가 `dst_node`를 + 직접 반환. +- 1 요청 → 1:1 모드에서 N개의 병렬 요청. -#### D-LA6. 1:1 mode detail +#### D-LA6. 1:1 모드 상세 -- One logical access → N physical requests (N = `channels_per_pe`). +- 하나의 논리 접근 → N개의 물리 요청 (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. +- 각 요청: 완전히 해석된 51-bit PA, 특정 채널 라우터를 대상으로 함 + (`{pe_prefix}.ch_r{channel_id}`). +- 채널별 링크가 대역폭 경쟁을 모델링. +- PE_DMA가 N개의 sub-transaction을 동시에 주입. -Example: `hbm_pseudo_channels=64`, `pes_per_cube=8` → `channels_per_pe=8`. -PE0 owns ch0-7. +예: `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 @@ -375,32 +360,32 @@ BAAW segment: { channel_size: 512, } -BAAW resolve result (8 requests): +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-transactions parallel inject - per-channel router → hbm_ctrl link (channel_bw_gbs) per channel - Total effective BW = 8 × channel_bw_gbs +PE_DMA: 8개 sub-transaction 병렬 주입 + 채널별 라우터 → hbm_ctrl 링크 (channel_bw_gbs) per channel + 전체 유효 BW = 8 × channel_bw_gbs ``` -Other N values: +다른 N 값: - `hbm_pseudo_channels=32`, `pes_per_cube=8` → `channels_per_pe=4`, - 4 requests + 4 요청 - `hbm_pseudo_channels=64`, `pes_per_cube=4` → `channels_per_pe=16`, - 16 requests + 16 요청 -#### D-LA7. n:1 mode detail +#### D-LA7. n:1 모드 상세 -- 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. +- 하나의 논리 접근 → 하나의 집계 요청. +- 대상: 집계 라우터 → 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 @@ -411,111 +396,108 @@ BAAW segment: { agg_node_id: "sip0.cube0.pe0.agg_router", } -BAAW resolve result: +BAAW resolve 결과: → 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) + 집계 라우터 → hbm_ctrl 링크 (256 GB/s) ``` -#### D-LA8. Kernel model preserved +#### D-LA8. 커널 모델 보존 -- 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. +- 커널은 여전히 단일 메모리 op(`tl.load`, `tl.store`, + `tl.composite`)을 발행한다. +- LA가 커널 코드에 노출되는 주소 체계이다. +- 채널 분해·집계는 PE_DMA의 BAAW 내부에서 일어난다. +- 커널 코드는 물리 채널 정보를 절대 보지 않는다. -#### Consequences (LA model, proposed) +#### Consequences (LA 모델, 제안됨) -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. +- 1:1 vs n:1 시맨틱이 한 곳(BAAW)에 모인다. +- 커널 추상화 보존 — 커널 코드 변경 없음. +- 토폴로지 기반 정책 제어 (yaml로 모드 전환). +- 시뮬레이션 모델의 정합성·디버깅 가능성 향상. +- Segment 기반 매핑이 페이지 테이블보다 단순하며 오버헤드도 적다. -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. +- 전체 VA/MMU 코드 리팩터가 필요하다. +- 요청 생성 경로가 더 복잡 (1:1 모드에서 N 요청). +- n:1 모드에서 채널별 가시성 감소. +- VA 관련 테스트 재작성 필요. --- ## 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. +- **PA → VA**는 확장이었다. PA 모드는 PE_DMA 내부의 PageFault fallback으로 + 유지된다. 전환은 PA 코드 제거를 요구하지 않는다. +- **VA → LA**는, 채택될 경우, 공존이 아닌 대체이다. VA 인프라 제거 + 목록은 D-LA1 참조. PA fallback은 테스트를 위해 PE_DMA 내부에 직교적으로 + 유지될 수 있다. -## Alternatives Considered (LA model) +## Alternatives Considered (LA 모델) -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. +1. **VA 유지 + MMU에서 fan-out**: MMU가 채널별 PA를 반환한다. + 기각: MMU의 역할이 변환을 넘어 요청 분해까지 확장되며, 집계(n:1)를 + 표현하기 어색해진다. +2. **채널 인지 커널 API**: 커널이 채널별 load/store를 직접 호출한다. + 기각: 추상화 누출, 이식성 손실, 모든 벤치마크 재작성 필요. +3. **항상 PA (LA 없음)**: 런타임이 커널에 채널별 PA를 직접 전달한다. + 기각: 집계와 양립 불가; 변환 시점이 불명확; 채널 정보가 커널로 누출. ## Test Requirements -### VA model (current, regression) +### VA 모델 (현재, 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. +- 설치된 매핑을 따라 cross-PE / cross-cube DMA 경로. +- 측정된 레이턴시를 동반한 `MmuMapMsg` / `MmuUnmapMsg`의 패브릭 순회. +- 접근당 TLB 오버헤드 타이밍. +- PageFault fallback 경로가 PA 단독 동작을 보존하는지. -### LA model (when implemented) +### LA 모델 (구현 시) -- 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. +- 1:1 모드: 동일 논리 접근 → N개의 채널별 요청. +- n:1 모드: 동일 논리 접근 → 1개의 집계 요청. +- 동일 워크로드에 대해 두 모드 사이의 대역폭 동치. +- 1:1 모드: 채널별 경쟁이 올바르게 모델링됨. +- n:1 모드: 집계된 대역폭이 올바르게 반영됨. +- 모드 전환에 걸쳐 커널 코드가 변경되지 않음. +- BAAW segment install / uninstall 정확성. +- 별개 segment 안의 여러 텐서가 충돌하지 않음. -## Implementation Order (LA, when scheduled) +## Implementation Order (LA, 일정 잡힐 때) -1. LA type (`policy/address/la_allocator.py`). -2. BAAW segment table (`policy/address/baaw.py`). +1. LA 타입 (`policy/address/la_allocator.py`). +2. BAAW segment 테이블 (`policy/address/baaw.py`). 3. `BaawSegmentInstallMsg` (`runtime_api/kernel.py`). -4. PE_DMA BAAW integration (`components/builtin/pe_dma.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. Remove VA/MMU code. -8. Remove `pe_mmu` from `topology.yaml`; add mapping mode settings. -9. Test migration: +7. VA/MMU 코드 제거. +8. `topology.yaml`에서 `pe_mmu` 제거; 매핑 모드 설정 추가. +9. 테스트 이전: -| 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 | +| `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 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) +- 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 (메모리 주소 지정) diff --git a/docs/adr-ko/ADR-0012-api-host-io-message-schema.md b/docs/adr-ko/ADR-0012-api-host-io-message-schema.md index 07d95c5..6911ac2 100644 --- a/docs/adr-ko/ADR-0012-api-host-io-message-schema.md +++ b/docs/adr-ko/ADR-0012-api-host-io-message-schema.md @@ -1,4 +1,4 @@ -# ADR-0012: Host ↔ IO_CPU Message Schema (PA-first, PE-tagged) +# ADR-0012: Host ↔ IO_CPU 메시지 스키마 (PA-우선, PE-태깅) ## Status @@ -6,63 +6,65 @@ Accepted ## Context -Phase 0 uses a PA-first memory model (ADR-0011): +Phase 0은 PA-우선 메모리 모델을 사용한다(ADR-0011): -- memory operations use device physical addresses (PA) only, -- VA/MMU/IOMMU is not modeled. +- 메모리 연산은 디바이스 물리 주소(PA)만 사용한다, +- VA/MMU/IOMMU는 모델링하지 않는다. -The host-facing runtime API interacts with the device via the IO_CPU endpoint. -We define stable, minimal message schemas for Host ↔ IO_CPU so that: +호스트 대면 runtime API는 IO_CPU 엔드포인트를 통해 디바이스와 +상호작용한다. 다음을 보장하기 위해 Host ↔ IO_CPU에 대한 안정적이고 +최소한의 메시지 스키마를 정의한다: -- benchmarks remain stable, -- IO_CPU-internal fan-out/aggregation can evolve independently, -- completion and failure propagation is deterministic. +- 벤치마크는 안정적으로 유지된다, +- IO_CPU 내부의 팬아웃/집계는 독립적으로 진화할 수 있다, +- 완료와 실패 전파는 결정론적이다. -We also require PE-tagging (A 방식): each shard explicitly carries (sip,cube,pe) -so IO_CPU can deterministically route/fan-out without relying on PA decoding. +또한 PE-태깅(A 방식)을 요구한다: 각 샤드는 (sip,cube,pe)를 명시적으로 +운반하여, IO_CPU가 PA 디코딩에 의존하지 않고 결정론적으로 +라우팅/팬아웃할 수 있도록 한다. --- ## Decision -### D1. Contract scope +### D1. 계약 범위 -This schema is the stable contract ONLY for Host ↔ IO_CPU. +본 스키마는 오직 Host ↔ IO_CPU에 대해서만 안정적인 계약이다. -Messages beyond IO_CPU (to M_CPU, PE_CPU, schedulers, engines) are component-internal -and are NOT part of this host contract in Phase 0. +IO_CPU를 넘어선 메시지(M_CPU, PE_CPU, 스케줄러, 엔진으로 가는 것)는 +컴포넌트 내부 사항이며 Phase 0에서 이 호스트 계약의 일부가 아니다. --- -### D2. Required message set +### D2. 필수 메시지 집합 -The runtime API MUST use only these message types for Host ↔ IO_CPU: +runtime API는 Host ↔ IO_CPU에 대해 오직 다음 메시지 타입만 사용해야 한다: - MemoryWrite - MemoryRead - KernelLaunch -All operations required by benchmarks (tensor init/copy, kernel run) MUST be expressible -with these messages. +벤치마크가 필요로 하는 모든 연산(텐서 초기화/복사, 커널 실행)은 이 +메시지들로 표현 가능해야 한다. --- -### D3. Common envelope (mandatory for all requests) +### D3. 공통 envelope (모든 요청에 필수) -All Host ↔ IO_CPU requests MUST include: +모든 Host ↔ IO_CPU 요청은 반드시 다음을 포함해야 한다: - `msg_type: str` - `correlation_id: str` - - generated by the host - - used to match responses deterministically + - 호스트에서 생성 + - 응답을 결정론적으로 매칭하는 데 사용 - `request_id: str` - - unique within a correlation_id + - correlation_id 내에서 고유함 - `target_device: str` - - device identifier (e.g., "sip:0") -- `timestamp_tag: str | None` (optional) - - debug tag only; MUST NOT affect determinism + - 디바이스 식별자(예: "sip:0") +- `timestamp_tag: str | None` (선택) + - 디버그 태그 전용; 결정성에 영향을 주어서는 안 됨 -All Host ↔ IO_CPU responses MUST include: +모든 Host ↔ IO_CPU 응답은 반드시 다음을 포함해야 한다: - `correlation_id: str` - `request_id: str` @@ -70,120 +72,123 @@ All Host ↔ IO_CPU responses MUST include: --- -### D4. Completion schema (mandatory) +### D4. Completion 스키마 (필수) -`Completion` MUST have: +`Completion`은 반드시 다음을 가져야 한다: - `ok: bool` - `error_code: str | None` - `error_message: str | None` -Rules: +규칙: -- If `ok == true` then `error_code` and `error_message` MUST be null. -- If `ok == false` then `error_code` MUST be non-null. -- Completion semantics MUST be deterministic. +- `ok == true`이면 `error_code`와 `error_message`는 반드시 null이어야 한다. +- `ok == false`이면 `error_code`는 반드시 null이 아니어야 한다. +- 완료 시맨틱은 결정론적이어야 한다. --- -### D5. MemoryWrite schema (PA-first, PE-tagged) +### D5. MemoryWrite 스키마 (PA-우선, PE-태깅) -`MemoryWrite` represents a host-initiated write/initialize operation to device memory. +`MemoryWrite`는 호스트에서 시작된 디바이스 메모리 쓰기/초기화 연산을 +나타낸다. -Mandatory fields: +필수 필드: -- common envelope fields (D3) -- destination placement tags (A 방식): +- 공통 envelope 필드 (D3) +- 목적지 배치 태그 (A 방식): - `dst_sip: int` - `dst_cube: int` - `dst_pe: int` - `dst_pa: int` - - destination physical address in the destination PE's address space + - 목적지 PE의 주소 공간 내 목적지 물리 주소 - `nbytes: int` - `src_kind: "pattern" | "host_buffer_ref"` - - Phase 0 MUST support "pattern" + - Phase 0은 반드시 "pattern"을 지원해야 한다 - `pattern: Pattern | None` - - required if `src_kind == "pattern"` + - `src_kind == "pattern"`인 경우 필수 -`Pattern` (Phase 0 mandatory support): +`Pattern` (Phase 0 필수 지원): - `pattern_kind: "zero" | "fill_u8" | "fill_u16" | "fill_u32" | "fill_fp16" | "fill_fp32"` - `value: number | None` - - required for fill_*; ignored for zero + - fill_*에 필요; zero에서는 무시됨 -Optional fields: +선택 필드: -- `dst_mem_kind: "HBM" | "TCM" | "AUTO"` (default "AUTO") +- `dst_mem_kind: "HBM" | "TCM" | "AUTO"` (기본값 "AUTO") - `debug_label: str | None` -Notes: +비고: -- This message MUST NOT embed bulk tensor data in Phase 0. -- All latency MUST come from explicit graph traversal and modeled components. +- 이 메시지는 Phase 0에서 대용량 텐서 데이터를 임베드해서는 안 된다. +- 모든 레이턴시는 명시적인 그래프 순회 및 모델링된 컴포넌트로부터 + 발생해야 한다. --- -### D6. MemoryRead schema (PA-first, PE-tagged) +### D6. MemoryRead 스키마 (PA-우선, PE-태깅) -`MemoryRead` represents a host-initiated read from device memory. +`MemoryRead`는 호스트에서 시작된 디바이스 메모리 읽기를 나타낸다. -Mandatory fields: +필수 필드: -- common envelope fields (D3) -- source placement tags (A 방식): +- 공통 envelope 필드 (D3) +- 소스 배치 태그 (A 방식): - `src_sip: int` - `src_cube: int` - `src_pe: int` - `src_pa: int` - `nbytes: int` -Optional fields: +선택 필드: -- `dst_kind: "host_sink" | "discard"` (default "host_sink") +- `dst_kind: "host_sink" | "discard"` (기본값 "host_sink") - `debug_label: str | None` -Response payload: +응답 페이로드: -- actual bytes are NOT required in Phase 0 (latency/traces focus) -- implementations MAY return lightweight stats or hashes later via a new ADR +- Phase 0에서는 실제 바이트는 필요하지 않다(레이턴시/트레이스 중심) +- 구현은 추후 새로운 ADR을 통해 가벼운 통계나 해시를 반환할 수 있다 --- -### D7. KernelLaunch schema (PA-first, PE-tagged shards) +### D7. KernelLaunch 스키마 (PA-우선, PE-태깅된 샤드) -`KernelLaunch` represents launching a kernel on a target device via IO_CPU. +`KernelLaunch`는 IO_CPU를 통해 대상 디바이스에서 커널을 런치하는 것을 +나타낸다. -Mandatory fields: +필수 필드: -- common envelope fields (D3) +- 공통 envelope 필드 (D3) - `kernel_ref: KernelRef` - `args: list[KernelArg]` -`KernelRef` MUST have: +`KernelRef`는 반드시 다음을 가져야 한다: - `name: str` - `kind: "deployed" | "builtin"` -- `deploy_pa: int | None` — PA where kernel binary was deployed (required for "deployed") -- `deploy_sip: int` — SIP where binary resides -- `deploy_cube: int` — cube where binary resides -- `deploy_pe: int` — PE where binary resides -- `nbytes_code: int` — kernel binary size (for BW modeling) +- `deploy_pa: int | None` — 커널 바이너리가 배치된 PA("deployed"에 필수) +- `deploy_sip: int` — 바이너리가 위치한 SIP +- `deploy_cube: int` — 바이너리가 위치한 큐브 +- `deploy_pe: int` — 바이너리가 위치한 PE +- `nbytes_code: int` — 커널 바이너리 크기(BW 모델링용) -Kernel binaries MUST be pre-deployed to device memory via MemoryWrite. -KernelLaunch MUST NOT embed kernel source code or IR in the launch message. +커널 바이너리는 MemoryWrite를 통해 디바이스 메모리에 사전 배치되어야 한다. +KernelLaunch는 커널 소스 코드나 IR을 런치 메시지에 임베드해서는 안 된다. -`KernelArg` supports tensor args by PA mapping and scalars by value. +`KernelArg`는 PA 매핑을 통한 텐서 인자와 값을 통한 스칼라 인자를 지원한다. -Tensor arg (mandatory): +텐서 인자 (필수): - `arg_kind: "tensor"` - `tensor_pa_map: TensorPAMap` -`TensorPAMap` MUST have: +`TensorPAMap`은 반드시 다음을 가져야 한다: - `shards: list[TensorShard]` -`TensorShard` MUST have (A 방식 강제): +`TensorShard`는 반드시 다음을 가져야 한다 (A 방식 강제): - `sip: int` - `cube: int` @@ -192,42 +197,43 @@ Tensor arg (mandatory): - `nbytes: int` - `offset_bytes: int` -Scalar arg (mandatory): +스칼라 인자 (필수): - `arg_kind: "scalar"` - `dtype: "i32" | "i64" | "fp16" | "fp32" | "bool"` - `value: number | bool` -Optional KernelLaunch fields: +KernelLaunch 선택 필드: - `grid: dict | None` - `meta: dict | None` -- `failure_policy: "fail_fast" | "collect_all"` (default "fail_fast") +- `failure_policy: "fail_fast" | "collect_all"` (기본값 "fail_fast") - `debug_label: str | None` -Notes: +비고: -- KernelLaunch MUST NOT embed bulk tensor data. -- KernelLaunch MUST be submitted only to the IO_CPU endpoint. -- IO_CPU MUST fan-out work internally using the shard (sip,cube,pe) tags. +- KernelLaunch는 대용량 텐서 데이터를 임베드해서는 안 된다. +- KernelLaunch는 오직 IO_CPU 엔드포인트에만 제출되어야 한다. +- IO_CPU는 샤드의 (sip,cube,pe) 태그를 사용하여 내부적으로 작업을 + 팬아웃해야 한다. --- ## Verification Notes -Tests SHOULD validate: +테스트는 다음을 검증해야 한다: -- schema validation rejects missing mandatory fields, -- deterministic correlation/response matching, -- MemoryWrite/Read/KernelLaunch produce explicit hop traces, -- all routed requests incur latency > 0. +- 스키마 검증이 필수 필드 누락을 거부함, +- 결정론적 correlation/응답 매칭, +- MemoryWrite/Read/KernelLaunch가 명시적인 홉 트레이스를 생성함, +- 라우팅된 모든 요청은 레이턴시 > 0을 가짐. --- ## Links -- 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) +- ADR-0011 (메모리 주소 지정 — PA / VA / LA) +- ADR-0007 (runtime_api와 sim_engine 경계) +- ADR-0009 (커널 실행 팬아웃/집계) +- ADR-0013 (검증 전략 — V1 메시지 스키마 검증) - SPEC R2, R7, R8 diff --git a/docs/adr-ko/ADR-0013-ver-verification-strategy.md b/docs/adr-ko/ADR-0013-ver-verification-strategy.md new file mode 100644 index 0000000..ecdf635 --- /dev/null +++ b/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 (커널 실행 시맨틱) diff --git a/docs/adr-ko/ADR-0014-dev-pe-pipeline-execution-model.md b/docs/adr-ko/ADR-0014-dev-pe-pipeline-execution-model.md index ccb63f3..93ce393 100644 --- a/docs/adr-ko/ADR-0014-dev-pe-pipeline-execution-model.md +++ b/docs/adr-ko/ADR-0014-dev-pe-pipeline-execution-model.md @@ -1,4 +1,4 @@ -# ADR-0014: PE Pipeline Execution Model +# ADR-0014: PE 파이프라인 실행 모델 ## Status @@ -6,153 +6,149 @@ Accepted ## Context -This ADR defines the PE-internal kernel execution model: +본 ADR은 PE 내부 커널 실행 모델을 정의한다: -- 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 내부 컴포넌트의 역할 분담 +- 명령 디스패치 경로 (simple / composite / epilogue를 포함한 multi-op composite) +- TileToken 기반 자가-라우팅 파이프라인 (스케줄러는 디스패치와 완료 처리만 담당) +- 레지스터 파일을 매개로 한 TCM 중심 데이터플로우 +- 엔진 자원 모델 +- 관측 가능성 및 트레이스 계약 +- 토폴로지 표현 -PE-internal structure (7 components in scope; 2 cross-referenced): +PE 내부 구조 (본 ADR 범위 7개 컴포넌트 + 외부 참조 2개): - `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 + `pe_tcm` — 본 ADR에서 정의 +- `pe_mmu` — VA 모델, ADR-0011 D-VA에서 정의 +- `pe_ipcq` — 집합 통신, ADR-0023에서 정의 -The goal is a deterministic, trace-friendly execution contract that keeps -each block independently swappable. +목표는 결정론적이고 트레이스 친화적인 실행 계약을 통해 각 블록이 독립적으로 +교체 가능하도록 유지하는 것이다. ## Decision -### D1. PE-internal component roles +### D1. PE 내부 컴포넌트의 역할 **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 명령을 생성하여 `PE_SCHEDULER`에 제출한다 (`PeInternalTxn`을 통해). +- 엔진 큐에 직접 작업을 넣지 않는다. **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 내부의 유일한 디스패처. +- `PE_CPU`로부터 명령을 수신한다. 명령 타입별 디스패치: + - Simple 명령 (`DmaReadCmd`, `DmaWriteCmd`, `GemmCmd`, `MathCmd`) + → 대상 엔진으로 직접 전달. + - `CompositeCmd` → `TilePlan`을 생성하고, 단일 `_feed_loop`를 통해 + 파이프라인에 타일을 공급한다 (D6). +- composite 내부의 stage-to-stage 체이닝에는 관여하지 않는다; + 이는 토큰 자가-라우팅(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). +- 큐브 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 ↔ 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. +- TCM ↔ 레지스터 파일 전송 유닛. +- 레지스터 파일 접근 시맨틱을 컴퓨트 엔진으로부터 격리하여 + GEMM/MATH가 순수한 컴퓨트 컴포넌트로 유지되도록 한다. +- BW 기반 레이턴시 모델; TCM 접근 경합은 `PE_TCM`의 BW 자원을 통해 + 자연스럽게 직렬화된다. **PE_GEMM** -- MAC array. Reads operands from the register file; writes results to - the register file. Does not touch `PE_TCM` directly. +- MAC 어레이. 레지스터 파일에서 피연산자를 읽고, 결과를 레지스터 파일에 + 쓴다. `PE_TCM`에 직접 접근하지 않는다. **PE_MATH** -- Element-wise / reduction / SIMD unit. Reads / writes the register file. +- 원소별 / 리덕션 / SIMD 유닛. 레지스터 파일을 읽고 쓴다. **PE_TCM** -- Tightly-coupled scratchpad with BW-serialized access. Two logical - regions partitioned by ownership (see D5). +- BW로 직렬화된 접근을 갖는 tightly-coupled 스크래치패드. 소유권에 따라 + 두 개의 논리 영역으로 분할된다 (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 +- `pe_mmu` — 접근마다 VA→PA 변환 (ADR-0011 D-VA). +- `pe_ipcq` — 집합 통신 링 버퍼와 피어 엔드포인트 메타데이터 (ADR-0023). -### D2. Command lifecycle and queues +### D2. 명령 생명주기와 큐 -`PE_SCHEDULER` maintains three logical structures: +`PE_SCHEDULER`는 세 개의 논리적 구조를 유지한다: -**SubmissionQueue** — written by `PE_CPU`; consumed by the scheduler. +**SubmissionQueue** — `PE_CPU`가 쓰고, 스케줄러가 소비한다. -**InflightTable** — owned and mutated only by `PE_SCHEDULER`; tracks -expanded sub-commands, dependency state, engine assignment, and -completion status. +**InflightTable** — `PE_SCHEDULER`만 소유하고 변경한다; 전개된 sub-command, +의존성 상태, 엔진 할당, 완료 상태를 추적한다. -**CompletionQueue** — written by `PE_SCHEDULER`; holds final completion -records. +**CompletionQueue** — `PE_SCHEDULER`가 쓴다; 최종 완료 레코드를 보관한다. -**Single-writer rule**: only `PE_SCHEDULER` mutates command completion -state. Engines report completion via explicit events / messages -consumed by the scheduler. +**Single-writer 규칙**: `PE_SCHEDULER`만이 명령 완료 상태를 변경한다. +엔진은 명시적 이벤트 / 메시지로 완료를 보고하며, 이는 스케줄러가 +소비한다. -**Command completion**: when all sub-commands complete, `PE_SCHEDULER` -publishes a completion record. +**명령 완료**: 모든 sub-command가 완료되면 `PE_SCHEDULER`가 완료 레코드를 +발행한다. -### D3. Dispatch modes +### D3. 디스패치 모드 -#### D3.1 Simple command +#### D3.1 Simple 명령 -A simple command expands to exactly one engine sub-command: +simple 명령은 정확히 하나의 엔진 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) +#### D3.2 Composite 명령 (단일-op 타일 파이프라인) -The default `CompositeCmd` runs a single compute op as a tile-pipelined -sequence: +기본 `CompositeCmd`는 단일 컴퓨트 op를 타일 파이프라인 시퀀스로 실행한다: ```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`. +`PE_SCHEDULER`는 DMA 페이로드를 하드웨어 타일로 분할하고, 단조 증가하는 +`tile_id`를 갖는 `TileToken`을 타일마다 하나씩 발행한다. -Tile dependency (within one tile `t`): +타일 의존성 (단일 타일 `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): +엔진 자원이 허용하는 한 타일 간 오버랩이 허용된다 +(D4가 제약을 규정): ```text DMA_READ(t+1) ∥ COMPUTE(t) DMA_WRITE(t-1) ∥ COMPUTE(t) ``` -#### D3.3 Multi-op composite (head + epilogue with scope) +#### D3.3 Multi-op composite (스코프를 갖는 head + epilogue) -A `CompositeCmd` MAY carry `ops: tuple[OpSpec, ...]` to express a -multi-op pipeline: +`CompositeCmd`는 `ops: tuple[OpSpec, ...]`를 운반하여 multi-op +파이프라인을 표현할 수 있다: ```python @dataclass(frozen=True) @@ -162,48 +158,46 @@ class OpSpec: ... ``` -- `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. +- `ops[0]` (head)이 타일 기하 구조를 정의한다 (예: head GEMM이 M/K/N + 분할을 결정). +- `ops[1:]` (epilogue)는 후속 stage이며 `scope`에 따라 실행 빈도가 + 결정된다: + - `per_k_tile` — 모든 K-리덕션 스텝마다. + - `per_output_tile` — 출력 타일당 한 번. + - `once` — 커널당 한 번. -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). +크로스-엔진 체인(예: GEMM head → MATH epilogue)은 자연스럽다 — +각 stage는 토큰 자가-라우팅(D6)을 통해 디스패치되므로, GEMM과 MATH는 +동일한 컴퓨트 슬롯(D4)을 공유하더라도 동일 composite 내에서 직렬적으로 +참여한다. -The empty-`ops` form is the legacy single-op path. +비어 있는 `ops` 형식은 레거시 단일-op 경로이다. -### D4. Engine resource model +### D4. 엔진 자원 모델 -**DMA engine**: +**DMA 엔진**: - `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. +- 두 채널은 동시에 실행된다 (READ ∥ WRITE 허용). +- 채널 내부에서는 요청이 직렬화된다 (READ ∥ READ 불가; WRITE도 동일). +- `vc_comm`은 IPCQ 트래픽을 위한 직교 채널로 ADR-0023 D8에서 정의됨 — + 본 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. +- `accel_slot`: `PE_GEMM`과 `PE_MATH`가 공유하는 `simpy.Resource(capacity=1)`. +- PE 내에서 동시에 최대 한 개의 컴퓨트 op만 실행된다. +- Multi-op composite 체인(D3.3)은 이 슬롯을 통해 컴퓨트 stage를 직렬로 + 실행한다; 토큰 자가-라우팅(D6)이 이전 컴퓨트가 슬롯을 해제한 후에만 + 다음 stage가 시작되도록 보장한다. -**Engine completion**: each engine emits a completion event consumed by -the scheduler / `PipelineContext` (D6). +**엔진 완료**: 각 엔진은 완료 이벤트를 발행하며, 이는 스케줄러 / +`PipelineContext`(D6)가 소비한다. -### D5. Dataflow +### D5. 데이터플로우 -**Input path (HBM source)**: +**입력 경로 (HBM 소스)**: ```text HBM → cube NOC → PE_DMA (DMA_READ) → PE_TCM @@ -211,60 +205,58 @@ PE_TCM → PE_FETCH_STORE → Register File Register File → PE_GEMM | PE_MATH ``` -**Input path (shared SRAM source)**: +**입력 경로 (공유 SRAM 소스)**: ```text Shared SRAM → cube NOC → PE_DMA (DMA_READ) → PE_TCM PE_TCM → PE_FETCH_STORE → Register File ``` -**Output path (HBM destination)**: +**출력 경로 (HBM 목적지)**: ```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. +GEMM/MATH는 `PE_TCM`에 직접 접근하지 않는다 — `PE_FETCH_STORE`가 +TCM↔레지스터 파일의 유일한 게이트웨이이다. 이를 통해 TCM BW 경합이 +명시적으로 드러나며, fetch 유닛 정책(예: 프리패치)을 컴퓨트 엔진과 +독립적으로 교체할 수 있다. -#### D5.1 PE_TCM partitioning +#### D5.1 PE_TCM 분할 -`PE_TCM` is split into two logical regions: +`PE_TCM`은 두 개의 논리 영역으로 분할된다: **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. +- `PE_SCHEDULER`가 단독으로 소유한다. +- composite 명령의 타일 버퍼를 보관한다. +- `PE_SCHEDULER`가 이 영역을 분할하고, DMA_READ / COMPUTE / DMA_WRITE + stage마다 버퍼를 할당하며, 입출력 분리를 보장하고, 타일-버퍼 수명을 + 관리한다. **AllocatableTCM** -- General-purpose region managed by `PEMemAllocator`. -- Used for host / DP-visible allocations. +- `PEMemAllocator`가 관리하는 범용 영역. +- 호스트 / DP 가시 할당에 사용된다. -**Visibility rule (hard isolation)**: `PEMemAllocator` MUST NOT see or -allocate inside `SchedulerReservedTCM`. The reserved region is excluded -from allocator-managed ranges by construction. +**가시성 규칙 (강한 격리)**: `PEMemAllocator`는 `SchedulerReservedTCM`을 +보거나 그 내부에 할당해서는 안 된다. 예약 영역은 구성 시점에 할당자가 +관리하는 범위에서 제외된다. -**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. +- 타일이 활성 수명 동안 `SchedulerReservedTCM` 내부의 입력 버퍼와 출력 + 버퍼는 겹쳐서는 안 된다. +- 타일 버퍼는 해당 `DMA_WRITE`가 완료될 때까지 유효하다. +- 버퍼 재사용은 소비하는 타일의 수명이 끝난 후에만 허용된다. -### D6. TileToken self-routing pipeline +### D6. TileToken 자가-라우팅 파이프라인 -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`: +composite의 stage-to-stage 진행은 스케줄러를 거치지 **않고** 일어난다. +각 컴포넌트는 토큰의 `plan`을 사용해 토큰을 다음 stage의 컴포넌트로 +직접 전달한다: ```text Scheduler → DMA → Fetch → GEMM → Math (epi) → Store → DMA_WB → (complete) @@ -272,8 +264,8 @@ Scheduler → DMA → Fetch → GEMM → Math (epi) → Store → DMA_WB → (co PipelineContext.complete_tile() ``` -This mirrors real-HW done-wire chains. The scheduler handles only -**initial dispatch + completion aggregation**. +이는 실제 HW의 done-wire 체인을 반영한다. 스케줄러는 **초기 디스패치 + +완료 집계**만 담당한다. #### TilePlan / Stage @@ -311,12 +303,12 @@ class TileToken: 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()`. +단일 소유자 불변식: 토큰은 한 시점에 정확히 한 컴포넌트가 소유한다. +생명주기: 스케줄러가 `stage_idx=0`으로 생성 → 컴포넌트 `_process()` → +`stage_idx` 증가 → 다음 stage의 `in_port`에 put → 마지막 stage가 +`pipeline_ctx.complete_tile()` 호출. -#### PipelineContext (exactly-once completion) +#### PipelineContext (정확히 한 번 완료) ```python @dataclass @@ -332,25 +324,24 @@ class PipelineContext: 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). +각 타일의 마지막 stage는 `complete_tile()`을 정확히 한 번 호출해야 +한다. 중복 호출은 버그이다 (SimPy `Event`는 최대 한 번만 succeed +가능). -#### Feed ordering +#### Feed 순서 -`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.** +`PE_SCHEDULER`는 `_pending_feeds` FIFO를 소비하는 `_feed_loop` 프로세스를 +정확히 하나 갖는다. composite 명령은 제출 순서대로 인큐되며, 한 명령의 +타일 feed는 다음 명령의 feed가 시작되기 전에 완료까지 실행된다. +**명령 간 타일-feed 인터리빙은 허용되지 않는다.** -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. +단일 명령의 타일들 내부에서는 다운스트림 파이프라인 오버랩이 자연스럽게 +발생한다 — 이전 타일이 후행 stage를 진행하는 동안 feeder는 남은 타일을 +첫 stage 큐로 계속 푸시한다 (SimPy Store 백프레셔가 흐름 제어를 +관장한다). 첫 stage 큐가 가득 차면 feeder만 블록되며, 스케줄러 워커의 +inbox 처리는 계속된다. -#### Token routing pattern (base class) +#### 토큰 라우팅 패턴 (기본 클래스) ```python def _pipeline_worker(self, env): @@ -367,12 +358,11 @@ def _pipeline_worker(self, env): token.pipeline_ctx.complete_tile() ``` -Each component implements only `_process()`; chaining lives in the -base class. +각 컴포넌트는 `_process()`만 구현한다; 체이닝은 기본 클래스에 존재한다. -### D7. Observability and trace contract +### D7. 관측 가능성 및 트레이스 계약 -The simulator emits deterministic trace events: +시뮬레이터는 결정론적 트레이스 이벤트를 발행한다: - `command_submitted` - `sub_command_dispatched` @@ -381,11 +371,11 @@ The simulator emits deterministic trace events: - `tile_ready` - `command_complete` -For identical inputs, trace ordering MUST be deterministic. +동일한 입력에 대해 트레이스 순서는 결정론적이어야 한다. -### D8. Topology representation +### D8. 토폴로지 표현 -PE-internal components are declared in `cube.pe_template`: +PE 내부 컴포넌트는 `cube.pe_template`에 선언된다: ```yaml pe_template: @@ -416,36 +406,36 @@ pe_template: 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). +템플릿은 PE마다 한 번 인스턴스화된다. PE 인스턴스는 `cube.pe_layout` +(코너 배치)으로부터 파생된다. 외부 연결성(PE_DMA ↔ cube NoC ↔ HBM 등)은 +큐브 수준에서 모델링된다 (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. +- 각 블록이 독립적인 토폴로지 노드이다 — DI(ADR-0015)를 통해 개별 + 교체 가능하다. +- PE 내부 구조가 토폴로지 그래프에 가시화된다. +- 컴포넌트는 자신의 다운스트림을 알지 못한다 — plan 기반 라우팅이 + 유연성을 제공한다 (예: epilogue 체인에 스케줄러 변경이 불필요). +- DMA와 컴퓨트가 SimPy Store 백프레셔를 통해 자연스럽게 오버랩된다. +- Multi-op composite가 융합 연산(예: GEMM + bias_add)을 엔진 수준 + 결합 없이 표현한다. +- TCM 접근 경합이 현실적이다 — `PE_FETCH_STORE`가 TCM↔RF의 유일한 + 게이트웨이이다. ### 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). +- PE 내부 컴포넌트 수가 더 거친 모델보다 많다 (기본 7개 + 외부 참조 + 2개) — 더 많은 토폴로지 노드/엣지. +- PE 내부 토큰 전달이 트레이스에 명시적으로 드러난다 (HW 충실도와의 + 허용 가능한 trade-off). ## 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) +- ADR-0011 D-VA (PE_MMU 컴포넌트, VA 변환) +- ADR-0015 D4 (컴포넌트 포트/와이어 모델) +- ADR-0020 (greenlet 커널 실행 / two-pass) +- ADR-0023 (PE_IPCQ + PE_DMA 가상 채널) - SPEC R3, R4 diff --git a/docs/adr-ko/ADR-0015-dev-component-port-wire-model.md b/docs/adr-ko/ADR-0015-dev-component-port-wire-model.md index 5f999af..561db20 100644 --- a/docs/adr-ko/ADR-0015-dev-component-port-wire-model.md +++ b/docs/adr-ko/ADR-0015-dev-component-port-wire-model.md @@ -1,4 +1,4 @@ -# ADR-0015: Component Port/Wire Model and Fabric Routing +# ADR-0015: 컴포넌트 포트/와이어 모델과 패브릭 라우팅 ## Status @@ -6,46 +6,44 @@ Accepted ## Context -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. +현실적인 하드웨어 모델링 — 큐, 경합, fan-out — 을 위해서는 +컴포넌트가 패브릭 순회를 소유하고, 시뮬레이션 엔진은 초기화와 완료 +관측만 처리해야 한다. 컴포넌트 간의 직접 메서드 호출이나 엔진 내부의 +경로 탐색은 큐잉과 경합 시맨틱을 무력화한다. -This ADR defines: +본 ADR은 다음을 정의한다: -- 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 engine's reduced role (wire init + completion observation only), -- M_CPU.DMA as an internal subcomponent of M_CPU. +- 컴포넌트가 타입드 포트 큐를 통해 통신하는 방식, +- 전파 지연을 모델링하는 방식 (BW 점유를 포함한 와이어 프로세스), +- Memory R/W (M_CPU 우회)와 Kernel Launch (M_CPU 경유)의 패브릭 경로, +- 엔진의 축소된 역할 (와이어 초기화 + 완료 관측만), +- M_CPU의 내부 서브컴포넌트로서의 M_CPU.DMA. --- ## Decision -### D1. Component port model +### D1. 컴포넌트 포트 모델 -Each component has typed input/output ports modeled as SimPy Stores: +각 컴포넌트는 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 ``` -Ports are created at engine initialization based on graph edges. -Each directed edge (src → dst) results in: +포트는 그래프 엣지를 기반으로 엔진 초기화 시 생성된다. +각 유향 엣지(src → dst)는 다음을 생성한다: -- `src.out_ports[dst]` — the sending end -- `dst.in_ports[src]` — the receiving end +- `src.out_ports[dst]` — 송신측 +- `dst.in_ports[src]` — 수신측 --- -### D2. Wire process (propagation delay + BW occupancy) +### D2. 와이어 프로세스 (전파 지연 + BW 점유) -For each directed edge (src, dst) in the topology graph, a SimPy wire process -models propagation delay and BW occupancy: +토폴로지 그래프의 각 유향 엣지 (src, dst)에 대해 SimPy 와이어 프로세스가 +전파 지연과 BW 점유를 모델링한다: ```python def wire_process(env, out_port, in_port, delay_ns, bw_gbs): @@ -63,39 +61,39 @@ def wire_process(env, out_port, in_port, delay_ns, bw_gbs): yield in_port.put(cmd) ``` -Wire processes are started at engine initialization. -Each directed edge maintains an `available_at` timestamp tracking when the link -becomes free for the next transaction. When a transaction occupies a link, the -next transaction on the same directed link must wait until occupancy clears -(back-to-back serialization). TX and RX directions are independent (separate -wire processes with separate `available_at` state). +와이어 프로세스는 엔진 초기화 시점에 시작된다. +각 유향 엣지는 링크가 다음 트랜잭션을 위해 비워지는 시점을 추적하는 +`available_at` 타임스탬프를 유지한다. 한 트랜잭션이 링크를 점유하는 동안, +동일 유향 링크의 다음 트랜잭션은 점유가 해제될 때까지 대기해야 한다 +(연속 직렬화). TX와 RX 방향은 독립적이다 (각각의 `available_at` 상태를 +갖는 별개의 와이어 프로세스). --- -### D3. Engine role (reduced) +### D3. 엔진 역할 (축소) -The simulation engine MUST: +시뮬레이션 엔진은 다음을 수행해야 한다: -- wire components at initialization (create port Stores, start wire processes), -- identify the entry component for each request type (PCIE_EP), -- put the request into the entry component's in_port, -- wait for a completion event. +- 초기화 시점에 컴포넌트 와이어링 (포트 Store 생성, 와이어 프로세스 시작), +- 각 요청 타입별 진입 컴포넌트 식별 (PCIE_EP), +- 진입 컴포넌트의 in_port에 요청을 put, +- 완료 이벤트 대기. -The simulation engine MUST NOT: +시뮬레이션 엔진은 다음을 해서는 안 된다: -- walk the topology path during request execution, -- call component `run()` methods directly, -- track per-hop latency or decompose fan-out. +- 요청 실행 중 토폴로지 경로 탐색, +- 컴포넌트 `run()` 메서드 직접 호출, +- hop별 레이턴시 추적이나 fan-out 분해. --- -### D4. Fabric paths for Memory R/W and Kernel Launch +### D4. Memory R/W와 Kernel Launch의 패브릭 경로 -Memory R/W and Kernel Launch use **different** fabric paths. -Memory operations bypass M_CPU and route directly to HBM via the crossbar. -Kernel Launch routes through M_CPU for PE fan-out. +Memory R/W와 Kernel Launch는 **서로 다른** 패브릭 경로를 사용한다. +메모리 연산은 M_CPU를 우회하여 크로스바를 통해 직접 HBM으로 라우팅된다. +Kernel Launch는 PE fan-out을 위해 M_CPU를 경유한다. -**Memory R/W forward path (pcie_ep → hbm_ctrl, M_CPU bypass):** +**Memory R/W forward 경로 (pcie_ep → hbm_ctrl, M_CPU 우회):** ```text pcie_ep → io_noc → io_ucie @@ -103,14 +101,14 @@ pcie_ep → io_noc → io_ucie → target cube: ucie_in → router mesh → hbm_ctrl ``` -**Memory R/W completion path:** +**Memory R/W 완료 경로:** ```text hbm_ctrl → router mesh → [transit cubes: ucie → router mesh → ucie] → io_ucie → io_noc → pcie_ep ``` -**Kernel Launch forward path (pcie_ep → io_cpu → M_CPU → PE):** +**Kernel Launch forward 경로 (pcie_ep → io_cpu → M_CPU → PE):** ```text pcie_ep → io_noc → io_cpu → io_noc → io_ucie @@ -118,7 +116,7 @@ pcie_ep → io_noc → io_cpu → io_noc → io_ucie → target cube: ucie_in → noc → M_CPU → PE[0..n] (parallel fan-out) ``` -**Kernel Launch completion path:** +**Kernel Launch 완료 경로:** ```text PE[0..n] all complete → M_CPU (aggregation) @@ -126,77 +124,79 @@ PE[0..n] all complete → M_CPU (aggregation) → io_ucie → io_noc → io_cpu → io_noc → pcie_ep ``` -**Rationale for M_CPU bypass on Memory R/W:** +**Memory R/W가 M_CPU를 우회하는 근거:** -Memory write/read operations do not require command interpretation or PE -dispatch — they are direct data transfers to/from HBM. Routing through M_CPU -would add unnecessary overhead (5ns) without functional benefit. The io_noc -inside the IO chiplet handles the routing decision: memory operations go -directly to cube fabric, while kernel launches are forwarded to io_cpu first. +메모리 write/read 연산은 명령 해석이나 PE 디스패치가 필요하지 않다 — +HBM으로의/로부터의 직접 데이터 전송이다. M_CPU를 경유하면 기능적 이득 +없이 불필요한 오버헤드(5ns)를 추가한다. IO 칩렛 내부의 io_noc가 라우팅 +결정을 처리한다: 메모리 연산은 큐브 패브릭으로 직접 가고, kernel +launch는 io_cpu로 먼저 전달된다. --- -### D5. M_CPU.DMA is an internal subcomponent of M_CPU +### D5. M_CPU.DMA는 M_CPU의 내부 서브컴포넌트이다 -M_CPU.DMA is NOT a separate topology node. -It is an internal subcomponent owned by the M_CPU component implementation. +M_CPU.DMA는 별개의 토폴로지 노드가 아니다. +M_CPU 컴포넌트 구현이 소유하는 내부 서브컴포넌트이다. -M_CPU.DMA: +M_CPU.DMA는: -- owns the DMA READ and DMA WRITE queues (capacity=1 each, per ADR-0014 D4), -- issues memory requests over the NOC to hbm_ctrl, -- receives completion from hbm_ctrl via the NOC, -- reports completion to M_CPU, -- is created and managed inside M_CPU's `__init__` and `run()`. +- DMA READ 및 DMA WRITE 큐를 소유한다 (각 capacity=1, ADR-0014 D4), +- NoC를 통해 hbm_ctrl에 메모리 요청을 발행한다, +- NoC를 통해 hbm_ctrl로부터 완료를 수신한다, +- M_CPU에 완료를 보고한다, +- M_CPU의 `__init__`과 `run()` 내부에서 생성·관리된다. -M_CPU.DMA does not appear as a node in the compiled topology graph. +M_CPU.DMA는 컴파일된 토폴로지 그래프에서 노드로 나타나지 않는다. --- -### D6. Transit cube forwarding +### D6. Transit 큐브 포워딩 -A cube that is not the target of a memory or kernel request acts as a transit node. -Transit cubes forward requests without consuming them: +메모리나 커널 요청의 대상이 아닌 큐브는 transit 노드로 동작한다. +Transit 큐브는 요청을 소비하지 않고 포워딩한다: ```text ucie_in (from upstream) → noc → ucie_out (to downstream) ``` -Transit forwarding is implemented entirely within the ucie_in component. -The noc and ucie_out components in a transit cube forward the packet without modification. +Transit 포워딩은 ucie_in 컴포넌트 내부에서 전적으로 구현된다. +transit 큐브의 noc와 ucie_out 컴포넌트는 패킷을 수정 없이 포워딩한다. --- -### D7. _formula_latency is preserved as a lower-bound cross-check +### D7. _formula_latency는 하한 교차 검증 용도로 유지된다 -The path-based formula latency function (`_formula_latency`) is preserved in the engine -as a lower bound for correctness verification. +경로 기반 공식 레이턴시 함수(`_formula_latency`)는 정확성 검증을 위한 +하한값으로 엔진 내에 유지된다. -Invariant: +불변식: - Phase 0: `_formula_latency == component model total_ns` -- Phase 1+: `_formula_latency <= component model total_ns` (contention adds queueing) +- Phase 1+: `_formula_latency <= component model total_ns` (경합이 + 큐잉을 추가) -This function is independent of the port/wire model and requires only the topology graph. -It is used for shard comparison in `_route_kernel` and as a regression guard. +이 함수는 포트/와이어 모델과 독립적이며 토폴로지 그래프만 요구한다. +`_route_kernel`의 샤드 비교와 회귀 가드로 사용된다. --- ## Consequences -- Components model realistic hardware behavior (queues, contention, fan-out). -- Propagation delay is modeled accurately per edge. -- Engine is decoupled from routing policy. -- Component implementations remain swappable via DI (ADR-0007 D3). +- 컴포넌트가 현실적인 하드웨어 동작(큐, 경합, fan-out)을 모델링한다. +- 전파 지연이 엣지마다 정확하게 모델링된다. +- 엔진이 라우팅 정책으로부터 분리된다. +- 컴포넌트 구현이 DI(ADR-0007 D3)를 통해 교체 가능하게 유지된다. --- ## Links -- 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) +- 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 (이러한 메커니즘 위에 구축된 레이턴시 모델 가정) diff --git a/docs/adr-ko/ADR-0016-dev-iochiplet-noc-and-memory-path.md b/docs/adr-ko/ADR-0016-dev-iochiplet-noc-and-memory-path.md index cb1e281..8e706fe 100644 --- a/docs/adr-ko/ADR-0016-dev-iochiplet-noc-and-memory-path.md +++ b/docs/adr-ko/ADR-0016-dev-iochiplet-noc-and-memory-path.md @@ -1,4 +1,4 @@ -# ADR-0016: IOChiplet NOC and Memory Data Path +# ADR-0016: IOChiplet NoC와 메모리 데이터 경로 ## Status @@ -6,72 +6,73 @@ Accepted ## Context -ADR-0003 D2 defines IO chiplets as SIP-level components providing PCIe-EP and -IO_CPU interfaces, but does not specify internal routing within the IO chiplet. -ADR-0015 D4 was updated to document the M_CPU bypass for Memory R/W, but the -IO chiplet's internal NOC architecture that enables this routing was not -formally documented. +ADR-0003 D2는 IO chiplet을 PCIe-EP 및 IO_CPU 인터페이스를 제공하는 SIP +레벨 컴포넌트로 정의하지만, IO chiplet 내부의 라우팅은 명세하지 않는다. +ADR-0015 D4는 Memory R/W에 대한 M_CPU 우회를 문서화하도록 갱신되었지만, +이 라우팅을 가능하게 하는 IO chiplet의 내부 NoC 아키텍처는 형식적으로 +문서화되지 않았다. -The IO chiplet needs an internal routing fabric (io_noc) to: +IO chiplet은 다음을 위해 내부 라우팅 패브릭(io_noc)을 필요로 한다: -- connect pcie_ep, io_cpu, and per-cube UCIe PHY ports -- route memory operations (MemoryWrite/Read) directly to cube fabric without - passing through io_cpu -- route kernel launch commands through io_cpu for command interpretation +- pcie_ep, io_cpu, 그리고 큐브당 UCIe PHY 포트들을 연결한다 +- 메모리 연산(MemoryWrite/Read)을 io_cpu를 거치지 않고 큐브 패브릭으로 + 직접 라우팅한다 +- 커널 런치 명령을 명령 해석을 위해 io_cpu를 통해 라우팅한다 ## Decision -### D1. IOChiplet internal NOC (io_noc) +### D1. IOChiplet 내부 NoC (io_noc) -Each IO chiplet instance contains an internal NOC node (`io_noc`) that connects: +각 IO chiplet 인스턴스는 다음을 연결하는 내부 NoC 노드(`io_noc`)를 +포함한다: -- `pcie_ep` — host-facing PCIe endpoint -- `io_cpu` — command processor for kernel launch interpretation -- `io_ucie-{PHY}.conn{N}` — per-PHY connection nodes to cube UCIe ports +- `pcie_ep` — 호스트 대면 PCIe 엔드포인트 +- `io_cpu` — 커널 런치 해석용 명령 프로세서 +- `io_ucie-{PHY}.conn{N}` — 큐브 UCIe 포트들로 가는 PHY별 연결 노드 -The io_noc is a forwarding-only fabric (`forwarding_v1` implementation) with -zero overhead. All routing decisions are made by the simulation engine based -on message type, not by io_noc itself. +io_noc은 오버헤드가 0인 포워딩 전용 패브릭(`forwarding_v1` 구현)이다. +모든 라우팅 결정은 io_noc 자체가 아니라 메시지 타입에 기반하여 시뮬레이션 +엔진이 내린다. -### D2. IOChiplet UCIe decomposition +### D2. IOChiplet UCIe 분해 -Each IO chiplet PHY port is decomposed into: +각 IO chiplet PHY 포트는 다음으로 분해된다: -- `io_ucie-{PHY}` — the UCIe protocol endpoint (overhead = 8ns) -- `io_ucie-{PHY}.conn{N}` — N connection nodes between io_noc and io_ucie +- `io_ucie-{PHY}` — UCIe 프로토콜 엔드포인트(overhead = 8ns) +- `io_ucie-{PHY}.conn{N}` — io_noc과 io_ucie 사이의 N개 연결 노드 -This mirrors the cube-side UCIe decomposition (ADR-0015 D1) and allows -multiple independent NOC-to-UCIe connections per PHY. +이는 큐브 측 UCIe 분해(ADR-0015 D1)를 미러링하며, PHY당 여러 독립적인 +NoC-UCIe 연결을 허용한다. -### D3. Memory R/W path (M_CPU bypass) +### D3. Memory R/W 경로 (M_CPU 우회) -Memory operations (MemoryWrite, MemoryRead) are routed directly from pcie_ep -through io_noc to the target cube, bypassing io_cpu entirely: +메모리 연산(MemoryWrite, MemoryRead)은 pcie_ep에서 io_noc을 거쳐 대상 +큐브로 직접 라우팅되며, io_cpu를 완전히 우회한다: ```text pcie_ep → io_noc → conn → io_ucie → [cube UCIe] → router mesh → hbm_ctrl ``` -This avoids the 10ns io_cpu overhead for pure data transfers. The simulation -engine's `_process_memory_direct()` method uses `find_memory_path()` which -resolves the shortest path from pcie_ep to the target HBM node. +이는 순수 데이터 전송에 대해 10ns의 io_cpu 오버헤드를 회피한다. +시뮬레이션 엔진의 `_process_memory_direct()` 메서드는 pcie_ep에서 대상 +HBM 노드까지의 최단 경로를 해석하는 `find_memory_path()`를 사용한다. -### D4. Kernel Launch path (via io_cpu) +### D4. 커널 런치 경로 (io_cpu 경유) -Kernel launch commands require io_cpu for command interpretation and PE -fan-out setup: +커널 런치 명령은 명령 해석 및 PE 팬아웃 설정을 위해 io_cpu를 필요로 +한다: ```text pcie_ep → io_noc → io_cpu → io_noc → conn → io_ucie → [cube UCIe] → noc → m_cpu → PE ``` -The engine's `_entry_points()` method routes KernelLaunchMsg through both -pcie_ep (entry) and io_cpu (command processing). +엔진의 `_entry_points()` 메서드는 KernelLaunchMsg를 pcie_ep(진입)와 +io_cpu(명령 처리) 양쪽 모두를 통해 라우팅한다. -### D5. IOChiplet-to-cube port mapping +### D5. IOChiplet-to-큐브 포트 매핑 -Each IO chiplet instance declares which cube ports it connects to: +각 IO chiplet 인스턴스는 자신이 연결되는 큐브 포트를 선언한다: ```yaml cube_ports: @@ -79,20 +80,20 @@ cube_ports: - { cube: {xy: [1,0]}, cube_side: N, phy: P1, distance_mm: 2.0 } ``` -The topology builder creates edges from io_ucie PHY nodes to the -corresponding cube UCIe port nodes, with the specified distance and -the IO chiplet's `per_connection_bw_gbs` as link bandwidth. +토폴로지 빌더는 io_ucie PHY 노드에서 해당 큐브 UCIe 포트 노드로의 엣지를 +지정된 거리 및 IO chiplet의 `per_connection_bw_gbs`를 링크 대역폭으로 +하여 생성한다. ## Consequences -- IO chiplet has a well-defined internal routing fabric -- Memory operations avoid unnecessary io_cpu overhead -- Kernel launch commands still get proper command interpretation -- The io_noc pattern is consistent with cube-level NOC design -- ADR-0003 D2 is extended (not contradicted) by this ADR +- IO chiplet은 잘 정의된 내부 라우팅 패브릭을 가진다 +- 메모리 연산은 불필요한 io_cpu 오버헤드를 회피한다 +- 커널 런치 명령은 여전히 적절한 명령 해석을 받는다 +- io_noc 패턴은 큐브 레벨 NoC 설계와 일관된다 +- ADR-0003 D2는 본 ADR에 의해 확장된다(모순되지 않는다) ## Links -- ADR-0003 D2 (IO chiplet definition) -- ADR-0015 D4 (fabric paths for Memory R/W and Kernel Launch) -- ADR-0012 D1 (host-to-IO_CPU message schema) +- ADR-0003 D2 (IO chiplet 정의) +- ADR-0015 D4 (Memory R/W 및 커널 런치의 패브릭 경로) +- ADR-0012 D1 (호스트-IO_CPU 메시지 스키마) diff --git a/docs/adr-ko/ADR-0017-dev-cube-noc-and-hbm-connectivity.md b/docs/adr-ko/ADR-0017-dev-cube-noc-and-hbm-connectivity.md index c442dde..1742d5e 100644 --- a/docs/adr-ko/ADR-0017-dev-cube-noc-and-hbm-connectivity.md +++ b/docs/adr-ko/ADR-0017-dev-cube-noc-and-hbm-connectivity.md @@ -1,4 +1,4 @@ -# ADR-0017: Cube NOC and HBM Connectivity +# ADR-0017: 큐브 NoC와 HBM 연결성 ## Status @@ -6,78 +6,74 @@ Accepted ## Context -The CUBE-level NOC is a 2D router mesh that carries every intra-cube -request: PE-to-HBM data, PE-to-PE traffic, command paths -(M_CPU↔PE_CPU), shared SRAM access, and inter-cube UCIe traffic. +CUBE 레벨의 NoC는 모든 큐브 내부 요청을 운반하는 2D 라우터 메시이다: +PE-HBM 데이터, PE-PE 트래픽, 명령 경로(M_CPU↔PE_CPU), 공유 SRAM 접근, +큐브 간 UCIe 트래픽. -The CUBE's HBM is exposed through per-PE controller endpoints attached -to PE routers. This per-PE partitioning makes local-vs-remote HBM -distinguishable by mesh distance: a PE's own HBM partition sits at its -own router (switching overhead only); another PE's HBM partition is -reachable by mesh hops to that PE's router. +CUBE의 HBM은 PE 라우터에 부착된 PE별 컨트롤러 엔드포인트를 통해 노출된다. +이러한 PE별 분할 덕분에 로컬-vs-원격 HBM이 메시 거리로 구분 가능하다: +PE 자신의 HBM 파티션은 자신의 라우터에 위치하고(스위칭 오버헤드만 발생), +다른 PE의 HBM 파티션은 해당 PE의 라우터로 메시 hop을 거쳐 도달 가능하다. -Two channel-mapping modes are supported in the design space: +설계 공간에서는 두 가지 채널 매핑 모드를 지원한다: -- **n:1 (default, implemented)** — each PE's HBM partition aggregates - `channels_per_pe` pseudo-channels into one endpoint. Effective - per-PE BW = N × per-channel BW. -- **1:1 (future)** — each PE router decomposes into per-channel - mini-routers; per-channel BW contention is modeled directly. +- **n:1 (default, 구현됨)** — 각 PE의 HBM 파티션이 `channels_per_pe` + pseudo-channel을 하나의 엔드포인트로 집계한다. 유효 PE당 BW = + N × per-channel BW. +- **1:1 (future)** — 각 PE 라우터가 채널별 미니 라우터로 분해된다; + 채널별 BW 경합을 직접 모델링한다. -In both modes the per-PE effective BW is identical; only the connectivity -granularity differs. +두 모드 모두 PE당 유효 BW는 동일하다; 연결 입도만 다르다. ## Decision -### D1. 2D router mesh +### D1. 2D 라우터 메시 -Each cube contains a 2D mesh of NOC routers generated by `mesh_gen.py`. +각 큐브는 `mesh_gen.py`가 생성하는 2D 라우터 메시를 포함한다. -- Node naming: `sip{S}.cube{C}.r{row}c{col}` (e.g., `sip0.cube0.r0c0`). -- Implementation: `forwarding_v1`. NOC `overhead_ns = 0`. -- Default 6×6 grid (sized from PE corner placement + UCIe attachment - count); larger PE counts scale the grid up. -- HBM exclusion zone: center rows/columns are excluded where HBM die - physically occupies space (e.g., r2c2, r2c3, r3c2, r3c3 for a 6×6). -- Latency = Manhattan distance × `ns_per_mm`. +- 노드 명명: `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 routing algorithm +### D2. XY 라우팅 알고리즘 -Deterministic XY routing: +결정론적 XY 라우팅: -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. +1. 수평 구간: 소스 X에서 목적지 X까지 소스 Y에서 라우팅. +2. 수직 구간: 소스 Y의 목적지 X에서 목적지 Y까지 라우팅. -Each directed segment carries a unique key: +각 유향 구간은 고유 키를 운반한다: -- Horizontal: `("H", y_band, x_min, x_max, direction)` -- Vertical: `("V", x_band, y_min, y_max, direction)` +- 수평: `("H", y_band, x_min, x_max, direction)` +- 수직: `("V", x_band, y_min, y_max, direction)` -Grid positions are snapped to the router grid, excluding the HBM zone. +그리드 위치는 HBM 영역을 제외하고 라우터 그리드에 스냅된다. -### D3. Per-segment contention model +### D3. 구간별 경합 모델 -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 — modelling link-level serialization in a wormhole-routed -mesh. +각 유향 XY 구간은 `simpy.Resource(capacity=1)`이다. 동일 구간을 공유하는 +트랜잭션(동일한 행 또는 열 밴드, 동일한 방향)은 자원을 두고 경합한다 — +wormhole 라우팅 메시에서의 링크 수준 직렬화를 모델링한다. -With no contention, NOC traversal latency equals Manhattan distance × -`ns_per_mm`. Under contention, SimPy's resource scheduling adds queueing -delay. +경합이 없을 때 NoC 순회 레이턴시는 Manhattan 거리 × `ns_per_mm`이다. +경합이 있을 때는 SimPy의 자원 스케줄링이 큐잉 지연을 추가한다. -### D4. NOC attachment points (per-PE HBM partition) +### D4. NoC 부착 지점 (PE별 HBM 파티션) -Every PE router carries three attachments: `pe{idx}.dma`, `pe{idx}.cpu`, -and `pe{idx}.hbm`. The last is the per-PE HBM controller endpoint — -`sip{S}.cube{C}.hbm_ctrl.pe{idx}` — which owns one slice of the cube's -HBM (one pseudo-channel group; see D8). +모든 PE 라우터는 세 개의 부착을 갖는다: `pe{idx}.dma`, `pe{idx}.cpu`, +그리고 `pe{idx}.hbm`. 마지막은 PE별 HBM 컨트롤러 엔드포인트로 +`sip{S}.cube{C}.hbm_ctrl.pe{idx}`이며, 큐브 HBM의 한 슬라이스를 +소유한다 (하나의 pseudo-channel 그룹; D8 참조). -Other attachments: +기타 부착: -- M_CPU and shared SRAM each occupy a dedicated edge router. -- UCIe endpoints (N/S/E/W) each expose 4 connection routers distributed - along that edge (see D6). +- M_CPU와 공유 SRAM은 각각 전용 edge 라우터를 점유한다. +- UCIe 엔드포인트(N/S/E/W)는 각각 해당 변에 분산된 4개의 연결 라우터를 + 노출한다 (D6 참조). ```text UCIe-N (conn x4) @@ -102,35 +98,34 @@ PE4.cpu <--+ +hbm.pe4| | +hbm.pe6+--< PE6.cpu UCIe-S (conn x4) ``` -Per-PE HBM partitioning is the key invariant that makes local vs -cross-PE HBM distinguishable by mesh distance (see D7). +PE별 HBM 분할은 로컬 vs 크로스-PE HBM을 메시 거리로 구분 가능하게 만드는 +핵심 불변식이다 (D7 참조). -### D5. NOC edge bandwidths and distances +### D5. NoC 엣지 대역폭과 거리 | Connection | BW (GB/s) | Distance | Notes | | ----------------------------- | ---------- | ------------- | ------------------------------------------- | -| PE_DMA → NOC | 256.0 | Physical (PE) | Matches local-HBM aggregate BW | -| NOC → PE_CPU | — | 0.0 mm | Command path only | -| Router ↔ hbm_ctrl.pe{idx} | 256.0 | 0.0 mm | Per PE router; N × per-channel BW (see D8) | -| NOC ↔ M_CPU | — | 0.0 mm | Command path | -| NOC ↔ SRAM | 128.0 × 4 | 0.0 mm | 512 GB/s aggregate | -| NOC ↔ UCIe conn | 128.0 | 0.0 mm | Per connection; 4 conn per port | +| 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` distances reflect the distributed nature of the NOC; actual -traversal distance is computed via Manhattan distance within the router -grid. +`0.0 mm` 거리는 NoC의 분산 특성을 반영한다; 실제 순회 거리는 라우터 +그리드 내에서 Manhattan 거리로 계산된다. -### D6. UCIe decomposition and inter-cube traffic +### D6. UCIe 분해와 큐브 간 트래픽 -Each of the 4 UCIe ports (N, S, E, W) decomposes into: +4개의 UCIe 포트(N, S, E, W) 각각은 다음으로 분해된다: -- 1 `ucie-{PORT}` node: UCIe protocol endpoint (`overhead = 8.0 ns`). -- 4 `ucie-{PORT}.conn{0-3}` nodes: connection bridges between NOC and UCIe. +- `ucie-{PORT}` 노드 1개: UCIe 프로토콜 엔드포인트 (`overhead = 8.0 ns`). +- `ucie-{PORT}.conn{0-3}` 노드 4개: NoC와 UCIe 간 연결 브리지. -This decomposition gives 4 independent NOC↔UCIe connections per port, -each with 128 GB/s bandwidth (512 GB/s aggregate per port). +이 분해로 포트당 4개의 독립 NoC↔UCIe 연결이 생성되며, 각각 128 GB/s +대역폭을 갖는다 (포트당 집계 512 GB/s). -Inter-cube traffic path: +큐브 간 트래픽 경로: ```text Source: PE_DMA → NOC → conn{i} → ucie-{PORT} @@ -138,56 +133,56 @@ Source: PE_DMA → NOC → conn{i} → ucie-{PORT} Target: ucie-{PORT} → conn{i} → r{x}c{y} → (mesh hops) → hbm_ctrl.pe{idx} ``` -UCIe overhead (8.0 ns) is applied at each `ucie-{PORT}` node, so a full -crossing incurs 16 ns (TX port + RX port). +UCIe 오버헤드(8.0 ns)는 각 `ucie-{PORT}` 노드에서 적용되므로 전체 횡단은 +16 ns(TX 포트 + RX 포트)가 소요된다. -### D7. Data paths through the NOC +### D7. NoC를 통한 데이터 경로 -All intra-cube traffic uses the same router mesh — no separate fast -paths. +모든 큐브 내부 트래픽은 동일한 라우터 메시를 사용한다 — 별도의 fast path는 +없다. -**Local HBM** (same PE's own partition; 0 mesh hops): +**로컬 HBM** (동일 PE의 자신 파티션; 0 메시 hop): ```text PE_DMA → r{x}c{y} → hbm_ctrl.pe{idx} (switching overhead only) ``` -**Cross-PE HBM within cube** (target PE's partition, reached by mesh): +**큐브 내 크로스-PE HBM** (대상 PE의 파티션, 메시로 도달): ```text PE_DMA → r{x}c{y} → (mesh hops) → r{x'}c{y'} → hbm_ctrl.pe{idx'} ``` -Example: PE0 (on `r0c0`) accessing PE2's HBM (PE2 on `r1c4`): +예시: PE0(`r0c0` 위)이 PE2의 HBM(PE2는 `r1c4` 위)에 접근: ```text PE0.pe_dma → r0c0 → r0c1 → r0c2 → r0c3 → r0c4 → r1c4 → hbm_ctrl.pe2 ``` -Dijkstra computes the shortest path within the mesh. +Dijkstra가 메시 내 최단 경로를 계산한다. -**Cross-cube HBM** (UCIe traversal): +**큐브 간 HBM** (UCIe 횡단): ```text PE_DMA → r{x}c{y} → conn → ucie-{PORT} → [seam] → ucie-{PORT'} → conn → r{x'}c{y'} → hbm_ctrl.pe{idx'} ``` -**Kernel launch command to PE**: +**PE로의 커널 launch 명령**: ```text [from io_noc] → ucie → conn → r{x}c{y} → (mesh) → M_CPU → (mesh) → PE_CPU ``` -**Shared SRAM access**: +**공유 SRAM 접근**: ```text PE_DMA → r{x}c{y} → (mesh) → SRAM ``` -### D8. HBM channel mapping mode +### D8. HBM 채널 매핑 모드 -Channel mapping is configured at cube scope: +채널 매핑은 큐브 범위에서 구성된다: ```yaml cube: @@ -200,37 +195,35 @@ cube: hbm_total_gb_per_cube: 48 ``` -**n:1 mode (default, implemented).** Each PE's HBM partition is a single -endpoint `hbm_ctrl.pe{idx}` that aggregates `channels_per_pe` pseudo- -channels. The `Router ↔ hbm_ctrl.pe{idx}` link bandwidth equals -`channels_per_pe × hbm_channel_bw_gbs`. Pseudo-channels are assumed to -interleave; only aggregate per-PE BW is modeled. No separate aggregated -router node exists — the per-PE router itself serves that role. +**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 mode (future).** Each PE router decomposes into N channel -mini-routers; per-channel routing carries fully-resolved PA + channel ID. -A `ChannelSplitter` resolves a logical access to N per-channel physical -requests. Per-channel link models BW contention. Cross-PE channel -access semantics are deferred to the implementation ADR. +**1:1 모드 (future).** 각 PE 라우터가 N개의 채널 미니 라우터로 +분해된다; 채널별 라우팅이 완전히 해석된 PA + channel ID를 운반한다. +`ChannelSplitter`가 논리적 접근을 N개의 채널별 물리 요청으로 해결한다. +채널별 링크가 BW 경합을 모델링한다. 크로스-PE 채널 접근 시맨틱은 +구현 ADR로 연기된다. -**BW math (defaults).** +**BW 계산 (default 값).** | Parameter | Value | | ---------------------------------- | -------------------------- | -| pseudo channels per cube | 64 (parameter) | -| PEs per cube | 8 (parameter) | -| channels per PE (N) | 64 / 8 = 8 | -| per-channel BW | 32 GB/s (parameter) | -| per-PE local BW | N × 32 = 256 GB/s | -| cube total HBM BW | 64 × 32 = 2048 GB/s | +| 큐브당 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 | -Both modes give the same per-PE effective BW; only the request shape and -contention model differ. +두 모드 모두 PE당 유효 BW는 동일하다; 요청 형태와 경합 모델만 다르다. -### D9. AddressResolver — per-PE HBM endpoint +### D9. AddressResolver — PE별 HBM 엔드포인트 -The address resolver decodes a PA's HBM offset to the owning PE's -partition: +주소 리졸버는 PA의 HBM 오프셋을 소유 PE의 파티션으로 디코딩한다: ```python # policy/routing/router.py @@ -241,51 +234,49 @@ if addr.kind == "hbm": return f"sip{s}.cube{d}.hbm_ctrl.pe{pe_id}" ``` -The pe_id computation is intrinsic to the routing layer (not a -topology-time concern). Any HBM PA falls within exactly one partition, -yielding deterministic routing. +pe_id 계산은 라우팅 레이어의 본질적 일부이다 (토폴로지 시점 관심사가 +아니다). 모든 HBM PA는 정확히 하나의 파티션에 속하므로 결정론적 라우팅이 +보장된다. -External callers (e.g., M_CPU DMA, Memory R/W from PCIE_EP) follow the -same resolver path — there is no separate fast path. +외부 호출자(예: M_CPU DMA, PCIE_EP로부터의 Memory R/W)도 동일한 리졸버 +경로를 따른다 — 별도의 fast path는 존재하지 않는다. -### D10. Mesh generation parameters +### D10. 메시 생성 파라미터 -`mesh_gen.py` produces `cube_mesh.yaml` from: +`mesh_gen.py`는 다음으로부터 `cube_mesh.yaml`을 생성한다: -- `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. +- `cube.pe_layout`: 코너 배치(NW, NE, SW, SE)와 코너당 PE 개수. +- `cube.geometry`: 큐브 물리 치수와 HBM 영역. +- `cube.ucie.n_connections`: UCIe 부착용 라우터 개수를 결정. -Output `mesh_data` dictionary contains: +출력 `mesh_data` 딕셔너리는 다음을 포함한다: -- Router grid with positions and HBM exclusion zones. -- PE-to-router attachments (`pe{idx}.dma`, `pe{idx}.cpu`, `pe{idx}.hbm` - per PE). -- UCIe-to-router attachments (N/S/E/W distributed across edge routers). -- M_CPU and SRAM router attachments. +- 위치 및 HBM 제외 영역을 갖는 라우터 그리드. +- PE-라우터 부착 (PE별 `pe{idx}.dma`, `pe{idx}.cpu`, `pe{idx}.hbm`). +- UCIe-라우터 부착 (N/S/E/W가 edge 라우터에 분산). +- M_CPU와 SRAM 라우터 부착. ## Consequences -- Local HBM (0 mesh hops, switching overhead only) and cross-PE HBM - (mesh hops) are naturally distinguishable, satisfying SPEC R5 - (multi-domain communication) and ADR-0002 (no zero-latency end-to-end - paths). -- All cube-internal traffic routes through one mesh — single contention - model, single layout, single set of edge BWs. -- Per-PE HBM partitioning maps cleanly to the LA model (ADR-0011): each - PE's partition is the n:1 aggregate of its assigned pseudo-channels. -- 1:1 mode extension is structurally natural — split each PE router into - N channel routers. -- Mesh generation is fully parameterised by `topology.yaml`; PE/cube - geometry changes propagate without code edits. +- 로컬 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 (Routing distance, ordering, no zero-latency paths) -- ADR-0003 D3 (cube-level NOC definition — extended here) -- ADR-0004 (Memory semantics, local HBM) -- ADR-0011 (Memory addressing — LA model consumes per-PE partition) -- ADR-0014 D1 (PE_DMA egress via router mesh) -- ADR-0015 D4 (fabric paths for Memory R/W and Kernel Launch) -- ADR-0016 (IOChiplet io_noc — analogous pattern at IO chiplet level) -- ADR-0033 (Latency model: per-PC parallelism, switch penalty) +- 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당 병렬성, 스위치 패널티) diff --git a/docs/adr-ko/ADR-0022-prog-program-id-2d-grid.md b/docs/adr-ko/ADR-0022-prog-program-id-2d-grid.md index 371bb49..94177c8 100644 --- a/docs/adr-ko/ADR-0022-prog-program-id-2d-grid.md +++ b/docs/adr-ko/ADR-0022-prog-program-id-2d-grid.md @@ -1,4 +1,4 @@ -# ADR-0022: 2D Grid program_id Semantics +# ADR-0022: 2D 그리드 program_id 시맨틱 ## Status @@ -6,38 +6,43 @@ Accepted ## Context -Triton kernels use `tl.program_id(axis)` to identify their position in a launch grid. -Our hardware has a 2-level hierarchy: **cubes** contain **PEs**. -The previous implementation ignored the `axis` parameter and always returned a flat PE index, -making it impossible for kernels to distinguish their cube-local position from their cube identity. +Triton 커널은 `tl.program_id(axis)`를 사용해 launch 그리드 내 자신의 +위치를 식별한다. 본 하드웨어는 2단계 계층을 갖는다: **큐브**가 **PE**를 +포함한다. 이전 구현은 `axis` 파라미터를 무시하고 항상 평탄화된 PE +인덱스를 반환했기 때문에, 커널이 큐브 내부 위치와 큐브 식별자를 구분할 +수 없었다. ## Decision -Map `tl.program_id` and `tl.num_programs` to the 2D hardware grid: +`tl.program_id`와 `tl.num_programs`를 2D 하드웨어 그리드에 매핑한다: | Call | Returns | Description | |------|---------|-------------| -| `tl.program_id(axis=0)` | `local_pe_id` | PE index within cube | -| `tl.program_id(axis=1)` | `cube_id` | Cube index | -| `tl.num_programs(axis=0)` | `num_pes_per_cube` | PEs per cube | -| `tl.num_programs(axis=1)` | `num_cubes` | Total cubes | +| `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` | 전체 큐브 개수 | -Global PID is derived as: +전역 PID는 다음과 같이 도출된다: ```python global_pid = tl.program_id(axis=1) * tl.num_programs(axis=0) + tl.program_id(axis=0) ``` -### Axis mapping rationale +### 축 매핑 근거 -- **axis=0 = PE (innermost)**: PEs within a cube share HBM and communicate via local NOC mesh. This is the fast, tightly-coupled dimension — analogous to threads within a block. -- **axis=1 = Cube (outer)**: Cross-cube communication goes through UCIe with higher latency. This is the coarser scheduling dimension — analogous to blocks in a grid. +- **axis=0 = PE (최내부)**: 큐브 내부 PE들은 HBM을 공유하고 로컬 NoC + 메시를 통해 통신한다. 빠르고 강하게 결합된 차원이다 — 블록 내부의 + 스레드와 유사하다. +- **axis=1 = 큐브 (외부)**: 큐브 간 통신은 더 높은 레이턴시의 UCIe를 + 통한다. 더 거친 스케줄링 차원이다 — 그리드 내의 블록과 유사하다. ## Implementation ### TLContext (`triton_emu/tl_context.py`) -Added `cube_id` and `num_cubes` constructor parameters. `program_id()` and `num_programs()` dispatch on `axis`: +`cube_id`와 `num_cubes` 생성자 파라미터를 추가했다. `program_id()`와 +`num_programs()`가 `axis`로 디스패치한다: ```python def program_id(self, axis: int = 0) -> int: @@ -53,18 +58,22 @@ def num_programs(self, axis: int = 0) -> int: ### PE_CPU (`components/builtin/pe_cpu.py`) -- Extracts `num_cubes` from `ctx.spec["system"]["sips"]["cubes_per_sip"]` -- Passes `cube_id` (already available as `self._cube_idx`) and `num_cubes` to TLContext +- `ctx.spec["system"]["sips"]["cubes_per_sip"]`에서 `num_cubes`를 + 추출한다. +- `cube_id`(이미 `self._cube_idx`로 사용 가능)와 `num_cubes`를 + TLContext에 전달한다. ### KernelRunner (`triton_emu/kernel_runner.py`) -- Receives `num_cubes` from PE_CPU -- Passes `cube_id` and `num_cubes` to TLContext in greenlet mode +- PE_CPU로부터 `num_cubes`를 수신한다. +- greenlet 모드에서 `cube_id`와 `num_cubes`를 TLContext에 전달한다. ## Backward Compatibility -- Existing code using `tl.program_id(0)` or `tl.program_id()` is unchanged — returns the same PE index as before. -- `cube_id` and `num_cubes` default to `0` and `1`, so callers that don't provide them (e.g. unit tests) continue to work. +- `tl.program_id(0)` 또는 `tl.program_id()`를 사용하는 기존 코드는 + 변경되지 않는다 — 이전과 동일한 PE 인덱스를 반환한다. +- `cube_id`와 `num_cubes`는 기본값이 `0`과 `1`이므로, 이를 제공하지 + 않는 호출자(예: 유닛 테스트)도 계속 동작한다. ## Usage Example @@ -74,7 +83,7 @@ def sharded_gemm_kernel(a_ptr, b_ptr, out_ptr, M, K, N, tl): cube_id = tl.program_id(axis=1) # which cube global_pid = cube_id * tl.num_programs(axis=0) + local_pid - # Column-wise sharding across global 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 @@ -86,5 +95,6 @@ def sharded_gemm_kernel(a_ptr, b_ptr, out_ptr, M, K, N, tl): ## Consequences -- Benchmarks can now express cube-aware sharding and addressing without hardcoding topology dimensions. -- Future axis=2 (SIP-level) can be added following the same pattern if needed. +- 벤치마크가 토폴로지 차원을 하드코딩하지 않고 큐브 인식 샤딩과 주소 + 지정을 표현할 수 있다. +- 필요 시 axis=2(SIP 레벨)를 동일한 패턴을 따라 향후 추가할 수 있다. diff --git a/docs/adr-ko/ADR-0032-algo-intercube-allreduce.md b/docs/adr-ko/ADR-0032-algo-intercube-allreduce.md index bb6ba3c..f89dd77 100644 --- a/docs/adr-ko/ADR-0032-algo-intercube-allreduce.md +++ b/docs/adr-ko/ADR-0032-algo-intercube-allreduce.md @@ -1,4 +1,4 @@ -# ADR-0032: Intercube All-Reduce — pe0 cube-mesh reduce + multi-SIP exchange +# ADR-0032: 큐브 간 All-Reduce — pe0 큐브-메시 리듀스 + 다중-SIP 교환 ## Status @@ -6,112 +6,109 @@ Accepted (supersedes ADR-0029). ## Context -### Goal +### 목표 -Define a single all-reduce algorithm that exploits the topology hierarchy: -cube mesh within each SIP (intercube) + inter-SIP exchange. One kernel, -one SFR configuration path, driven by `topology.yaml` and `ccl.yaml`. +토폴로지 계층을 활용하는 단일 all-reduce 알고리즘을 정의한다: 각 SIP +내부의 큐브 메시(큐브 간) + SIP 간 교환. 단일 커널, 단일 SFR 구성 +경로이며 `topology.yaml`과 `ccl.yaml`로 구동된다. -### Why replace ADR-0029 (hierarchical 3-level) +### ADR-0029(계층적 3-레벨)를 대체하는 이유 -ADR-0029 proposed a 3-level (intra-cube → inter-cube → inter-SIP) algorithm -where every PE in the system participates. In practice this adds the -intra-cube PE-to-PE stage complexity (bidirectional reduce + chain broadcast) -without matching the common workload pattern where the tensor is sharded -**per cube** (not per PE within a cube). +ADR-0029는 시스템의 모든 PE가 참여하는 3-레벨(큐브 내 → 큐브 간 → +SIP 간) 알고리즘을 제안했다. 실제로는 텐서가 큐브 내 PE 단위가 아니라 +**큐브 단위로 샤딩되는** 일반적 워크로드 패턴과 맞지 않으면서, 큐브 내 +PE-PE stage 복잡성(양방향 reduce + 체인 브로드캐스트)을 추가한다. -Moreover, the hierarchical design required: -- per-PE neighbor graph installation (`_build_pe_installs` multi-level) -- multi-level topology schema (`hierarchical_3level`) -- `all_pes` mapper + `multi_pe_sip_local` validator infrastructure +또한 계층적 설계는 다음을 요구했다: +- PE별 이웃 그래프 설치 (`_build_pe_installs` 다중 레벨) +- 다중 레벨 토폴로지 스키마 (`hierarchical_3level`) +- `all_pes` 매퍼 + `multi_pe_sip_local` 검증자 인프라 -The intercube algorithm below removes all of that: **pe0-only same-lane -intercube reduce on the 4×4 cube mesh**, then inter-SIP exchange on the -root cube, then broadcast back. Simpler kernel, simpler wiring, same -bandwidth characteristics for the common per-cube DP workload. +아래의 큐브 간 알고리즘은 이 모든 것을 제거한다: **4×4 큐브 메시 위에서 +pe0만의 same-lane 큐브 간 reduce**, 그 다음 루트 큐브에서 SIP 간 교환, +그 다음 다시 브로드캐스트. 더 단순한 커널, 더 단순한 와이어링, +일반적인 큐브당 DP 워크로드에 대해 동일한 대역폭 특성을 갖는다. -### Current state +### 현재 상태 -- `src/kernbench/ccl/algorithms/intercube_allreduce.py` — kernel +- `src/kernbench/ccl/algorithms/intercube_allreduce.py` — 커널 - `src/kernbench/ccl/sfr_config.py` — `configure_sfr_intercube_multisip` -- `src/kernbench/runtime_api/distributed.py` — `AhbmCCLBackend` wires this - automatically at `init_process_group` time. -- Old `ring_allreduce`, `mesh_allreduce`, `tree_allreduce`, - `hierarchical_allreduce` modules and their tests are **removed**. +- `src/kernbench/runtime_api/distributed.py` — `AhbmCCLBackend`가 + `init_process_group` 시점에 자동으로 와이어링한다. +- 기존 `ring_allreduce`, `mesh_allreduce`, `tree_allreduce`, + `hierarchical_allreduce` 모듈과 그 테스트는 **제거됨**. --- ## Decision -### D1. Algorithm structure — 5 phases +### D1. 알고리즘 구조 — 5단계 -For each SIP (launched concurrently by `mp.spawn`): +각 SIP에 대해 (`mp.spawn`으로 동시에 launch): ``` -Phase 1 — Row reduce W → E (cube mesh, pe0 only): - col=0 sends E → col=1 accumulates, sends E → ... → col=3 holds row sum. +Phase 1 — Row reduce W → E (큐브 메시, pe0만): + col=0이 E로 송신 → col=1이 누적, E로 송신 → ... → col=3이 row sum 보유. -Phase 2 — Col reduce N → S on rightmost column (pe0, col = mesh_w-1): - row=0 sends S → row=1 accumulates, sends S → ... → root cube (15) - holds the full SIP sum. +Phase 2 — 최우측 열에서 Col reduce N → S (pe0, col = mesh_w-1): + row=0이 S로 송신 → row=1이 누적, S로 송신 → ... → 루트 큐브 (15)가 + 전체 SIP sum 보유. -Phase 3 — Inter-SIP exchange on root cube (pe0 of root cube only): +Phase 3 — 루트 큐브에서 SIP 간 교환 (루트 큐브의 pe0만): Ring / torus-2d row+col ring / mesh-2d chain reduce+broadcast — - selected by sip_topo_kind (from topology.yaml sips.topology). + sip_topo_kind(topology.yaml의 sips.topology)로 선택. -Phase 4 — Col broadcast S → N on rightmost column. +Phase 4 — 최우측 열에서 Col 브로드캐스트 S → N. -Phase 5 — Row broadcast E → W across the cube mesh. +Phase 5 — 큐브 메시 전반에 걸친 Row 브로드캐스트 E → W. ``` -After all phases every cube's pe0 holds the global sum. +모든 단계가 끝나면 모든 큐브의 pe0이 전역 sum을 보유한다. -The kernel is a single function parameterised by `sip_topo_kind ∈ {0, 1, 2}` -(ring_1d, torus_2d, mesh_2d_no_wrap). Phases 1-2 and 4-5 are identical -across topologies; only phase 3 branches. Helper functions -`_inter_sip_ring`, `_inter_sip_torus_2d`, `_inter_sip_mesh_2d` encode the -three exchange patterns. +커널은 `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. Tensor layout (rank = SIP, per-worker) +### D2. 텐서 레이아웃 (rank = SIP, 워커별) -Per ADR-0024 rank = SIP at the process-group level. Each worker allocates -its own cube-mesh-spanning tensor: +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) ``` -Shard layout: 16 shards per SIP, one per cube on pe0. The kernel addresses -each cube's shard as `pe_addr = t_ptr + cube_id * n_elem * 2`. +샤드 레이아웃: SIP당 16개 샤드, 큐브별 pe0에 하나씩. 커널은 각 큐브의 +샤드를 `pe_addr = t_ptr + cube_id * n_elem * 2`로 주소 지정한다. -### D3. SFR / IPCQ wiring — `configure_sfr_intercube_multisip` +### D3. SFR / IPCQ 와이어링 — `configure_sfr_intercube_multisip` -Replaces the rank-to-2-PE install from ADR-0024. Wires PE_IPCQ neighbor -tables for **every cube's pe0 across every SIP** — regardless of which -cube is the root or which SIP topology is selected. This lets the kernel -elect the root cube at runtime and supports topology switches without -re-wiring. +ADR-0024의 rank-to-2-PE 설치를 대체한다. 어느 큐브가 루트인지 또는 어느 +SIP 토폴로지가 선택되었는지와 무관하게 **모든 SIP의 모든 큐브의 pe0**에 +대해 PE_IPCQ 이웃 테이블을 와이어링한다. 이를 통해 커널이 런타임에 루트 +큐브를 선출할 수 있고, 재와이어링 없이 토폴로지 전환을 지원한다. | Level | Direction labels | Scope | |---|---|---| -| Intercube within SIP | N / S / E / W | pe0 of every cube → pe0 of mesh neighbors (no wrap) | -| Inter-SIP (all cubes) | global_E / global_W / global_N / global_S | pe0 of cube c on sip A → pe0 of cube c on peer SIP per `sips.topology` | +| 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 | -Inter-SIP directions use the `global_*` prefix to keep the namespace -disjoint from intercube directions. ADR-0025's `_OPPOSITE_DIR` is extended -with `global_E ↔ global_W` and `global_N ↔ global_S` so the reverse- -direction resolver handles 2-SIP bidirectional rings correctly. +SIP 간 방향은 `global_*` 접두사를 사용하여 큐브 간 방향과 네임스페이스를 +분리한다. ADR-0025의 `_OPPOSITE_DIR`은 `global_E ↔ global_W` 및 +`global_N ↔ global_S`로 확장되어, 2-SIP 양방향 ring에 대한 역방향 +리졸버가 올바르게 처리되도록 한다. -Internally the function calls `install_ipcq` with: +내부적으로 이 함수는 다음 인자로 `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)]` -- A closure-captured `neighbors()` function that builds the map above. +- 위 매핑을 생성하는 클로저로 캡처된 `neighbors()` 함수. -This `world_size` is internal to IPCQ wiring and does not leak to the -process-group rank. +이 `world_size`는 IPCQ 와이어링 내부적이며 프로세스-그룹 rank로 유출되지 +않는다. -### D4. SIP topology — from `topology.yaml` +### D4. SIP 토폴로지 — `topology.yaml`에서 ```yaml system: @@ -120,35 +117,36 @@ system: topology: ring_1d # or torus_2d, mesh_2d_no_wrap ``` -- `ring_1d`: n_sips-1 rounds of `send global_E / recv global_W`. -- `torus_2d`: sqrt(n_sips)×sqrt(n_sips) wrapping mesh. Row ring on - `global_E/W` then col ring on `global_S/N`. -- `mesh_2d_no_wrap`: square mesh without wrap-around. Chain reduce + - broadcast per dimension. +- `ring_1d`: n_sips-1 라운드의 `send global_E / recv global_W`. +- `torus_2d`: sqrt(n_sips)×sqrt(n_sips) 랩핑 메시. `global_E/W`에서 + row ring, 이어서 `global_S/N`에서 col ring. +- `mesh_2d_no_wrap`: 랩어라운드 없는 정사각형 메시. 차원별 chain + reduce + 브로드캐스트. -2D variants require `n_sips` to be a perfect square. +2D 변형은 `n_sips`가 완전 제곱수여야 한다. -### D5. Process-group integration — `AhbmCCLBackend` +### D5. 프로세스-그룹 통합 — `AhbmCCLBackend` -At `init_process_group` time the backend: +`init_process_group` 시점에 백엔드는: -1. Loads `ccl.yaml` + `topology.yaml`. -2. Derives `sip_topo_kind, sip_topo_w, sip_topo_h` from - `system.sips.topology` using the algorithm module's `TOPO_NAME_TO_KIND`. -3. Calls `configure_sfr_intercube_multisip(engine, spec, cfg)` — one-time - SFR wiring, mirrors NCCL communicator creation. +1. `ccl.yaml` + `topology.yaml`을 로드한다. +2. 알고리즘 모듈의 `TOPO_NAME_TO_KIND`를 사용하여 + `system.sips.topology`로부터 `sip_topo_kind, sip_topo_w, sip_topo_h`를 + 도출한다. +3. `configure_sfr_intercube_multisip(engine, spec, cfg)`를 호출한다 — + 일회성 SFR 와이어링, NCCL 커뮤니케이터 생성을 모방한다. -At each `dist.all_reduce(tensor)` call: +각 `dist.all_reduce(tensor)` 호출 시: -1. Resolves `kernel_fn` from `cfg["module"]`. -2. Builds args: `(n_elem, cube_w, cube_h, n_sips)` from - `kernel_args(world_size, n_elem)`. -3. Appends `(sip_rank, sip_topo_kind, sip_topo_w, sip_topo_h)` where - `sip_rank` is the current greenlet's bound rank. -4. Launches with `_defer_wait=True`; the main scheduler drains pending - handles after all workers submit (per ADR-0027 D0.4). +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. Config schema +### D6. 구성 스키마 `ccl.yaml`: @@ -178,38 +176,39 @@ sip: cube_mesh: { w: 4, h: 4 } ``` -### D7. Algorithm module contract +### D7. 알고리즘 모듈 계약 -Modules loaded via `cfg["module"]` must export: +`cfg["module"]`로 로드되는 모듈은 다음을 export해야 한다: | Name | Purpose | |---|---| -| `kernel` | callable, signature `(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` | returns the first 4 scalar args (per-tensor) | -| `TOPO_NAME_TO_KIND: dict[str, int]` | maps `system.sips.topology` name to kernel branch code | -| `SIP_TOPO_RING`, `SIP_TOPO_TORUS`, `SIP_TOPO_MESH` | integer constants (0, 1, 2) | +| `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 protocol (neighbor table, send/recv, credit return). -- **ADR-0024**: rank = SIP launcher, `mp.spawn`, greenlet-local rank. -- **ADR-0025**: Address-based IPCQ direction matching; extended - `_OPPOSITE_DIR` with `global_*` pairs. -- **ADR-0027**: Worker-wait / collective-pending drain in main scheduler. +- **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 -- **Per-PE allreduce** (intra-cube PE-to-PE reduce). Out of scope — the - workload for this algorithm is per-cube DP. -- **Asymmetric SIP topologies** (non-square mesh/torus). `torus_2d` and - `mesh_2d_no_wrap` require `n_sips = k²`. -- **Pipelined chunks**: single-tile per cube, no pipelining yet. -- **Root cube runtime election**: the kernel currently uses - `root_cube = (mesh_h - 1) * mesh_w + (mesh_w - 1)` hardcoded to the SE - corner. SFR wiring covers all cubes, so runtime election is a pure kernel - change when needed. +- **PE별 allreduce** (큐브 내 PE-PE reduce). 범위 밖 — 본 알고리즘의 + 워크로드는 큐브당 DP이다. +- **비대칭 SIP 토폴로지** (정사각형이 아닌 메시/토러스). + `torus_2d`와 `mesh_2d_no_wrap`은 `n_sips = k²`를 요구한다. +- **파이프라인 청크**: 큐브당 단일 타일, 아직 파이프라이닝 없음. +- **루트 큐브의 런타임 선출**: 커널은 현재 SE 코너로 하드코딩된 + `root_cube = (mesh_h - 1) * mesh_w + (mesh_w - 1)`을 사용한다. SFR + 와이어링이 모든 큐브를 커버하므로, 필요해질 때 런타임 선출은 순수 + 커널 변경이다. --- @@ -217,24 +216,24 @@ Modules loaded via `cfg["module"]` must export: ### Positive -- **Single kernel, single install path** for all-reduce — replaces four - removed modules (`ring`, `mesh`, `tree`, `hierarchical`). -- **Topology-agnostic kernel**: ring / torus / mesh selected via one - integer param, no kernel duplication. -- **Automatic via `dist.all_reduce`**: no bench-level or user-level - algorithm selection needed; config-driven end-to-end. -- **Full SFR wiring**: every cube on every SIP has inter-SIP links - available — supports future dynamic root-cube election. +- **단일 커널, 단일 설치 경로**로 all-reduce를 처리 — 제거된 네 개의 + 모듈(`ring`, `mesh`, `tree`, `hierarchical`)을 대체한다. +- **토폴로지 무관 커널**: ring / torus / mesh를 정수 파라미터 하나로 + 선택, 커널 중복 없음. +- **`dist.all_reduce`를 통한 자동화**: 벤치 레벨이나 사용자 레벨의 + 알고리즘 선택 불필요; end-to-end 구성 기반. +- **완전한 SFR 와이어링**: 모든 SIP의 모든 큐브가 SIP 간 링크를 보유 — + 향후 동적 루트 큐브 선출을 지원한다. ### Negative -- **Not suitable for per-PE sharded tensors**: TP-layer-style tensors that - shard within one cube across 8 PEs are not addressable by this kernel. - Such workloads would need a separate intra-cube all-reduce path (not - yet implemented). -- **`configure_sfr_intercube_multisip` always wires all pe0s**: even if a - given run only needs a subset (e.g. 1 SIP, ring only). Install cost is - small but not zero. +- **PE별 샤딩된 텐서에 부적합**: 큐브 하나 내부에서 8개 PE에 걸쳐 + 샤딩되는 TP-레이어 스타일 텐서는 본 커널로 주소 지정할 수 없다. 이러한 + 워크로드에는 별도의 큐브 내 all-reduce 경로가 필요하다 (아직 구현되지 + 않음). +- **`configure_sfr_intercube_multisip`는 항상 모든 pe0을 와이어링**: + 주어진 실행이 부분집합(예: 1 SIP, ring만)만 필요하더라도. 설치 비용은 + 작지만 영(zero)은 아니다. --- @@ -242,15 +241,15 @@ Modules loaded via `cfg["module"]` must export: | File | Change | |---|---| -| `src/kernbench/ccl/algorithms/intercube_allreduce.py` (new) | Kernel + `_inter_sip_*` helpers + `TOPO_NAME_TO_KIND` | -| `src/kernbench/ccl/sfr_config.py` (new) | `configure_sfr_intercube_multisip` | -| `src/kernbench/ccl/topologies.py` | Added `torus_2d`, `mesh_2d_no_wrap` | -| `src/kernbench/ccl/install.py` | Extended `_OPPOSITE_DIR` with `global_*` pairs | -| `src/kernbench/runtime_api/distributed.py` | `AhbmCCLBackend` uses `configure_sfr_intercube_multisip` + appends sip_rank/topo args | -| `ccl.yaml` | Single `intercube_allreduce` entry | -| `topology.yaml` | Added `system.sips.topology` | -| `benches/ccl_allreduce.py` | Row-wise cube-mesh tensor layout | -| `tests/test_allreduce_multidevice.py` (new) | Config-driven ring/torus/mesh | -| `tests/test_distributed_intercube_allreduce.py` (new) | Full `dist.all_reduce` path | -| `tests/test_intercube_sfr_config.py` (new) | SFR wiring verification | -| Removed | `ring_allreduce.py`, `mesh_allreduce.py`, `tree_allreduce.py`, `hierarchical_allreduce.py`, `hello_send.py`, `testing.py` and their tests | +| `src/kernbench/ccl/algorithms/intercube_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` | 단일 `intercube_allreduce` 항목 | +| `topology.yaml` | `system.sips.topology` 추가 | +| `benches/ccl_allreduce.py` | Row-wise 큐브-메시 텐서 레이아웃 | +| `tests/test_allreduce_multidevice.py` (신규) | 구성 기반 ring/torus/mesh | +| `tests/test_distributed_intercube_allreduce.py` (신규) | 전체 `dist.all_reduce` 경로 | +| `tests/test_intercube_sfr_config.py` (신규) | SFR 와이어링 검증 | +| 제거 | `ring_allreduce.py`, `mesh_allreduce.py`, `tree_allreduce.py`, `hierarchical_allreduce.py`, `hello_send.py`, `testing.py` 및 그 테스트 | diff --git a/docs/adr-ko/ADR-0033-lat-latency-model-assumptions.md b/docs/adr-ko/ADR-0033-lat-latency-model-assumptions.md index 13ca1f9..97c7ba3 100644 --- a/docs/adr-ko/ADR-0033-lat-latency-model-assumptions.md +++ b/docs/adr-ko/ADR-0033-lat-latency-model-assumptions.md @@ -1,4 +1,4 @@ -# ADR-0033 — Latency Model: Assumptions and Known Simplifications +# ADR-0033 — 레이턴시 모델: 가정 및 알려진 단순화 ## Status @@ -6,157 +6,147 @@ Accepted ## Context -The simulator is an analytical, event-driven performance model — not a -cycle-accurate or RTL-level simulator. Many real-HW effects are approximated -or omitted by design. To keep the model auditable and reviewable as a whole, -this ADR consolidates the assumptions in one place. Individual component ADRs -(ADR-0015, ADR-0017, ADR-0004) define the *mechanisms*; this document defines -the *limits of fidelity*. +이 시뮬레이터는 분석적·이벤트 기반 성능 모델이지, 사이클 정확(cycle-accurate) +시뮬레이터나 RTL 수준 시뮬레이터가 아니다. 실제 HW의 많은 효과들이 설계상 +근사되거나 생략되었다. 모델 전체를 감사·리뷰할 수 있도록 유지하기 위해, +본 ADR은 그런 가정들을 한 곳에 통합한다. 개별 컴포넌트 ADR(ADR-0015, +ADR-0017, ADR-0004)들이 *메커니즘*을 정의하고, 본 문서는 *충실도의 한계*를 +정의한다. ## Decisions -### D1. Modeled precisely +### D1. 정밀하게 모델링되는 것 -- **Per-directed-edge BW occupancy** (FIFO serialization via `available_at`) — +- **방향 에지별 BW 점유** (`available_at`을 통한 FIFO 직렬화) — ADR-0015 D2. -- **Per-component switching/overhead latency** (`overhead_ns` attr). -- **HBM per-pseudo-channel parallelism** via stateless `pc_avail[N]` array - with address-based PC selection (ADR-0034 D3). Burst granularity tunable - (`burst_bytes`, default 256B). Read and write share each PC's - `available_at` (real HW command bus is per-PC shared). -- **HBM direction switching penalty mechanism**: per-PC last-direction - tracking + configurable `switch_penalty_ns`. Default 0 — see D2. -- **Wire chunk-streaming (Phase 2c)**: each wire decomposes Transactions - with payload into `Flit` objects of `flit_bytes` (default = HBM - `burst_bytes` = 256B). The wire emits each flit individually after - `prop_ns + flit_nbytes/bw_gbs` so the link's bandwidth throttles - flit arrival rate per real-HW wormhole semantics. -- **Separate Stores per directed edge** (Phase 2c key fix): the wire - is the *only* conduit between `src.out_ports[dst]` and - `dst.in_ports[src]`. Earlier the two were aliased to the same - `simpy.Store`; when the wire put a chunkified flit back, the - destination's `fan_in` could pull it before the wire applied - bandwidth delay, leaving half the flits bypassing the bottleneck. -- **Flit-aware pass-through** (`TransitComponent`, `HbmCtrlComponent`): - forward each flit serially with per-transaction overhead applied - ONCE on the first-flit arrival (header decode model). Subsequent - flits pipeline through with no extra delay. Wormhole emerges - naturally across multi-hop paths. -- **HBM CTRL per-flit PC commit**: each flit arriving at HBM CTRL - schedules a PC commit at `max(env.now, pc_avail[pc]) + chunk_time`, - with the `is_last` flit waiting for the last PC commit before - signaling `txn.done`. -- **Non-flit-aware components (default) reassemble flits at - ``_fan_in``** before the legacy `_forward_txn` path runs. This - preserves backward compatibility for components that have not yet - been migrated to flit-aware processing (e.g., `MCpuComponent`, - `IoCpuComponent` sub-txn generators). Such components reassemble - *once per leg boundary*, NOT per hop — multi-hop wormhole timing - through a chain of flit-aware routers is preserved. +- **컴포넌트별 스위칭/오버헤드 레이턴시** (`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. Approximated (with known directional error) +### D2. 근사됨 (알려진 방향성 오차와 함께) -| Effect | Real HW | Our model | Error direction | +| 효과 | 실제 HW | 본 모델 | 오차 방향 | |--------|---------|-----------|----------------| -| Router output port arbitration | Round-robin / weighted | Wire edge FIFO + serial worker | Fair when one txn per cycle; multi-stream sharing not modeled at flit level | -| HBM scheduler / write buffer | FR-FCFS + watermark drain | FIFO, no reordering | Pessimistic for mixed R/W when alternations are dense — default `switch_penalty_ns = 0` assumes ideal scheduler amortizes | -| Flit ↔ burst granularity | 32B flit < 256B burst | `flit_bytes = burst_bytes = 256B` | Sub-flit fine-grained timing noise; affects very small wire arbitration windows only | -| Wire-level RR fairness | Per-cycle multi-flow arbitration on shared link | Single serial wire process per edge | Fair only when one transaction is in flight on a given edge at a time. Multi-stream concurrent traffic on the same edge serializes by FIFO order | +| 라우터 출력 포트 중재 | 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. Ignored (out of scope) +### D3. 무시됨 (범위 외) -- Bank-level row buffer conflict penalty (assume no conflicts — best case; - the model has no per-bank state within a PC, so same-bank reuse cannot be - detected). -- HBM tRP / tRCD / tFAW / tRC timing constraints (absorbed into the steady-state - `burst_time = burst_bytes / pc_bw_gbs`). -- Refresh, ECC, thermal throttling, power gating. -- Clock domain crossings, PLL lock time. -- Upstream backpressure due to downstream buffer occupancy (input ports use - unbounded `simpy.Store`). -- Sub-flit cycle-level arbitration at routers (flit granularity is our - smallest unit). +- 뱅크 수준의 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. Workload sensitivity +### D4. 워크로드 민감도 -Workloads where the above simplifications meaningfully affect results: +위 단순화들이 결과에 의미 있게 영향을 미치는 워크로드: -- **Random scatter/gather**: bank conflict ignored → model optimistic. -- **Heavy mixed R/W intensive** (e.g., GEMM bias accumulation): HBM scheduler - absent. With default `switch_penalty_ns = 0` we assume ideal amortization; - setting it non-zero models pessimistic per-alternation cost. -- **High concurrency (>10 active flows on one link)**: HoL blocking and VC - limits not modeled → model optimistic. -- **Very small (sub-flit) transactions**: flit quantization noise. -- **Concurrent multi-flow on a single wire**: wire is serial FIFO at the - flit level, so per-flow fairness within a single edge is not modeled. - Pre-edge merging (multiple sources arriving at a router and being - forwarded to the same downstream wire) is correctly modeled via the - flit-aware router's serial worker. +- **무작위 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. Verification policy +### D5. 검증 정책 -For workloads in D4, cross-check against real HW or a cycle-accurate -simulator before drawing absolute-magnitude conclusions. The model remains -accurate for **relative comparisons** within the modeled regime. +D4의 워크로드에 대해 절대값 결론을 내리기 전에 실제 HW나 사이클 정확 +시뮬레이터와 cross-check 할 것. 모델은 모델링된 영역 내에서의 **상대적 +비교**에 대해서는 여전히 정확하다. -### D6. Future work +### D6. 향후 작업 -Note: multi-stream merging at routers IS modeled correctly — each -in_port has its own fan_in process, all push to a shared inbox, and -the router worker forwards in inbox FIFO order. Flits from different -upstream streams naturally interleave at flit granularity. The items -below are different concerns, ordered by expected workload impact. +참고: 라우터에서의 multi-stream 병합은 올바르게 모델링되고 있다 — 각 +in_port가 자신의 fan_in 프로세스를 가지며 모두 공유 인박스로 push하고, +라우터 워커가 인박스 FIFO 순서로 전달한다. 서로 다른 상위 스트림의 flit들이 +flit granularity에서 자연스럽게 인터리브된다. 아래 항목들은 별개의 관심사이며, +예상되는 워크로드 영향 순으로 정렬되어 있다. -**Higher impact (workload accuracy gap)**: +**영향이 큼 (워크로드 정확도 격차)**: -- [ ] **Bank-level conflict modeling** within a PC (opt-in via - `track_banks: true`). Currently we assume no same-bank reuse; - random scatter/gather workloads are optimistic here. -- [ ] **HBM scheduler** with write buffer + watermark drain (Tier 2 - from the design discussion). Default `switch_penalty_ns=0` is the - ideal-amortization stand-in; bursty mixed R/W workloads benefit - from explicit modeling. -- [ ] **Backpressure** modeling for finite component buffers. Matters - at high concurrency / sustained saturation where buffer occupancy - causes upstream stalls. -- [ ] **Op_log integration with chunk-streaming**: currently op_log - fires on PE-internal command messages (DmaReadCmd, DmaWriteCmd, - GemmCmd, MathCmd) which are not chunkified. Integration would - require flit-aware components to also emit op_log start/end hooks - per transaction (start on first flit, end on is_last). +- [ ] 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)을 함께 방출하도록 요구한다. -**Lower impact (academic / specific use cases)**: +**영향이 작음 (학술적 / 특정 use case)**: -- [ ] **Cycle-accurate router arbitration policies** (RR with - priorities, age, iSLIP). The FIFO inbox is already approximately - fair when flit arrival times differ slightly between streams (the - common case for similar-rate workloads). True impact appears only - for: (a) priority/QoS modeling, (b) per-stream tail latency - analysis under sustained saturation. Not critical for makespan or - average-latency studies. -- [ ] **Sub-flit (32B) granularity** for finer wire arbitration - cycles. Our `flit_bytes` equals burst (256B); real HW arbitrates - per 32B flit. Effect is small for most workloads (sub-flit timing - noise on small messages). +- [ ] **사이클 정확 라우터 중재 정책** (우선순위·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 -- Single review point for all model fidelity questions. Each future PR - touching latency must update the relevant section here. -- Workload-specific magnitude error envelopes are explicit. -- Builder-side derivation of `pc_bw_gbs = hbm_to_router_bw_gbs / num_pcs` - enforces the ADR-0017 D8 invariant in code rather than relying on yaml - manual consistency. -- Wire transfer time is charged once per bottleneck-link transit (Phase 2c - per-flit timing) rather than via terminal `drain_ns` injection. Single - transactions land at `drain + commit_time + small_overheads`; multi-hop - preserves wormhole pipelining; multi-stream merge correctly serializes - at the shared wire's FIFO. +- 모든 모델 충실도 질문에 대한 단일 리뷰 지점. 레이턴시를 건드리는 향후 + 모든 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 — component / port / wire model. -- ADR-0017 — Cube NOC architecture and HBM connectivity. -- ADR-0004 — memory semantics, local HBM. -- ADR-0034 — HBM controller internal design. +- ADR-0015 — 컴포넌트 / 포트 / 와이어 모델. +- ADR-0017 — 큐브 NOC 아키텍처 및 HBM 연결성. +- ADR-0004 — 메모리 시맨틱, 로컬 HBM. +- ADR-0034 — HBM 컨트롤러 내부 설계. diff --git a/docs/adr-ko/ADR-0034-dev-hbm-controller-internal-design.md b/docs/adr-ko/ADR-0034-dev-hbm-controller-internal-design.md index b7d3e8f..6a45e84 100644 --- a/docs/adr-ko/ADR-0034-dev-hbm-controller-internal-design.md +++ b/docs/adr-ko/ADR-0034-dev-hbm-controller-internal-design.md @@ -1,4 +1,4 @@ -# ADR-0034: HBM Controller Internal Design +# ADR-0034: HBM 컨트롤러 내부 설계 ## Status @@ -6,111 +6,108 @@ Accepted ## Context -`HbmCtrlComponent` is the per-PE HBM partition endpoint at the leaf of -the cube NOC. One instance is created per PE under the topology node -`sip{S}.cube{C}.hbm_ctrl.pe{idx}` and attaches to that PE's router -(ADR-0017 D4). The component models per-pseudo-channel (PC) scheduling, -burst-granular commit timing, address-based PC selection, and response -routing back to the requester. +`HbmCtrlComponent`는 큐브 NOC의 말단(leaf)에 위치하는 PE별 HBM +파티션 엔드포인트이다. 토폴로지 노드 +`sip{S}.cube{C}.hbm_ctrl.pe{idx}` 아래에 PE마다 하나의 인스턴스가 +생성되며 해당 PE의 라우터에 연결된다 (ADR-0017 D4). 본 컴포넌트는 +의사 채널(PC, pseudo-channel)별 스케줄링, 버스트 단위 커밋 타이밍, +주소 기반 PC 선택, 그리고 응답을 요청자에게 되돌리는 라우팅을 +모델링한다. -This ADR documents the component as currently implemented. ADR-0017 D4/D8 -defines *where* HBM CTRL attaches and *what* aggregate BW it must -deliver. ADR-0033 D1/D2 defines *what fidelity* of HBM modelling is in -scope. This ADR fills the gap between those two — the per-instance -internal scheduling model. +본 ADR은 현재 구현된 컴포넌트를 문서화한다. ADR-0017 D4/D8은 HBM CTRL이 +*어디에* 부착되는지와 *어떤* 집계 대역폭을 제공해야 하는지를 정의한다. +ADR-0033 D1/D2는 HBM 모델링의 *어떤 정밀도(fidelity)*가 범위에 포함되는지를 +정의한다. 본 ADR은 그 둘 사이의 공백 — 인스턴스별 내부 스케줄링 모델을 +채운다. ## Decision -### D1. Role +### D1. 역할 -`HbmCtrlComponent` is a per-PE HBM partition endpoint. One instance per -PE (default 8 per cube, set by `cube.memory_map.hbm_slices_per_cube`) -attaches to that PE's router via the `peX.hbm` attachment list in -`cube_mesh.yaml` (ADR-0017 D4). In the default n:1 channel mapping -(ADR-0017 D8) the instance aggregates `channels_per_pe` pseudo-channels -into one endpoint. +`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`개의 의사 채널을 하나의 엔드포인트로 집계한다. -The component models: +본 컴포넌트는 다음을 모델링한다: -- Per-PC scheduling (D2) with R/W command-bus sharing. -- Address-based PC selection (D3). -- Burst-granular commit timing (D4). -- Flit-aware per-flit PC commit and async finalize (D5, D6). -- Command-only Transaction handling for read-data drain (D7). -- Response routing back to the requester (D8). +- PC별 스케줄링(D2) 및 R/W 명령 버스 공유. +- 주소 기반 PC 선택(D3). +- 버스트 단위 커밋 타이밍(D4). +- Flit 인지 per-flit PC 커밋 및 비동기 finalize(D5, D6). +- 읽기 데이터 드레인(drain)을 위한 명령 전용 Transaction 처리(D7). +- 요청자에게 되돌리는 응답 라우팅(D8). -It does not model: +다음은 모델링하지 않는다: -- Bank-level row-buffer conflicts, refresh, ECC, thermal throttling +- Bank 수준의 row-buffer 충돌, refresh, ECC, 열 스로틀링 (ADR-0033 D3). -- Cross-PE HBM contention beyond its own router edge (handled by the - router mesh — ADR-0017 D3). -- 1:1 channel mode (ADR-0017 D8 future work). +- 자신의 라우터 엣지를 넘어가는 PE 간 HBM 경합(라우터 메시가 처리 — + ADR-0017 D3). +- 1:1 채널 모드(ADR-0017 D8 향후 작업). -### D2. Per-PC scheduling model +### D2. PC별 스케줄링 모델 -Per-instance state initialised in `start()`: +`start()`에서 초기화되는 인스턴스별 상태: -- `_pc_avail: list[float]` — earliest sim-time each PC is free; length - `num_pcs`, initial 0.0. -- `_pc_last_dir: list["R"|"W"|None]` — direction of the last commit on - each PC, used for switch-penalty detection (D4); initial `None`. +- `_pc_avail: list[float]` — 각 PC가 다음에 자유로워지는 가장 빠른 + 시뮬레이션 시각; 길이 `num_pcs`, 초기값 0.0. +- `_pc_last_dir: list["R"|"W"|None]` — 각 PC의 마지막 커밋 방향, 스위치 + 페널티 감지에 사용(D4); 초기값 `None`. -`num_pcs` and `burst_bytes` must each be a positive power of two so -that address-based PC selection (D3) reduces to a shift-and-mask. +`num_pcs`와 `burst_bytes`는 각각 양의 2의 거듭제곱이어야 주소 기반 PC +선택(D3)이 시프트와 마스크로 축약된다. -Read and write requests share the same `_pc_avail` slot per PC — the -real HW per-PC command bus is shared between read and write traffic, so -issuing a write to PC k blocks a subsequent read to PC k by exactly the -burst time. +읽기와 쓰기 요청은 PC별로 동일한 `_pc_avail` 슬롯을 공유한다 — 실제 HW에서 +PC별 명령 버스는 읽기와 쓰기 트래픽이 공유하므로, PC k에 쓰기를 발행하면 +PC k에 대한 후속 읽기가 정확히 버스트 시간만큼 블록된다. -Direction `dir` for a request is inferred from the request type: +요청의 방향 `dir`은 요청 타입으로부터 추론된다: - `MemoryWriteMsg` → `"W"`. -- `PeDmaMsg` with `is_write=True` → `"W"`. -- All others (`MemoryReadMsg`, `PeDmaMsg` read) → `"R"`. +- `is_write=True`인 `PeDmaMsg` → `"W"`. +- 그 외 전부(`MemoryReadMsg`, 읽기 `PeDmaMsg`) → `"R"`. -### D3. Address-based PC selection +### D3. 주소 기반 PC 선택 -PC index for an access is derived from the access address by shift and -mask: +접근에 대한 PC 인덱스는 접근 주소로부터 시프트와 마스크로 도출된다: ```text -pc_shift = log2(burst_bytes) # default 8 (burst=256B) -pc_mask = num_pcs - 1 # default 7 (8 PCs) +pc_shift = log2(burst_bytes) # 기본값 8 (burst=256B) +pc_mask = num_pcs - 1 # 기본값 7 (8 PCs) pc = (address >> pc_shift) & pc_mask ``` -Computed once in `start()` from topology config so alternative -`(burst_bytes, num_pcs)` pairs stay consistent. For the canonical -default `(256, 8)` this places the PC select field at bits `[10:8]` of -the HBM byte offset: bits `[7:0]` are within-burst (same PC), bits -`[10:8]` are the 3-bit PC index, bits `[36:11]` are row/bank/column -within the PC slice (see `phyaddr.py` comment). +대안적인 `(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` 주석 참조). -Address-based striping — as opposed to address-blind global -round-robin — preserves PC parallelism for offset-disjoint concurrent -transfers: each transfer's bursts land deterministically on the PC set -implied by its byte addresses, so multi-PE workloads accessing disjoint -regions do not collide on a single PC. +주소 기반 스트라이핑은 — 주소를 보지 않는 전역 라운드로빈과 달리 — +오프셋이 분리된 동시 전송들에 대해 PC 병렬성을 보존한다: 각 전송의 +버스트는 자신의 바이트 주소가 함의하는 PC 집합 위에 결정론적으로 +떨어지므로, 분리된 영역에 접근하는 멀티 PE 워크로드가 단일 PC에서 +충돌하지 않는다. -### D4. Burst granularity and PC commit timing +### D4. 버스트 단위 시간 및 PC 커밋 타이밍 -A single PC commit takes: +단일 PC 커밋에 걸리는 시간: ```text chunk_time = burst_bytes / pc_bw_gbs # ns ``` -- `burst_bytes` (default 256) is the burst granularity matching the - flit size (ADR-0033 D1). -- `pc_bw_gbs` is **builder-derived** from - `hbm_to_router_bw_gbs / num_pcs` (`topology/builder.py`), enforcing - the ADR-0017 D8 invariant that aggregate per-PE BW equals the - router-to-HBM link BW. +- `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의 + 불변식을 강제한다. -Per-PC commit scheduling for an arriving access on PC `pc` with -direction `dir`: +방향 `dir`로 PC `pc`에 도착한 접근에 대한 PC별 커밋 스케줄링: ```text switch_cost = switch_penalty_ns @@ -121,33 +118,32 @@ pc_avail[pc] = finish pc_last_dir[pc] = dir ``` -Default `switch_penalty_ns = 0` — Tier 0 assumption that an ideal HBM -scheduler amortises R/W switching cost (ADR-0033 D2). Non-zero values -model pessimistic per-alternation cost. +기본 `switch_penalty_ns = 0` — 이상적인 HBM 스케줄러가 R/W 스위칭 +비용을 분할 상환한다는 Tier 0 가정(ADR-0033 D2). 0이 아닌 값은 +교차마다 발생하는 비관적 비용을 모델링한다. -### D5. Flit-aware per-flit PC commit (primary path) +### D5. Flit 인지 per-flit PC 커밋 (주 경로) -`_handle_flit` is the primary worker path. For each arriving `Flit`: +`_handle_flit`이 주 워커 경로이다. 각 도착 `Flit`에 대해: -1. On the **first** flit of a transaction (`tid = id(txn)` not in - `_txn_state`): - - Apply `overhead_ns` once via `run(env, nbytes)` — header decode - model, first-flit overhead pattern (ADR-0033 D1). - - Initialise `_txn_state[tid] = {"last_finish": env.now}`. -2. Compute `pc = _pc_for_address(flit.address)` (D3). -3. Apply the per-PC schedule (D4) using the request direction (D2). -4. Update `state["last_finish"] = max(state["last_finish"], finish)`. -5. If `flit.is_last`: pop `_txn_state[tid]` and spawn `_finalize_txn` - (D6). +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 address-aware commit is the mechanism that lets concurrent -multi-PE traffic to disjoint HBM offsets pipeline through distinct PCs -in parallel. +per-flit 주소 인지 커밋이 분리된 HBM 오프셋으로 향하는 동시 멀티 PE +트래픽이 서로 다른 PC를 통해 병렬로 파이프라인되도록 하는 메커니즘이다. -### D6. Async finalize per transaction +### D6. 트랜잭션별 비동기 finalize -When a transaction's last flit has been scheduled, finalisation runs in -a separately-spawned process: +트랜잭션의 마지막 flit이 스케줄링되고 나면, finalize는 별도로 spawn된 +프로세스에서 실행된다: ```python def _finalize_txn(env, txn, last_finish): @@ -157,115 +153,111 @@ def _finalize_txn(env, txn, last_finish): yield from _send_response(env, txn) ``` -`_handle_flit` spawns this via `env.process(...)` and returns -immediately, so the worker can pick up the next inbox message while the -last PC commit drains. +`_handle_flit`은 이를 `env.process(...)`로 spawn한 뒤 즉시 반환하므로, +마지막 PC 커밋이 드레인되는 동안에도 워커는 다음 inbox 메시지를 집어들 +수 있다. -Without this split — i.e. if the worker itself did -`yield env.timeout(wait)` — concurrent single-flit transactions whose -addresses hit distinct PCs would still serialise at `chunk_time` each -inside the worker, hiding the PC parallelism that D3 and D5 are -designed to expose. +이 분리가 없다면 — 즉 워커 자신이 `yield env.timeout(wait)`를 한다면 — +서로 다른 PC에 떨어지는 주소를 가진 동시 단일 flit 트랜잭션들도 결국 +워커 내부에서 각각 `chunk_time`만큼 직렬화되어, D3와 D5가 노출하려고 +설계한 PC 병렬성을 숨겨버린다. -### D7. Non-flit fallback for command-only transactions +### D7. 명령 전용 트랜잭션을 위한 non-flit 폴백 -`_handle_txn` runs when the inbox delivers a `Transaction` rather than a -`Flit`. This is the path for command-only requests that the wire does -not chunk into flits — most notably `MemoryReadMsg` whose command txn -carries `nbytes=0` (data drain is modelled at HBM CTRL post-processing, -not as inbound flits). +`_handle_txn`은 inbox가 `Flit`이 아닌 `Transaction`을 전달할 때 실행된다. +이는 와이어가 flit으로 분할하지 않는 명령 전용 요청에 대한 경로로 — +대표적으로 명령 트랜잭션이 `nbytes=0`을 운반하는 `MemoryReadMsg`가 +해당한다(데이터 드레인은 HBM CTRL 후처리에서 모델링되며, 인바운드 +flit으로 모델링되지 않는다). -Procedure: +절차: 1. `work_bytes = txn.nbytes if txn.nbytes > 0 else int(request.nbytes or 0)` - — for read commands, work is sized by the request. -2. `n_chunks = ceil(work_bytes / burst_bytes)` if `work_bytes > 0` else - 0. -3. `chunk_interval = drain_ns / n_chunks` (when both > 0) — chunks are - scheduled over time at `drain/n_chunks` ns intervals to model the - bottleneck-link's data arrival rate (ADR-0033 D1 chunk-loop drain). -4. Apply `run(env, txn.nbytes)` once for `overhead_ns`. -5. For each chunk `i`, advance `chunk_interval` ns then apply the D4 - schedule with `pc = _pc_for_address(base_address + i * burst_bytes)`. -6. After scheduling all chunks, wait `last_finish - env.now` then call - `_send_response`. + — 읽기 명령의 경우 작업량은 요청으로 결정된다. +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` shares the same `_pc_avail` / `_pc_last_dir` state with -`_handle_flit` — there is exactly one source of PC scheduling truth -across both paths. +`_handle_txn`은 `_handle_flit`과 동일한 `_pc_avail` / `_pc_last_dir` +상태를 공유한다 — 두 경로에 걸쳐 PC 스케줄링의 단일 진실 원천이 정확히 +하나만 존재한다. -### D8. Response routing +### D8. 응답 라우팅 -`_send_response` dispatches on request type and path geometry: +`_send_response`는 요청 타입과 경로 형상에 따라 디스패치한다: -| Case | Trigger | Response | +| 경우 | 트리거 | 응답 | | --- | --- | --- | -| PE_DMA | `isinstance(txn.request, PeDmaMsg)` | New reverse-path Transaction (`is_response=True`, `nbytes=0`), same `done` | -| Bypass — Memory Read | `"m_cpu" not in any(txn.path)` AND `MemoryReadMsg` | Reverse-path Transaction with `nbytes=request.nbytes` (data return) | -| Bypass — Memory Write | `"m_cpu" not in any(txn.path)` AND not Memory Read | `txn.done.succeed()` (write completes locally) | -| Default | otherwise | New `ResponseMsg(correlation_id, request_id, src_cube, src_pe, success=True)` on reverse path | +| 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)` | -The "bypass" classification matches the Memory R/W fabric path defined -in ADR-0015 D4 (PCIE_EP → io_noc → ucie → cube router → hbm_ctrl, -without M_CPU). The PE_DMA case is its own dedicated reverse-path to -keep the inner-loop DMA fast (PE_DMA reads/writes do not synthesise a -ResponseMsg envelope). +"bypass" 분류는 ADR-0015 D4에서 정의된 Memory R/W 패브릭 경로(PCIE_EP → +io_noc → ucie → 큐브 라우터 → hbm_ctrl, M_CPU 미경유)와 일치한다. +PE_DMA 케이스는 내부 루프 DMA를 빠르게 유지하기 위한 전용 역방향 경로이다 +(PE_DMA 읽기/쓰기는 ResponseMsg 봉투를 합성하지 않는다). -In all reverse-path cases, the response Transaction is put onto -`out_ports[reverse_path[1]]` — the first hop back along the recorded -forward path. If `reverse_path` has fewer than 2 entries (degenerate -path), the original `txn.done` is signalled directly. +모든 역방향 경로 케이스에서, 응답 Transaction은 +`out_ports[reverse_path[1]]` — 기록된 정방향 경로를 따라 되돌아가는 첫 +홉 — 에 put된다. `reverse_path`의 엔트리가 2개 미만이면(축퇴된 경로), +원래의 `txn.done`이 직접 시그널된다. -### D9. Configurable attributes +### D9. 설정 가능한 속성 -| Attribute | Default | Source | Notes | +| 속성 | 기본값 | 출처 | 비고 | | --- | --- | --- | --- | -| `num_pcs` | 8 | topology cube `hbm_ctrl.attrs` | Must be power of 2 | -| `pc_bw_gbs` | 32.0 | builder-derived: `hbm_to_router_bw_gbs / num_pcs` | Enforces ADR-0017 D8 invariant | -| `burst_bytes` | 256 | topology attrs | Must be power of 2; equals `flit_bytes` (ADR-0033 D1) | -| `switch_penalty_ns` | 0.0 | topology attrs | Tier 0 default; non-zero models pessimistic R/W switching | -| `efficiency` | 1.0 | topology attrs | Applied at builder time to `hbm_to_router_bw_gbs` (router-edge BW scaling only) | -| `overhead_ns` | 0.0 | topology attrs | First-flit decode overhead (D5) | +| `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` is derived by `topology/builder.py` rather than configured -directly so the aggregate per-PE BW matches the router-to-HBM link BW -without yaml-side duplication. +`pc_bw_gbs`는 yaml 측 중복 없이 PE당 집계 대역폭을 라우터-HBM 링크 +대역폭과 일치시키기 위해 직접 설정되지 않고 `topology/builder.py`에서 +도출된다. ## Consequences ### Positive -- Address-based PC selection preserves multi-stream HBM parallelism - that an address-blind round-robin would collapse — important for - multi-PE workloads with disjoint HBM regions. -- Flit-aware path (D5) + async finalize (D6) preserves wormhole - pipelining and exposes PC parallelism for back-to-back single-flit - transactions. -- Single source of PC scheduling truth (D4 mechanism, used by both D5 - flit path and D7 chunk-loop path). -- Builder-derived `pc_bw_gbs` enforces ADR-0017 D8 in code, not yaml - discipline. +- 주소 기반 PC 선택은 주소를 보지 않는 라운드로빈이 무너뜨릴 멀티 스트림 + HBM 병렬성을 보존한다 — 분리된 HBM 영역을 갖는 멀티 PE 워크로드에서 + 중요하다. +- Flit 인지 경로(D5) + 비동기 finalize(D6)는 웜홀 파이프라이닝을 + 보존하며, 연속적인 단일 flit 트랜잭션에 대해 PC 병렬성을 노출한다. +- PC 스케줄링의 단일 진실 원천(D4 메커니즘이 D5 flit 경로와 D7 청크 루프 + 경로 모두에서 사용됨). +- 빌더 도출 `pc_bw_gbs`가 yaml 규율이 아닌 코드에서 ADR-0017 D8을 + 강제한다. ### Negative -- No bank-level conflict modelling within a PC; address-blind to - bank/row-buffer reuse (ADR-0033 D3). -- No HBM scheduler (FR-FCFS / write-buffer / watermark drain); fixed - FIFO per PC. Bursty mixed R/W is approximated by `switch_penalty_ns` +- 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` is a regular dict keyed by `id(txn)`; in-flight state - accumulates per concurrent transaction and is removed only on - `is_last`. Adequate for current workloads. +- `_txn_state`는 `id(txn)`로 키를 잡는 일반 dict이다; 동시 트랜잭션마다 + in-flight 상태가 누적되며 `is_last` 시에만 제거된다. 현재 워크로드에는 + 충분하다. ## Links -- ADR-0001 (Physical address layout — PC bit field comment) -- ADR-0015 D4 (Memory R/W fabric path — bypass response case) -- ADR-0017 D4 (Per-PE HBM partitioning — attachment to PE routers) -- ADR-0017 D8 (HBM channel mapping mode — n:1 aggregate this ADR - implements) -- ADR-0017 D9 (AddressResolver — `hbm_ctrl.pe{pe_id}` endpoint - resolution) -- ADR-0033 D1 (Modelled precisely — per-PC parallelism, switch penalty, - flit-aware PC commit, first-flit overhead, chunk-loop drain) -- ADR-0033 D2 (Switch-penalty default 0 — ideal scheduler amortisation) +- 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 — 이상적 스케줄러의 분할 상환) diff --git a/docs/adr-ko/ADR-0035-dev-m-cpu-and-m-cpu-dma-component-model.md b/docs/adr-ko/ADR-0035-dev-m-cpu-and-m-cpu-dma-component-model.md index 8d22a7b..699758b 100644 --- a/docs/adr-ko/ADR-0035-dev-m-cpu-and-m-cpu-dma-component-model.md +++ b/docs/adr-ko/ADR-0035-dev-m-cpu-and-m-cpu-dma-component-model.md @@ -1,4 +1,4 @@ -# ADR-0035: M_CPU and M_CPU.DMA Component Model +# ADR-0035: M_CPU 및 M_CPU.DMA 컴포넌트 모델 ## Status @@ -6,51 +6,47 @@ Accepted ## Context -M_CPU is the cube-level command processor. It receives commands from -IO_CPU (or from PCIE_EP when the engine routes Memory R/W through -M_CPU as a fallback), fans them out to the PEs in its cube, and -aggregates per-PE responses into a single ResponseMsg sent back to -IO_CPU on the reverse path. +M_CPU는 큐브 수준의 명령 프로세서이다. IO_CPU로부터(또는 엔진이 +Memory R/W를 폴백으로 M_CPU를 거쳐 라우팅할 때 PCIE_EP로부터) 명령을 +수신하여 자신의 큐브 내 PE들로 팬아웃하고, PE별 응답을 단일 ResponseMsg로 +집계하여 역방향 경로를 통해 IO_CPU로 되돌려 보낸다. -M_CPU.DMA is the cube-level DMA channel pair that handles Memory R/W -fan-out. Per ADR-0015 D5 it is **not** a separate topology node — -it lives as internal state of `MCpuComponent`. +M_CPU.DMA는 Memory R/W 팬아웃을 처리하는 큐브 수준의 DMA 채널 쌍이다. +ADR-0015 D5에 따라 별도의 토폴로지 노드가 **아니다** — `MCpuComponent`의 +내부 상태로서 존재한다. -This ADR documents the M_CPU component implementation that realizes -those responsibilities, including the three distinct fan-out paths -(Memory R/W, Kernel Launch, MMU Map/Unmap), the M_CPU.DMA resource -model, and the response aggregation contract. +본 ADR은 위의 책임들을 실현하는 M_CPU 컴포넌트 구현을 문서화한다. 여기에는 +세 가지 구별되는 팬아웃 경로(Memory R/W, Kernel Launch, MMU Map/Unmap), +M_CPU.DMA 자원 모델, 그리고 응답 집계 계약이 포함된다. ## Decision -### D1. Role +### D1. 역할 -M_CPU has three responsibilities: +M_CPU는 세 가지 책임을 갖는다: -1. **Transit forwarding** — when not the terminal hop (e.g., on the - reverse response path PE → M_CPU → IO_CPU), forwards Transactions - to `next_hop` in their pre-computed path. -2. **Multi-PE fan-out at terminal hop** — dispatches to one of three - fan-out paths based on request type (D2). -3. **Response aggregation** — collects per-PE responses, sends a - single aggregate ResponseMsg back to IO_CPU on the reverse path. +1. **Transit 포워딩** — 종단 홉이 아닐 때(예: 역방향 응답 경로 PE → + M_CPU → IO_CPU), 사전 계산된 경로의 `next_hop`으로 Transaction을 + 전달한다. +2. **종단 홉에서의 멀티 PE 팬아웃** — 요청 타입에 따라 세 팬아웃 경로 + 중 하나로 디스패치한다(D2). +3. **응답 집계** — PE별 응답을 수집하여 역방향 경로를 통해 단일 집계 + ResponseMsg를 IO_CPU로 되돌려 보낸다. -Per invocation (`run()`): applies `overhead_ns` once per incoming -Transaction. +호출당(`run()`): 들어오는 Transaction마다 `overhead_ns`를 한 번 적용한다. -M_CPU does **not**: +M_CPU는 다음을 하지 **않는다**: -- Decide routing — paths are pre-computed by the router (ADR-0002). -- Handle PE-internal execution — PE_CPU / PE_SCHEDULER / engines - (ADR-0014). -- Decode addresses — `ctx.resolver.resolve(pa)` returns the per-PE - `hbm_ctrl.pe{X}` directly (ADR-0017 D9). -- Interpret tensor or kernel semantics — fan-out dispatch by Python - isinstance check only. +- 라우팅 결정 — 경로는 라우터에 의해 사전 계산된다(ADR-0002). +- PE 내부 실행 처리 — PE_CPU / PE_SCHEDULER / 엔진들이 담당(ADR-0014). +- 주소 디코드 — `ctx.resolver.resolve(pa)`가 PE별 `hbm_ctrl.pe{X}`를 + 직접 반환한다(ADR-0017 D9). +- 텐서 또는 커널 의미 해석 — 팬아웃 디스패치는 Python isinstance + 체크만으로 이루어진다. -### D2. Three fan-out paths dispatched by request type +### D2. 요청 타입으로 디스패치되는 세 가지 팬아웃 경로 -At the terminal hop the worker dispatches by request type: +종단 홉에서 워커는 요청 타입에 따라 디스패치한다: ```python elif self.ctx is not None and txn.request is not None: @@ -62,99 +58,95 @@ elif self.ctx is not None and txn.request is not None: env.process(self._dma_fanout(env, txn)) ``` -Each path uses a different router method: +각 경로는 서로 다른 라우터 메서드를 사용한다: -- `_dma_fanout` uses `ctx.router.find_mcpu_dma_path()` — the - M_CPU-specific DMA path that avoids PE pipeline nodes. -- `_kernel_launch_fanout` uses `ctx.router.find_node_path()` — the - generic NOC command path to PE_CPU. -- `_mmu_msg_fanout` uses `ctx.router.find_node_path()` — NOC command - path to PE_MMU. +- `_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 internal subcomponent (ADR-0015 D5) +### D3. M_CPU.DMA 내부 서브 컴포넌트 (ADR-0015 D5) -`MCpuComponent.start()` initializes two SimPy resources: +`MCpuComponent.start()`는 두 개의 SimPy 자원을 초기화한다: ```python self._dma_write = simpy.Resource(env, capacity=1) # MemoryWriteMsg self._dma_read = simpy.Resource(env, capacity=1) # MemoryReadMsg ``` -Properties: +특성: -- **Not a topology node** — managed entirely inside `MCpuComponent`; - does not appear in `topology.yaml` or in the compiled graph. -- **Independent read and write channels** — concurrent in-flight - Memory R/W is allowed. -- **Capacity=1 per channel** serializes the **dispatch step** - (`yield self.out_ports[...].put(...)`) of concurrent in-flight Memory - R/W requests at this M_CPU. Actual fabric transfer time is modeled - by wire processes between components (ADR-0015 D2) and by - `drain_ns` at terminal hops; the DMA resource does not gate - transfer duration. +- **토폴로지 노드가 아님** — 전적으로 `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 자원은 전송 지속 시간을 + 게이팅하지 않는다. -Resource selection is request-type-based: +자원 선택은 요청 타입에 기반한다: ```python dma_res = self._dma_write if isinstance(request, MemoryWriteMsg) else self._dma_read ``` -### D4. Transit forwarding at non-terminal hops +### D4. 비종단 홉에서의 transit 포워딩 -When `txn.next_hop` is not None — typical for the reverse response -path (PE → M_CPU → IO_CPU) — the worker forwards normally: +`txn.next_hop`이 None이 아닐 때 — 전형적으로 역방향 응답 경로(PE → +M_CPU → IO_CPU)에서 — 워커는 정상적으로 전달한다: ```python if next_hop: yield self.out_ports[next_hop].put(txn.advance()) ``` -The fan-out branches fire only at the terminal hop. The same component -therefore serves both forward command dispatch and reverse response -relay roles. +팬아웃 분기는 종단 홉에서만 발화한다. 따라서 동일한 컴포넌트가 정방향 +명령 디스패치 역할과 역방향 응답 중계 역할을 모두 수행한다. -### D5. DMA fan-out (`_dma_fanout` — Memory R/W) +### D5. DMA 팬아웃 (`_dma_fanout` — Memory R/W) -For each Memory R/W request at terminal hop: +종단 홉에서 각 Memory R/W 요청에 대해: -1. `_resolve_dma_destinations(request)` returns a per-PE - `hbm_ctrl.pe{X}` derived from the request's PA via - `ctx.resolver.resolve(PhysAddr.decode(pa))` (ADR-0017 D9). -2. For each destination: - - Acquire the appropriate DMA resource (`_dma_write` or - `_dma_read`) via `with dma_res.request() as req`. - - Resolve path via `ctx.router.find_mcpu_dma_path()`. - - Compute `drain_ns = ctx.compute_drain_ns(path, nbytes)`. - - Create sub-Transaction carrying `drain_ns` and dispatch to - `path[1]`. -3. Track `max_drain_ns` across destinations and record it as - `txn.result_data["xfer_ns"]` after all responses arrive. -4. After all per-PE responses are collected (D8), send an aggregate - ResponseMsg on the reverse command path back to IO_CPU. +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 decode fallback (`f"{cube_prefix}.hbm_ctrl"`) is legacy dead code — -no such node exists after ADR-0017 D4's per-PE partitioning. Kept -defensively but does not route to a real destination. +PA 디코드 폴백(`f"{cube_prefix}.hbm_ctrl"`)은 레거시 데드 코드이다 — +ADR-0017 D4의 PE별 파티셔닝 이후로 그러한 노드는 존재하지 않는다. +방어적으로 남겨두었으나 실제 목적지로 라우팅되지는 않는다. -### D6. Kernel launch fan-out (`_kernel_launch_fanout`) +### D6. Kernel launch 팬아웃 (`_kernel_launch_fanout`) -For `KernelLaunchMsg` at terminal hop: +종단 홉에서 `KernelLaunchMsg`에 대해: -1. `_resolve_pe_ids(target_pe)` → list of PE ids in this cube. -2. For each PE: find path to `f"{cube_prefix}.pe{pe_id}.pe_cpu"` via - `ctx.router.find_node_path()`. -3. **`target_start_ns` handling** (ADR-0009 D5): - - If the request already carries `target_start_ns` (stamped by - IO_CPU per ADR-0036 D3): **pass through unchanged**. - - If absent (direct-to-M_CPU launch in unit tests): compute a - per-cube barrier `env.now + max(per-PE leg latency)` and stamp - via `dataclasses.replace`. -4. Dispatch sub-Transactions with `nbytes=0` (kernel launch is a - control message; preserving nbytes=0 keeps fan-out off the shared - first-hop fabric BW, mirroring ADR-0036 D4). -5. After all per-PE responses arrive (D8), aggregate per-PE metrics - from each sub-Transaction's `result_data` into the parent - transaction: +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)) @@ -162,125 +154,120 @@ For `KernelLaunchMsg` at terminal hop: txn.result_data["compute_ns"] = max(existing, max(compute_values)) ``` - The max-merge with the existing value matters because cross-cube - IO_CPU fan-out shares the same parent `result_data`; merging - prevents one cube from clobbering another's metric. -6. Send aggregate ResponseMsg on reverse path back to IO_CPU. + 기존 값과의 max 병합이 중요한 이유는 크로스 큐브 IO_CPU 팬아웃이 + 동일한 부모 `result_data`를 공유하기 때문이다; 병합을 통해 한 큐브가 + 다른 큐브의 메트릭을 덮어쓰는 일을 방지한다. +6. IO_CPU로 되돌아가는 역방향 경로로 집계 ResponseMsg를 전송. -### D7. MMU map/unmap fan-out (`_mmu_msg_fanout`) +### D7. MMU map/unmap 팬아웃 (`_mmu_msg_fanout`) -For `MmuMapMsg` / `MmuUnmapMsg` at terminal hop: +종단 홉에서 `MmuMapMsg` / `MmuUnmapMsg`에 대해: -1. `_resolve_pe_ids(target_pe)` → PE ids. -2. For each PE: find path to `f"{cube_prefix}.pe{pe_id}.pe_mmu"` via - `find_node_path()`. -3. Dispatch sub-Transactions with `nbytes=0`. -4. PE_MMU is a terminal node — it does **not** send a ResponseMsg - back. Instead, the sub-Transaction's own `sub_done` event is the - completion signal. -5. Wait for all `sub_done` events in-line (does **not** use - `_pending` counter — D8 is for response-bearing fan-out only). -6. Send aggregate ResponseMsg on reverse path back to IO_CPU. +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. Response aggregation (`_pending` + `_parent_txns`) +### D8. 응답 집계 (`_pending` + `_parent_txns`) -For DMA and kernel-launch fan-out (which expect per-PE ResponseMsg -arriving on the reverse path): +DMA 및 kernel-launch 팬아웃(역방향 경로로 도착하는 PE별 ResponseMsg를 +예상함)에 대해: ```python self._pending: dict[str, tuple[int, int, simpy.Event]] = {} self._parent_txns: dict[str, Any] = {} ``` -- On dispatch: register `(expected, received=0, all_done)` and - remember the parent transaction. -- `_worker` recognises responses by `is_response=True` and routes - them to `_collect_response`, which increments `received` and - signals `all_done` when `received >= expected`. -- After `yield all_done`, the fan-out path constructs the aggregate - ResponseMsg: +- 디스패치 시: `(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 aggregate, not a single PE - success=True, # no failure semantics implemented + src_pe=-1, # -1 = M_CPU 집계, 단일 PE가 아님 + success=True, # 실패 의미는 구현되어 있지 않음 ) ``` -- The response Transaction travels on `list(reversed(txn.path))` - back to IO_CPU. +- 응답 Transaction은 `list(reversed(txn.path))`를 따라 IO_CPU로 + 되돌아간다. -MMU fan-out (D7) uses a simpler in-line list of `sub_done` events -because PE_MMU is terminal — there is no ResponseMsg path to -intercept. +MMU 팬아웃(D7)은 PE_MMU가 종단이므로 더 단순한 `sub_done` 이벤트의 +인라인 리스트를 사용한다 — 가로챌 ResponseMsg 경로가 없다. -### D9. Helpers and configurable attribute +### D9. 헬퍼와 설정 가능한 속성 `_resolve_pe_ids(target_pe)`: - `int` → `[target_pe]` - `tuple[int, ...]` → `list(target_pe)` -- `"all"` → `range(n_slices)` where `n_slices` comes from cube - `memory_map.hbm_slices_per_cube` (default 8). +- `"all"` → `range(n_slices)`, 여기서 `n_slices`는 큐브 + `memory_map.hbm_slices_per_cube`(기본 8)에서 가져온다. -Used by kernel-launch and MMU fan-out paths. +Kernel-launch 및 MMU 팬아웃 경로에서 사용된다. -Single configurable attribute drives per-instance latency: +인스턴스별 레이턴시를 결정하는 단일 설정 가능 속성: -| Site | impl name | overhead_ns | +| 사이트 | impl 이름 | overhead_ns | | --- | --- | --- | -| Cube `m_cpu` | `builtin.m_cpu` | 5.0 | +| 큐브 `m_cpu` | `builtin.m_cpu` | 5.0 | -Applied once in `run()` per Transaction — models command -interpretation and dispatch-decision time at M_CPU. +Transaction마다 `run()`에서 한 번 적용 — M_CPU에서의 명령 해석 및 +디스패치 결정 시간을 모델링한다. ## Consequences ### Positive -- Three fan-out paths are clearly separated by request type — adding - a new request kind is an isinstance branch + one fan-out method. -- M_CPU.DMA channels are independent (read and write run concurrently) - and serialize only the dispatch step at capacity=1. -- Transit-vs-terminal behavior is a single `if next_hop` check, so - the same component handles forward dispatch and reverse response - relay without role duplication. -- `target_start_ns` passthrough (D6) preserves the cross-cube barrier - established by IO_CPU (ADR-0036 D3), while the fallback computation - keeps direct-to-M_CPU unit tests working. -- Per-PE metric `max`-merge against existing parent `result_data` - values is robust to cross-cube IO_CPU fan-out sharing the same - parent. +- 세 가지 팬아웃 경로가 요청 타입에 의해 명확히 분리됨 — 새로운 요청 + 종류 추가는 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 -- No partial-failure semantics — a missing per-PE response stalls the - parent `all_done` indefinitely. Acceptable for simulation; not - suitable as a production-style endpoint. -- `_resolve_dma_destinations`'s cube-wide hbm_ctrl fallback is dead - code (no such node exists post-ADR-0017 D4). Kept defensively; - invites confusion and merits a follow-up cleanup. -- DMA resource serialization applies only at dispatch (the `put` call - is instantaneous in unbounded stores). The capacity=1 channel - models "one request in flight at a time at this M_CPU", not - "transfer duration serialization" — readers must consult wire - processes (ADR-0015 D2) and `drain_ns` for actual transfer - parallelism. +- 부분 실패 의미가 없음 — 누락된 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 fan-out and aggregation completion semantics) -- ADR-0009 D5 (`target_start_ns` — passed through unchanged when - present; computed as per-cube barrier when absent) -- ADR-0011 D-VA3 (MmuMapMsg fabric path includes M_CPU as PE fan-out - point) -- ADR-0014 D4 (DMA engine capacity=1; M_CPU.DMA mirrors the same - contract at cube level) -- ADR-0015 D5 (M_CPU.DMA is internal subcomponent of M_CPU, not a - topology node) -- ADR-0017 D9 (AddressResolver returns per-PE `hbm_ctrl.pe{X}`) -- ADR-0036 D3 / D4 (IO_CPU stamps `target_start_ns`; M_CPU passes - through unchanged; nbytes=0 invariant preserved through fan-out) +- 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 불변식 보존) diff --git a/docs/adr-ko/ADR-0036-dev-io-cpu-component-model.md b/docs/adr-ko/ADR-0036-dev-io-cpu-component-model.md index b79e9ad..7e5a008 100644 --- a/docs/adr-ko/ADR-0036-dev-io-cpu-component-model.md +++ b/docs/adr-ko/ADR-0036-dev-io-cpu-component-model.md @@ -1,4 +1,4 @@ -# ADR-0036: IO_CPU Component Model +# ADR-0036: IO_CPU 컴포넌트 모델 ## Status @@ -6,72 +6,70 @@ Accepted ## Context -IO_CPU is the IO chiplet's host-facing endpoint inside the simulation -graph. PCIE_EP receives host messages from the runtime API and routes -them via the io_noc; for command-bearing requests (KernelLaunch, -MmuMap/Unmap) the io_noc forwards to IO_CPU, which: +IO_CPU는 시뮬레이션 그래프 내부의 IO 칩렛 호스트 대향 엔드포인트이다. +PCIE_EP는 런타임 API로부터 호스트 메시지를 수신하여 io_noc를 통해 +라우팅한다; 명령을 동반하는 요청(KernelLaunch, MmuMap/Unmap)의 경우 +io_noc는 IO_CPU로 전달하며, IO_CPU는 다음을 수행한다: -- Fans out the request to per-cube M_CPUs. -- Aggregates per-cube responses into a single host-visible completion. -- For kernel launches, stamps a global `target_start_ns` barrier so - every PE across every targeted cube begins kernel body execution at - the same simulated time (ADR-0009 D5). +- 요청을 큐브별 M_CPU로 팬아웃. +- 큐브별 응답을 단일 호스트 가시 완료로 집계. +- 커널 런치의 경우, 타깃이 된 모든 큐브의 모든 PE가 동일한 시뮬레이션 + 시각에 커널 본체 실행을 시작하도록 전역 `target_start_ns` 배리어를 + 스탬프함(ADR-0009 D5). -Memory R/W traffic bypasses IO_CPU per ADR-0015 D4 / ADR-0016 D3; -this component therefore handles only command-plane traffic in normal -operation. +Memory R/W 트래픽은 ADR-0015 D4 / ADR-0016 D3에 따라 IO_CPU를 우회한다; +따라서 본 컴포넌트는 정상 동작에서 명령 평면 트래픽만을 처리한다. -This ADR documents the IO_CPU component implementation that realizes -those responsibilities. +본 ADR은 위의 책임을 실현하는 IO_CPU 컴포넌트 구현을 문서화한다. ## Decision -### D1. Role +### D1. 역할 -IO_CPU is the host-facing endpoint of the IO chiplet. It has two -primary responsibilities: +IO_CPU는 IO 칩렛의 호스트 대향 엔드포인트이다. 두 가지 주요 책임을 +갖는다: -1. **Multi-cube fan-out** — distribute KernelLaunchMsg / MmuMapMsg / - MmuUnmapMsg to per-cube M_CPUs. -2. **Response aggregation** — collect per-cube ResponseMsg, signal - parent `txn.done` when all targeted cubes have responded. +1. **멀티 큐브 팬아웃** — KernelLaunchMsg / MmuMapMsg / MmuUnmapMsg를 + 큐브별 M_CPU로 분배. +2. **응답 집계** — 큐브별 ResponseMsg를 수집하고, 타깃이 된 모든 큐브가 + 응답한 후 부모 `txn.done`을 시그널. -A third, narrower responsibility applies only to KernelLaunchMsg: -**`target_start_ns` global barrier stamping** (D3). +세 번째이자 더 좁은 책임은 KernelLaunchMsg에만 적용된다: +**`target_start_ns` 전역 배리어 스탬핑**(D3). -The component does **not**: +본 컴포넌트는 다음을 하지 **않는다**: -- Decide routing — paths are pre-computed by the router (ADR-0002). -- Decode tensor or kernel internals — those concerns belong to - M_CPU / PE_CPU / engines. -- Handle PE-level fan-out — M_CPU fans out within a cube (ADR-0009 D3). -- Handle Memory R/W data path — those bypass IO_CPU per ADR-0015 D4 - and ADR-0016 D3 (Memory R/W resolution code in - `_resolve_cube_targets` exists as a defensive fallback only). +- 라우팅 결정 — 경로는 라우터에 의해 사전 계산된다(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 해석 코드는 + 방어적 폴백으로만 존재). -Per invocation (`run()`): applies the configured `overhead_ns` once -per incoming Transaction (D8). +호출당(`run()`): 들어오는 Transaction마다 설정된 `overhead_ns`를 한 번 +적용한다(D8). -### D2. Forward path — multi-cube fan-out +### D2. 정방향 경로 — 멀티 큐브 팬아웃 -When a non-response Transaction arrives, the worker: +응답이 아닌 Transaction이 도착하면, 워커는: -1. Pays `overhead_ns` via `run()`. -2. Calls `_resolve_cube_targets` to derive the list of `(sip, cube)` - targets from the request (D5). -3. For each target: - - Resolves M_CPU node id via `ctx.resolver.find_m_cpu(sip, cube)`. - - Resolves the path via `ctx.router.find_node_path(io_cpu, m_cpu)`. - - Creates a per-cube sub-Transaction with `path` populated and - forwards it to `path[1]` (the first hop on the io_noc). -4. Registers aggregation state: `_pending[request_id] = (expected, - received=0, parent_done)`. +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` global barrier (ADR-0009 D5) +### D3. KernelLaunch `target_start_ns` 전역 배리어 (ADR-0009 D5) -IO_CPU is the canonical stamper for `target_start_ns`. When the -request is a `KernelLaunchMsg`, IO_CPU computes a single global -barrier covering every targeted PE across every targeted cube: +IO_CPU는 `target_start_ns`의 정규 스탬퍼이다. 요청이 +`KernelLaunchMsg`일 때, IO_CPU는 타깃이 된 모든 큐브의 모든 PE를 포괄하는 +단일 전역 배리어를 계산한다: ```text for (sip, cube) in cube_targets: @@ -85,132 +83,123 @@ for (sip, cube) in cube_targets: target_start_ns = env.now + global_max ``` -The request is then replaced (via `dataclasses.replace`) so the -stamped value propagates through the fan-out. +이후 요청은 (`dataclasses.replace`를 통해) 교체되어 스탬프된 값이 팬아웃 +전반에 전파된다. -Two overhead corrections: +두 가지 오버헤드 보정: -- `io_overhead_ns` is subtracted because IO_CPU has already paid it - in `run()` before this method runs. -- `m_overhead_ns` is subtracted once because it appears as the - endpoint of leg1 *and* the start of leg2 in path latency, but - M_CPU pays it only once at run time. +- `io_overhead_ns`는 차감되는데, IO_CPU가 본 메서드 실행 전에 `run()`에서 + 이미 지불했기 때문이다. +- `m_overhead_ns`는 한 번 차감되는데, 경로 레이턴시에서 leg1의 종단점인 + 동시에 leg2의 시작점으로 두 번 등장하지만 M_CPU는 런타임에 단 한 번만 + 지불하기 때문이다. -Every downstream PE_CPU yields until `target_start_ns` before -beginning kernel body execution; all PEs therefore start at the same -simulated time regardless of how long their individual dispatch path -took. +모든 다운스트림 PE_CPU는 커널 본체 실행을 시작하기 전 `target_start_ns` +까지 yield한다; 이를 통해 개별 디스패치 경로가 얼마나 오래 걸렸는지와 +무관하게 모든 PE가 동일한 시뮬레이션 시각에 시작한다. -### D4. KernelLaunch sub-Transactions carry `nbytes=0` +### D4. KernelLaunch 서브 Transaction은 `nbytes=0`을 운반 -Per-cube sub-Transactions for KernelLaunchMsg force `nbytes=0`, -overriding the parent `txn.nbytes`: +KernelLaunchMsg의 큐브별 서브 Transaction은 부모 `txn.nbytes`를 무시하고 +`nbytes=0`을 강제한다: -- Kernel launch is a control message; payload size is irrelevant at - the data-fabric level. -- If `nbytes > 0`, every per-cube sub-txn occupies fabric BW on the - io_noc's shared first hop. With 16 cubes this serializes fan-out, - pushing far M_CPUs past `target_start_ns` and breaking the D3 - invariant. +- 커널 런치는 제어 메시지이다; 데이터 패브릭 수준에서 페이로드 크기는 + 무관하다. +- `nbytes > 0`이면 모든 큐브별 서브 트랜잭션이 io_noc의 공유 first-hop + 패브릭 BW를 점유한다. 16개 큐브에서는 이로 인해 팬아웃이 직렬화되어 + 먼 M_CPU들이 `target_start_ns`를 지나치게 되고 D3 불변식이 깨진다. -Non-KernelLaunch sub-Transactions preserve `txn.nbytes` (only relevant -for the defensive Memory R/W fallback path, which carries actual -payload sizes). +KernelLaunch가 아닌 서브 Transaction은 `txn.nbytes`를 보존한다(실제 +페이로드 크기를 운반하는 방어적 Memory R/W 폴백 경로에만 관련됨). -### D5. Per-request-type cube target resolution +### D5. 요청 타입별 큐브 타깃 해석 -`_resolve_cube_targets` dispatches by request type: +`_resolve_cube_targets`는 요청 타입에 따라 디스패치한다: -| Request type | Source of `(sip, cube)` | `target_cubes="all"` semantics | +| 요청 타입 | `(sip, cube)`의 출처 | `target_cubes="all"` 의미 | | --- | --- | --- | -| `MemoryWriteMsg` | `dst_sip`, `dst_cube` (or `PhysAddr.decode(dst_pa).die_id` fallback) | single cube derived from PA decode | -| `MemoryReadMsg` | `src_sip`, `src_cube` (or `PhysAddr.decode(src_pa).die_id` fallback) | single cube derived from PA decode | -| `KernelLaunchMsg` | tensor shards filtered by `shard.sip == my_sip` | every cube that owns a shard on this SIP | -| `MmuMapMsg` / `MmuUnmapMsg` | `target_cubes` list, filtered to this SIP | `range(cubes_per_sip)` from spec | +| `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)` | -Each IO_CPU instance fans out only within its own SIP — `_my_sip()` -parses the SIP id from the node id (e.g., `sip0.io0.io_cpu` → 0). +각 IO_CPU 인스턴스는 자기 SIP 내에서만 팬아웃한다 — `_my_sip()`이 +노드 id에서 SIP id를 파싱한다(예: `sip0.io0.io_cpu` → 0). -The Memory R/W rows exist for defensive completeness; the engine's -normal path routes Memory R/W via `_process_memory_direct()` / -`find_memory_path()`, bypassing IO_CPU entirely (ADR-0015 D4 / -ADR-0016 D3). +Memory R/W 행은 방어적 완전성을 위해 존재한다; 엔진의 정상 경로는 +Memory R/W를 `_process_memory_direct()` / `find_memory_path()`로 +라우팅하여 IO_CPU를 완전히 우회한다(ADR-0015 D4 / ADR-0016 D3). -### D6. Response aggregation +### D6. 응답 집계 `_pending: dict[request_id → (expected, received, parent_done)]`: -- On dispatch: register `(len(cube_targets), 0, txn.done)`. -- `_worker` recognises responses by `is_response=True` and routes - them to `_collect_response`. -- `_collect_response` increments `received`; when `received >= - expected`, `parent_done.succeed()` is invoked and the entry is - removed from `_pending`. +- 디스패치 시: `(len(cube_targets), 0, txn.done)`을 등록. +- `_worker`는 `is_response=True`로 응답을 인식하여 `_collect_response`로 + 라우팅한다. +- `_collect_response`는 `received`를 증가시키며, `received >= expected`가 + 되면 `parent_done.succeed()`를 호출하고 엔트리를 `_pending`에서 + 제거한다. -This is a simple per-request counter. There is no per-cube identity -tracking and no partial-failure handling — a missing response -indefinitely stalls the parent done. Production-style failure paths -are out of scope for the current simulator model. +이는 단순한 요청별 카운터이다. 큐브별 정체성 추적이나 부분 실패 처리는 +없다 — 누락된 응답은 부모 done을 무기한 스톨시킨다. 프로덕션 스타일의 +실패 경로는 현재 시뮬레이터 모델의 범위 밖이다. -### D7. `target_pe` resolution helper +### D7. `target_pe` 해석 헬퍼 `_resolve_pe_ids(target_pe)`: - `int` → `[target_pe]`. - `tuple[int, ...]` → `list(target_pe)`. -- `"all"` → `range(n_slices)`, where `n_slices` comes from cube - `memory_map.hbm_slices_per_cube` (default 8). +- `"all"` → `range(n_slices)`, 여기서 `n_slices`는 큐브 + `memory_map.hbm_slices_per_cube`(기본 8)에서 가져온다. -Used in D3's barrier computation to enumerate every PE target per -cube. +D3의 배리어 계산에서 큐브별로 모든 PE 타깃을 열거하는 데 사용된다. -### D8. Configurable `overhead_ns` +### D8. 설정 가능한 `overhead_ns` -A single attribute drives per-instance latency: +단일 속성이 인스턴스별 레이턴시를 결정한다: -| Site | impl name | overhead_ns | +| 사이트 | impl 이름 | overhead_ns | | --- | --- | --- | -| IO chiplet `io_cpu` | `builtin.io_cpu` | 10.0 | +| IO 칩렛 `io_cpu` | `builtin.io_cpu` | 10.0 | -Applied once in `run()` per Transaction. Models command -interpretation + dispatch-decision time at IO_CPU. +Transaction마다 `run()`에서 한 번 적용된다. IO_CPU에서의 명령 해석 및 +디스패치 결정 시간을 모델링한다. ## Consequences ### Positive -- Cross-cube and cross-SIP kernel launches share a single global - barrier (D3 + D4) — no per-cube divergence in start time. -- nbytes=0 invariant keeps fan-out off the shared first-hop fabric - BW, preserving the barrier's accuracy at scale (16 cubes). -- Response aggregation via a single counter → minimal state, - deterministic ordering of completion. -- Per-SIP scoping (`_my_sip()`) keeps IO_CPUs in different SIPs - cleanly independent. +- 크로스 큐브 및 크로스 SIP 커널 런치가 단일 전역 배리어를 공유한다 + (D3 + D4) — 시작 시각의 큐브별 분기가 없다. +- nbytes=0 불변식이 팬아웃을 공유 first-hop 패브릭 BW로부터 떼어내, + 대규모(16 큐브)에서도 배리어의 정확도를 보존한다. +- 단일 카운터를 통한 응답 집계 → 최소 상태, 결정론적 완료 순서. +- SIP별 스코핑(`_my_sip()`)이 서로 다른 SIP의 IO_CPU들을 깨끗이 + 독립시킨다. ### Negative -- No partial-failure semantics — a missing per-cube response - indefinitely stalls the parent. Adequate for simulation but not - suitable as a production-style endpoint. -- `_pending` is a regular dict; in-flight requests accumulate state. - Acceptable for current benchmark workloads (few concurrent - outstanding launches); unbounded in principle. -- The Memory R/W resolution branches in `_resolve_cube_targets` are - dead code in the normal engine path. Kept defensively but invite - drift if the bypass path ever changes. +- 부분 실패 의미가 없음 — 누락된 큐브별 응답은 부모를 무기한 + 스톨시킨다. 시뮬레이션 용도로는 충분하나 프로덕션 스타일의 + 엔드포인트로는 적합하지 않다. +- `_pending`은 일반 dict이다; in-flight 요청이 상태로 누적된다. 현재 + 벤치마크 워크로드(미해결 런치가 적음)에는 허용 가능하나, 원리적으로는 + 무한하다. +- `_resolve_cube_targets`의 Memory R/W 해석 분기는 정상 엔진 경로에서 + 데드 코드이다. 방어적으로 남겨두었으나 우회 경로가 변경되면 드리프트 + 위험을 초래한다. ## Links -- ADR-0002 (Routing distance — path computation) -- ADR-0009 D1 (Kernel launch is an endpoint request to IO_CPU) -- ADR-0009 D3 (M_CPU fans out within a cube; IO_CPU fans out across - cubes) -- ADR-0009 D5 (target_start_ns canonical stamping at IO_CPU) -- ADR-0011 D-VA3 (MmuMapMsg routes through IO_CPU for cube fan-out) -- ADR-0012 (Host ↔ IO_CPU message schema) -- ADR-0015 D4 (Memory R/W bypasses IO_CPU; Kernel Launch via IO_CPU) -- ADR-0016 D1 (IO chiplet io_noc — IO_CPU attaches here) -- ADR-0016 D3 (Memory R/W path bypasses IO_CPU) -- ADR-0016 D4 (Kernel Launch path through IO_CPU for command - interpretation) +- 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 경유 커널 런치 경로) diff --git a/docs/adr-ko/ADR-0037-dev-forwarding-component.md b/docs/adr-ko/ADR-0037-dev-forwarding-component.md index 193dbe0..9063d7f 100644 --- a/docs/adr-ko/ADR-0037-dev-forwarding-component.md +++ b/docs/adr-ko/ADR-0037-dev-forwarding-component.md @@ -1,4 +1,4 @@ -# ADR-0037: Forwarding Component (forwarding_v1) +# ADR-0037: Forwarding 컴포넌트 (forwarding_v1) ## Status @@ -6,86 +6,78 @@ Accepted ## Context -The simulation graph has many node positions that exist purely to model -fabric traversal — NOC mesh routers, switches, UCIe protocol endpoints, -IO chiplet io_noc, transit cubes. These share a common pattern: receive -a message, apply per-component overhead (modeling header decode + -routing decision time), forward to the next hop along the pre-computed -path. +시뮬레이션 그래프에는 순전히 패브릭 통과를 모델링하기 위해 존재하는 노드 +위치들이 많다 — NOC 메시 라우터, 스위치, UCIe 프로토콜 엔드포인트, IO +칩렛 io_noc, transit 큐브. 이들은 공통 패턴을 공유한다: 메시지를 수신하고, +컴포넌트별 오버헤드(헤더 디코드 + 라우팅 결정 시간을 모델링)를 적용하며, +사전 계산된 경로를 따라 다음 홉으로 전달한다. -This ADR defines the contract for these transit nodes: a single -component type (`TransitComponent`) that handles flit-aware forwarding -with wormhole cut-through semantics, used under multiple impl names -according to the conceptual role each instance plays. +본 ADR은 이러한 transit 노드에 대한 계약을 정의한다: 웜홀 cut-through +의미로 flit 인지 포워딩을 처리하는 단일 컴포넌트 타입(`TransitComponent`)이며, +각 인스턴스가 수행하는 개념적 역할에 따라 여러 impl 이름 아래에 사용된다. ## Decision -### D1. Role +### D1. 역할 -The Forwarding component (`TransitComponent` class) is a **stateless -transit node** in the simulation graph. It models any fabric position -where a message physically traverses but no semantic processing -happens. +Forwarding 컴포넌트(`TransitComponent` 클래스)는 시뮬레이션 그래프의 +**상태 없는 transit 노드**이다. 메시지가 물리적으로 통과하지만 의미론적 +처리는 일어나지 않는 모든 패브릭 위치를 모델링한다. -Per traversal, the component: +통과당 컴포넌트는: -1. Reads an incoming Transaction or Flit from an `in_port`. -2. Applies the configured per-component overhead (`overhead_ns`), - applied **once per Transaction** even across multi-flit payloads - (see D2). -3. Looks up the next hop along the Transaction's pre-computed `path`. -4. Forwards to the corresponding `out_port`; at the terminal node - (no next hop), signals `txn.done` once the `is_last` flit arrives. +1. `in_port`에서 들어오는 Transaction 또는 Flit을 읽는다. +2. 설정된 컴포넌트별 오버헤드(`overhead_ns`)를 적용한다. 멀티 flit + 페이로드라도 **Transaction당 한 번** 적용된다(D2 참조). +3. Transaction의 사전 계산된 `path`를 따라 다음 홉을 조회한다. +4. 해당 `out_port`로 전달한다; 종단 노드(다음 홉 없음)에서는 `is_last` + flit이 도착하면 `txn.done`을 시그널한다. -The component **does NOT**: +본 컴포넌트는 다음을 하지 **않는다**: -- Decide routing — paths are pre-computed by the router (ADR-0002 / - ADR-0017 D2). Forwarding only executes the per-hop step. -- Model wire propagation or bandwidth occupancy — separate wire - processes between components handle that (ADR-0015 D2). -- Resolve addresses — the AddressResolver does that (ADR-0017 D9). -- Aggregate completion — terminal endpoints (IO_CPU, M_CPU, HBM_CTRL) - handle that. +- 라우팅 결정 — 경로는 라우터에 의해 사전 계산된다(ADR-0002 / + ADR-0017 D2). Forwarding은 홉별 단계만 실행한다. +- 와이어 전파나 대역폭 점유 모델링 — 컴포넌트 사이의 별도 와이어 + 프로세스가 처리한다(ADR-0015 D2). +- 주소 해석 — AddressResolver가 담당한다(ADR-0017 D9). +- 완료 집계 — 종단 엔드포인트(IO_CPU, M_CPU, HBM_CTRL)가 담당한다. -### D2. First-flit overhead model (header decode) +### D2. First-flit 오버헤드 모델 (헤더 디코드) -Per-Transaction `overhead_ns` is applied **exactly once**, at first -flit arrival: +Transaction별 `overhead_ns`는 첫 flit 도착 시 **정확히 한 번** 적용된다: -- `_txn_decoded: set[int]` tracks which Transactions have already - paid the overhead at this node. -- On first-flit arrival for a Transaction: `yield self.run(env, - msg.txn.nbytes)` — pays the overhead. -- Subsequent flits of the same Transaction skip the overhead — they - pipeline through with no extra delay. -- On `is_last` flit: remove the Transaction from `_txn_decoded`. +- `_txn_decoded: set[int]`이 본 노드에서 이미 오버헤드를 지불한 + Transaction들을 추적한다. +- 어떤 Transaction의 첫 flit 도착 시: `yield self.run(env, msg.txn.nbytes)` + — 오버헤드를 지불한다. +- 동일 Transaction의 후속 flit들은 오버헤드를 건너뛰고 추가 지연 없이 + 파이프라인 통과한다. +- `is_last` flit 시: Transaction을 `_txn_decoded`에서 제거한다. -This models the real-HW behavior where header decode and routing -decision happen once on first flit; payload flits then stream through -the same path (wormhole cut-through). Multi-hop pipelining emerges -naturally — each hop adds its own first-flit overhead, but flits -after the first do not re-pay overhead at any hop they have already -passed first. +이는 실제 HW의 동작 — 헤더 디코드와 라우팅 결정이 첫 flit에서 한 번 +일어나고, 이후 페이로드 flit들은 같은 경로로 스트리밍되는(웜홀 +cut-through) — 을 모델링한다. 멀티 홉 파이프라이닝은 자연스럽게 +발현된다 — 각 홉이 자신의 first-flit 오버헤드를 추가하지만, 첫 flit +이후의 flit들은 이미 첫 flit이 통과한 어떤 홉에서도 오버헤드를 다시 +지불하지 않는다. -### D3. Serial worker forwarding (preserves order) +### D3. 직렬 워커 포워딩 (순서 보존) -The component's worker is a single SimPy process that consumes flits -from `_inbox` and forwards them serially in arrival order. The -component does NOT spawn `env.process(...)` per flit. +본 컴포넌트의 워커는 `_inbox`에서 flit을 소비하여 도착 순서대로 직렬 +포워딩하는 단일 SimPy 프로세스이다. 컴포넌트는 flit마다 +`env.process(...)`를 spawn하지 **않는다**. -Rationale: if the first flit yields on `overhead_ns` while subsequent -flits run in parallel processes, the later flits can overtake the -first. This produces out-of-order delivery and lets the `is_last` -flit arrive at the destination before the first flit — corrupting -both the transaction's completion semantics and any flit-index-based -processing downstream. +근거: 첫 flit이 `overhead_ns`에서 yield하는 동안 후속 flit이 병렬 +프로세스에서 실행되면, 후속 flit이 첫 flit을 추월할 수 있다. 이는 순서가 +어긋난 전달을 낳고, `is_last` flit이 첫 flit보다 먼저 목적지에 도착하게 +하여 — 트랜잭션의 완료 의미와 다운스트림의 flit 인덱스 기반 처리 모두를 +손상시킨다. -### D4. Path-based next-hop routing +### D4. 경로 기반 next-hop 라우팅 -Routing is **not** a Forwarding-component concern. The Transaction -arrives with a pre-computed `path` (built by the router; ADR-0002 / -ADR-0017 D2). The component just looks up its own position in the -path and forwards to `path[index + 1]`: +라우팅은 Forwarding 컴포넌트의 관심사가 **아니다**. Transaction은 라우터에 +의해 사전 계산된 `path`(ADR-0002 / ADR-0017 D2)와 함께 도착한다. +컴포넌트는 단지 자신의 경로상 위치를 찾아 `path[index + 1]`로 전달한다: ```python def _next_hop_in_path(self, txn): @@ -97,104 +89,97 @@ def _next_hop_in_path(self, txn): return None ``` -If `next_hop` is found and present in `out_ports`, the flit is -forwarded. Otherwise (terminal node), `txn.done.succeed()` is -invoked when the `is_last` flit arrives. +`next_hop`이 발견되고 `out_ports`에 존재하면 flit이 전달된다. 그렇지 +않으면(종단 노드) `is_last` flit이 도착할 때 `txn.done.succeed()`가 +호출된다. -### D5. Flit-aware mode with Non-Flit fallback +### D5. Flit 인지 모드와 Non-Flit 폴백 -`_FLIT_AWARE = True` opts this component out of the base class's -flit-reassembly logic in `_fan_in`. Flits are placed directly on -`_inbox` (no reassembly), enabling per-flit handling in the worker -loop (D2, D3). +`_FLIT_AWARE = True`는 본 컴포넌트가 베이스 클래스의 `_fan_in` 내 flit +재조립 로직에서 제외되도록 한다. Flit은 재조립 없이 `_inbox`에 직접 +놓이며, 이는 워커 루프(D2, D3)에서의 per-flit 처리를 가능케 한다. -Non-Flit messages — zero-byte control Transactions and other -non-chunkified payloads — fall through to the base class's legacy -`_forward_txn` path via `env.process`. This preserves backward -compatibility for control-plane traffic that does not benefit from -flit-level processing. +Non-Flit 메시지 — 0바이트 제어 Transaction이나 그 외 청크화되지 않는 +페이로드 — 는 `env.process`를 통해 베이스 클래스의 레거시 `_forward_txn` +경로로 빠진다. 이는 flit 수준 처리의 이득이 없는 제어 평면 트래픽에 +대한 하위 호환성을 보존한다. -### D6. Multi-stream merging at the base class +### D6. 베이스 클래스에서의 멀티 스트림 병합 -Multi-stream FIFO merging at routers is the base class's -responsibility, not Forwarding's. The base class's `_fan_in` spawns -one process per `in_port`; all push to a single shared `_inbox`. -Flits from different upstream streams therefore interleave at -flit granularity in `_inbox`'s FIFO order. +라우터에서의 멀티 스트림 FIFO 병합은 Forwarding이 아닌 베이스 클래스의 +책임이다. 베이스 클래스의 `_fan_in`은 `in_port`마다 하나의 프로세스를 +spawn한다; 모두가 공유된 단일 `_inbox`에 push한다. 따라서 서로 다른 +업스트림 스트림의 flit들은 `_inbox`의 FIFO 순서로 flit 단위에서 +인터리브된다. -The Forwarding worker simply consumes `_inbox` in arrival order — -correctly modeling per-router multi-flow arbitration as -fair-FIFO over the shared inbox. +Forwarding 워커는 단지 `_inbox`를 도착 순서대로 소비할 뿐이다 — +공유 inbox 위의 공정 FIFO로 라우터별 멀티 플로우 중재를 올바르게 +모델링한다. -### D7. Single implementation under multiple impl names +### D7. 여러 impl 이름 아래의 단일 구현 -A single `TransitComponent` class is registered under four impl names -in `components.yaml`: +단일 `TransitComponent` 클래스가 `components.yaml`에서 네 가지 impl +이름으로 등록된다: -- `builtin.forwarding` — generic forwarding (e.g., `io_noc`, - `noc_router`, UCIe conn bridges) -- `builtin.switch` — tray-level switch -- `builtin.noc` — cube-level NOC fabric (legacy singleton; current - NOC routers use `builtin.forwarding`) -- `builtin.ucie` — UCIe protocol endpoint +- `builtin.forwarding` — 범용 forwarding (예: `io_noc`, `noc_router`, + UCIe conn 브리지) +- `builtin.switch` — 트레이 수준 스위치 +- `builtin.noc` — 큐브 수준 NOC 패브릭(레거시 싱글톤; 현재 NOC + 라우터는 `builtin.forwarding`을 사용) +- `builtin.ucie` — UCIe 프로토콜 엔드포인트 -All four aliases instantiate the same class with the same behavior. -Per-instance differentiation lives only in `attrs.overhead_ns`. -Separate impl names exist as intent tags for readability and to -allow future divergence without backward-incompatible config -changes. +네 별칭 모두 동일한 동작을 갖는 동일한 클래스를 인스턴스화한다. +인스턴스별 차별화는 `attrs.overhead_ns`에만 존재한다. 별도 impl 이름이 +존재하는 것은 가독성을 위한 의도 태그이자, 하위 호환을 깨지 않고 향후 +분기를 허용하기 위함이다. -### D8. Configurable `overhead_ns` +### D8. 설정 가능한 `overhead_ns` -A single attribute drives per-instance latency: +단일 속성이 인스턴스별 레이턴시를 결정한다: -| Usage site | impl name | overhead_ns | +| 사용 사이트 | impl 이름 | overhead_ns | | --- | --- | --- | -| Tray-level switch | `builtin.switch` | 5.0 | -| Cube NOC router | `builtin.forwarding` | 2.0 | -| IO chiplet io_noc | `builtin.forwarding` | 0.0 | -| UCIe protocol endpoint (`ucie-{N,S,E,W}`) | `builtin.ucie` | 8.0 | -| UCIe conn bridge (`ucie-{PORT}.conn{N}`) | `builtin.forwarding` | 0.0 | +| 트레이 수준 스위치 | `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 | -Default is 0.0. The attribute is read at each `run()` invocation, so -dynamic reconfiguration is possible but not currently used. +기본값은 0.0이다. 속성은 매 `run()` 호출에서 읽히므로 동적 재설정이 +가능하나 현재는 사용되지 않는다. ## Consequences ### Positive -- A single class handles all transit-node roles in the simulation - graph — minimal code surface for a high-population component type. -- Flit-aware processing + serial worker preserves wormhole semantics - across multi-hop paths without per-flit process overhead. -- `overhead_ns` is the only per-instance tunable; routing, BW, and - address resolution stay cleanly separated in their own components / - modules. -- Multi-stream merging emerges from the base-class structure; no - router-specific logic duplicates fair-FIFO arbitration. -- Non-Flit fallback path keeps control-plane traffic working without - forcing every message into the flit framework. +- 단일 클래스가 시뮬레이션 그래프의 모든 transit 노드 역할을 처리한다 + — 개체 수가 많은 컴포넌트 타입에 대한 최소 코드 표면. +- Flit 인지 처리 + 직렬 워커는 per-flit 프로세스 오버헤드 없이 멀티 홉 + 경로 전반에 걸쳐 웜홀 의미를 보존한다. +- `overhead_ns`만이 유일한 인스턴스별 튜너블이다; 라우팅, 대역폭, 주소 + 해석은 자체 컴포넌트/모듈에서 깨끗이 분리되어 있다. +- 멀티 스트림 병합이 베이스 클래스 구조에서 자연스럽게 발현된다; 라우터 + 전용 로직이 공정 FIFO 중재를 중복 구현하지 않는다. +- Non-Flit 폴백 경로는 모든 메시지를 flit 프레임워크로 강제하지 않고도 + 제어 평면 트래픽이 계속 동작하도록 한다. ### Negative -- The single class hides usage-site intent inside `attrs.overhead_ns` - configuration; readers must consult `topology.yaml` + - `components.yaml` to see which impl name maps to which behavior - class. -- Per-flit serial worker is a bottleneck if `overhead_ns` is large - and many concurrent transactions arrive at the same router; current - values (0–8 ns) make this negligible. +- 단일 클래스가 사용 사이트의 의도를 `attrs.overhead_ns` 설정 안에 + 숨긴다; 어떤 impl 이름이 어떤 동작 클래스로 매핑되는지 보려면 독자가 + `topology.yaml` + `components.yaml`을 참조해야 한다. +- per-flit 직렬 워커는 `overhead_ns`가 크고 같은 라우터에 다수의 동시 + 트랜잭션이 도착할 때 병목이 된다; 현재 값(0–8 ns)에서는 무시할 만한 + 수준이다. ## Links -- ADR-0002 (Routing distance — path computation) -- ADR-0015 D1 (Component port model) -- ADR-0015 D2 (Wire process — BW + propagation, separate from this - component) -- ADR-0015 D6 (Transit cube forwarding pattern) -- ADR-0016 D1 (IO chiplet io_noc — uses this component) -- ADR-0017 D1 (Cube NOC routers — use this component) -- ADR-0017 D6 (UCIe decomposition — `ucie-{PORT}` instances use this - component) -- ADR-0033 D1 (Flit-aware pass-through, first-flit overhead, - multi-stream merge semantics) +- 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 오버헤드, 멀티 스트림 병합 + 의미) diff --git a/docs/adr-ko/ADR-0013-ver-verification_strategy.md b/docs/adr/ADR-0013-ver-verification-strategy.md similarity index 100% rename from docs/adr-ko/ADR-0013-ver-verification_strategy.md rename to docs/adr/ADR-0013-ver-verification-strategy.md diff --git a/docs/adr/ADR-0013-ver-verification_strategy.md b/docs/adr/ADR-0013-ver-verification_strategy.md deleted file mode 100644 index dd7e958..0000000 --- a/docs/adr/ADR-0013-ver-verification_strategy.md +++ /dev/null @@ -1,139 +0,0 @@ -# ADR-0013: Verification Strategy and Phase 1 Test Plan - -## Status - -Accepted - -## Context - -KernBench is a system-level simulator whose correctness is defined by: - -- adherence to SPEC-defined invariants, -- determinism and debuggability, -- explicit modeling of routing and latency. - -Given the evolving implementation, we need a stable verification strategy -that prevents architectural drift while allowing incremental development. - -This ADR defines the Phase 1 verification plan and what constitutes -"correct behavior" for early implementations. - ---- - -## Decision - -### D1. Verification is contract-based - -Verification MUST be derived from: - -- SPEC requirements, -- accepted ADRs. - -Tests MUST validate architectural contracts, not incidental implementation details. - ---- - -### D2. Phase 1 verification scope - -Phase 1 verification focuses on: - -- message contract validity (ADR-0012), -- routing and fan-out semantics at the IO_CPU boundary (ADR-0009), -- PA-first memory addressing and shard tagging (ADR-0011), -- core latency and trace invariants (SPEC 0.1, R2). - -Microarchitectural accuracy, bandwidth contention, and cycle-level behavior -are explicitly out of scope in Phase 1. - ---- - -### D3. Required Phase 1 verification cases - -The following verification cases MUST be supported by the implementation: - -#### V1. Message schema validation - -- KernelLaunch requests missing `(sip, cube, pe)` in any tensor shard MUST be rejected. -- MemoryWrite/MemoryRead requests missing destination/source placement tags MUST be rejected. -- Completion results MUST follow the `ok / error_code / error_message` contract. - -#### V2. IO_CPU fan-out and aggregation - -Given: - -- a topology with one SIP, one CUBE, and two PEs, -- a KernelLaunch request containing two tensor shards targeting different PEs, - -The system MUST: - -- submit a single KernelLaunch to IO_CPU, -- fan-out work internally to both PEs, -- aggregate completion and return a single deterministic completion to the host. - -#### V3. Latency and trace invariants - -For any valid request: - -- the hop-by-hop trace MUST be non-empty, -- total latency MUST be greater than zero, -- repeated runs with identical inputs MUST produce identical traces. - -#### V4. Topology independence and cross-domain coverage - -Verification cases MUST pass for multiple topology shapes, including: - -- minimal: (1 SIP, 1 CUBE, 1 PE) -- multi-PE: (1 SIP, 1 CUBE, N PEs) -- multi-CUBE within a SIP: (1 SIP, M CUBEs, ≥1 PE per CUBE) -- multi-SIP tray: (K SIPs, ≥1 CUBE per SIP, ≥1 PE per CUBE) - -For multi-CUBE and multi-SIP topologies, Phase 1 verification focuses on: - -- explicit connectivity (required links exist), -- deterministic routing and control-path traversal, -- non-empty traces and latency > 0 for representative cross-domain requests - (inter-CUBE and inter-SIP paths). - -Tests MUST NOT hardcode topology sizes, node ids, or link counts. -Instead, tests MUST derive expectations from the compiled topology metadata ---- - -### D4. Phase 1 artifacts - -Phase 1 MAY include: - -- verification-only test code, -- topology fixtures, -- trace inspection utilities. - -Phase 1 MUST NOT require: - -- production code changes solely to satisfy tests, -- weakening or removing tests to allow progress. - ---- - -### D5. Phase 2 enforcement - -Phase 2 (Apply) MUST: - -- run the Phase 1 verification cases, -- rollback all changes if any verification fails, -- preserve tests as authoritative contracts. - ---- - -## Consequences - -- Architectural correctness is enforced early. -- Tests serve as executable documentation of system behavior. -- Implementation remains flexible without losing rigor. - ---- - -## Links - -- SPEC 0.1, R2, R6 -- ADR-0011 (Memory Addressing — PA / VA / LA) -- ADR-0012 (Host ↔ IO_CPU message schema) -- ADR-0009 (Kernel execution semantics) diff --git a/tests/test_verify_adr_lang_pairs.py b/tests/test_verify_adr_lang_pairs.py index 0ad2881..1ec7876 100644 --- a/tests/test_verify_adr_lang_pairs.py +++ b/tests/test_verify_adr_lang_pairs.py @@ -52,7 +52,35 @@ def test_status_mismatch_fails(tmp_path: Path) -> None: _make_adr(tmp_path / "docs/adr/ADR-0001-foo-bar.md", "0001", status="Accepted") _make_adr(tmp_path / "docs/adr-ko/ADR-0001-foo-bar.md", "0001", status="Proposed") errs = v.verify(tmp_path) - assert any("Status block mismatch" in e for e in errs) + assert any("Status keyword mismatch" in e for e in errs) + + +def test_status_parenthetical_translation_passes(tmp_path: Path) -> None: + """Parenthetical commentary in Status may be translated; only the + lifecycle keyword needs to match across EN/KO.""" + _make_adr( + tmp_path / "docs/adr/ADR-0001-foo-bar.md", "0001", + status="Accepted (Revision 2 - concrete bit layout)", + ) + _make_adr( + tmp_path / "docs/adr-ko/ADR-0001-foo-bar.md", "0001", + status="Accepted (Revision 2 - 비트 레이아웃 확정)", + ) + assert v.verify(tmp_path) == [] + + +def test_status_list_after_keyword_passes(tmp_path: Path) -> None: + """Trailing bullet list in Status may be translated; only first-line + keyword is compared.""" + _make_adr( + tmp_path / "docs/adr/ADR-0011-foo-bar.md", "0011", + status="Accepted.\n\n- VA model: implemented (default).", + ) + _make_adr( + tmp_path / "docs/adr-ko/ADR-0011-foo-bar.md", "0011", + status="Accepted.\n\n- VA 모델: 구현됨 (기본값).", + ) + assert v.verify(tmp_path) == [] def test_title_id_mismatch_fails(tmp_path: Path) -> None: @@ -77,6 +105,21 @@ def test_multiline_status_with_parenthetical_passes(tmp_path: Path) -> None: assert v.verify(tmp_path) == [] +def test_superseded_keyword_must_match(tmp_path: Path) -> None: + """Multi-word lifecycle keywords like 'Superseded by ADR-NNNN' must + match in full (the ADR ref is part of the keyword).""" + _make_adr( + tmp_path / "docs/adr/ADR-0001-foo-bar.md", "0001", + status="Superseded by ADR-0031", + ) + _make_adr( + tmp_path / "docs/adr-ko/ADR-0001-foo-bar.md", "0001", + status="Superseded by ADR-0099", + ) + errs = v.verify(tmp_path) + assert any("Status keyword mismatch" in e for e in errs) + + def test_crlf_normalization(tmp_path: Path) -> None: """KO has CRLF, EN has LF; Status content is otherwise identical -> pass.""" en = tmp_path / "docs/adr/ADR-0001-foo-bar.md" @@ -105,9 +148,10 @@ def test_em_dash_title_separator_recognized(tmp_path: Path) -> None: def test_underscore_in_slug_recognized(tmp_path: Path) -> None: - """ADR-0013 uses an underscore in its slug; the regex must accept it.""" - _make_adr(tmp_path / "docs/adr/ADR-0013-ver-verification_strategy.md", "0013") - _make_adr(tmp_path / "docs/adr-ko/ADR-0013-ver-verification_strategy.md", "0013") + """Defensive: slug regex must accept underscores even though current + ADRs use kebab-case (in case future ADRs reintroduce them).""" + _make_adr(tmp_path / "docs/adr/ADR-0099-cat-foo_bar.md", "0099") + _make_adr(tmp_path / "docs/adr-ko/ADR-0099-cat-foo_bar.md", "0099") assert v.verify(tmp_path) == [] diff --git a/tools/verify_adr_lang_pairs.py b/tools/verify_adr_lang_pairs.py index 5661cf3..a909173 100644 --- a/tools/verify_adr_lang_pairs.py +++ b/tools/verify_adr_lang_pairs.py @@ -10,8 +10,10 @@ Checks: - every docs/adr/.md has a matching docs/adr-ko/.md - every docs/adr-ko/.md has a matching docs/adr/.md (no orphans) - title line `# ADR-NNNN:` of each pair matches the filename's NNNN - - `## Status` block content is byte-equal (after CRLF/LF normalization) - between EN and KO + - `## Status` keyword (e.g., Accepted / Proposed / Superseded by ADR-NNNN) + matches between EN and KO. Parenthetical commentary and trailing + list items in the Status block may differ in wording (e.g., be + translated); only the leading lifecycle keyword is compared. Exit code: 0 if all OK, 1 if any mismatch. """ @@ -70,6 +72,24 @@ def extract_status_block(text: str) -> str | None: return "\n".join(collected).strip() +def normalize_status_keyword(block: str) -> str: + """Reduce a Status block to its lifecycle keyword for cross-language compare. + + Drops parenthetical commentary (which is allowed to be translated) + and trailing list items / extra lines, then strips trailing + punctuation. e.g.: + 'Accepted' -> 'Accepted' + 'Accepted.' -> 'Accepted' + 'Accepted (Revision 2 - foo)' -> 'Accepted' + 'Accepted (supersedes ADR-0029)' -> 'Accepted' + 'Accepted.\\n\\n- VA model: ...' -> 'Accepted' + 'Superseded by ADR-0031' -> 'Superseded by ADR-0031' + """ + no_parens = re.sub(r"\([^)]*\)", "", block).strip() + first_line = next((l for l in no_parens.splitlines() if l.strip()), "") + return first_line.strip().rstrip(".,;:").strip() + + def verify(root: Path) -> list[str]: errors: list[str] = [] en_dir = root / "docs" / "adr" @@ -110,12 +130,15 @@ def verify(root: Path) -> list[str]: errors.append(f"{name}: EN missing `## Status` section") if ko_status is None: errors.append(f"{name}: KO missing `## Status` section") - if en_status is not None and ko_status is not None and en_status != ko_status: - errors.append( - f"{name}: Status block mismatch\n" - f" EN: {en_status!r}\n" - f" KO: {ko_status!r}" - ) + if en_status is not None and ko_status is not None: + en_kw = normalize_status_keyword(en_status) + ko_kw = normalize_status_keyword(ko_status) + if en_kw != ko_kw: + errors.append( + f"{name}: Status keyword mismatch\n" + f" EN: {en_kw!r}\n" + f" KO: {ko_kw!r}" + ) return errors