diff --git a/docs/adr/ADR-0018-Logical Address.md b/docs/adr/ADR-0018-Logical Address.md new file mode 100644 index 0000000..c8325f4 --- /dev/null +++ b/docs/adr/ADR-0018-Logical Address.md @@ -0,0 +1,440 @@ +# ADR-0018: LA 기반 메모리 주소 추상화 및 HBM Channel Mapping Mode 도입 + +## Status + +Proposed + +## Context + +Kernbench는 CUBE 내부에서 PE_DMA와 Local-HBM 간의 메모리 접근을 시뮬레이션한다. +현재는 VA 기반 접근 경로를 사용하고 있으나, 다음 두 가지 channel mapping 모델을 +일관되게 표현하기 어렵다. + +### 배경: Local-HBM pseudo channel 구조 + +CUBE의 HBM은 32개 또는 64개의 pseudo channel로 구성된다. +PE-Local-HBM 모델에서는 각 PE가 동일한 수의 pseudo channel을 담당한다. + +예: 64 pseudo channel, 8 PE per cube → 각 PE가 8개 pseudo channel을 local HBM으로 접근 + +pseudo channel 수와 PE 수는 모두 topology 파라미터이다. +`N = hbm_pseudo_channels / pes_per_cube` (= channels_per_pe)가 +PE당 local channel 수를 결정한다. + +각 pseudo channel의 BW(예: 32 GB/s)만큼 DMA와 pseudo channel 사이의 라우팅 경로 BW도 +맞춰지므로, PE가 N개 채널에 동시 request를 보내면 최대 메모리 BW를 활용할 수 있다. + +### 현재 VA 모델의 한계 + +채널을 8개로 나누면 request도 채널별로 생성되어 DMA에 보내져야 한다. +그러나 현재 구조에서는 커널이 VA를 가지고 request를 생성한 뒤(`tl.load`) +DMA에 바로 전달하므로, PE_CPU가 채널별 DMA request를 생성하기 어렵다. + +따라서 VA 대신 **Logical Address(LA)** 를 사용하고, +PE_DMA 내부의 **BAAW(Logical-to-Physical Mapping Unit)** 가 +segment-based mapping을 기반으로 LA → PA 또는 PA 리스트로 변환하는 구조를 제안한다. + +### 두 가지 channel mapping mode + +- **1:1 mode**: 채널별 request를 만들어 실행. 정밀한 per-channel 모델링 +- **n:1 mode (default)**: local HBM 채널 간 인터리빙 가정. aggregated BW 모델링 + +두 모드를 지원하여 n:1 모드의 오버헤드를 측정/검토할 수 있게 한다. + +### 핵심 요구사항 + +- PE_DMA → HBM_CTRL의 effective bandwidth semantics는 두 모드에서 동일해야 한다 +- 차이는 request 표현 방식과 resource 모델링 방식에만 있어야 한다 +- kernel programming model은 변경하지 않는다 +- physical channel 정보는 kernel에 노출되지 않아야 한다 + +### 기존 Physical Address + +현재 시스템의 51-bit Physical Address는 `policy/address/phyaddr.py`에 정의되어 있다: + +``` +[50:47] rack_id (4 bit) +[46:43] sip_id (4 bit) +[42:38] cube_id (5 bit, sip_seg) +[37] hbm_selector (1=HBM window) +[36:0] hbm_offset (37 bit, 128GB per cube) +``` + +PA는 최종 라우팅 가능한 canonical physical destination을 표현하는 데 사용되며, +이 역할은 유지된다. +하지만 logical access → physical request 변환 시점과 정책이 명확히 분리되어 있지 않다. + +--- + +## Decision + +### D1. LA (Logical Address) 도입 — VA를 대체 + +기존 VA(Virtual Address) 인프라를 LA(Logical Address)로 대체한다. + +#### LA의 특징 + +- VA처럼 Tensor를 연속적인 메모리 공간에 매핑할 수 있다 +- logical buffer + offset을 표현 +- physical channel 정보를 직접 포함하지 않음 +- physical resolution 이전까지 유지되는 중간 추상화 +- 커널 코드(`tl.load`, `tl.store`, `tl.composite`)가 사용하는 유일한 주소 체계 + +#### LA 공간 정의 + +| 항목 | 값 | +|------|-----| +| LA 시작 주소 | `0x1_0000_0000` (4 GB, 기존 VA 시작점 유지) | +| LA 공간 크기 | PE당 64 GB | +| 정렬 단위 | segment 단위 (아래 D3 참조) | + +LA는 PE-local 주소 공간이다. +서로 다른 PE가 동일한 LA 값을 사용해도 BAAW의 segment table이 다르므로 +서로 다른 PA로 resolve된다. + +#### VA 인프라 제거 범위 + +LA 도입에 따라 다음 기존 코드를 대체/제거한다: + +| 제거 대상 | 대체 | +|-----------|------| +| `policy/address/va_allocator.py` (VirtualAllocator) | LA allocator (동일 free-list 방식, 이름/역할 변경) | +| `policy/address/pe_mmu.py` (PeMMU) | BAAW segment table (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 mapping install | LA alloc + BAAW segment install | +| `runtime_api/tensor.py`: `va_base` 필드 | `la_base` 필드 | +| `topology.yaml`: pe_mmu 컴포넌트 항목 | 제거 | + +--- + +### D2. Mapping Mode 설정 + +topology.yaml의 cube 레벨에서 mapping mode를 설정한다: + +```yaml +cube: + memory_map: + hbm_mapping_mode: n_to_one # one_to_one | n_to_one + hbm_pseudo_channels: 64 # 전체 pseudo channel 수 + hbm_channels_per_pe: 8 # PE당 local channel 수 + hbm_channel_bw_gbs: 32.0 # per-channel bandwidth +``` + +이 설정은 graph compiler(topology builder)와 BAAW 초기화 시 참조된다. + +--- + +### D3. Segment 및 BAAW + +#### Segment 정의 + +Segment는 LA space를 partition하여, 각 segment가 특정 HBM channel 또는 +channel group에 매핑되도록 하는 logical allocation 단위이다. + +Segment는 runtime allocator가 tensor deploy 시 생성하며, +BAAW는 이를 기반으로 LA를 physical request로 변환한다. + +#### BAAW Segment Table Entry + +```python +@dataclass +class BaawSegment: + 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 # 이 segment에 할당된 channel 수 (e.g., 8) + pa_bases: list[int] # per-channel PA 시작 주소 리스트 (len = channel_count) + channel_ids: list[int] # per-channel 논리적 ID (e.g., [0,1,2,...,7]) + channel_size: int # per-channel 크기 (la_size // channel_count) + # n:1 mode fields + agg_pa_base: int # aggregated PA 시작 주소 + agg_node_id: str # aggregated router node_id (for routing) +``` + +#### Segment 라이프사이클 + +1. **할당 시점** (tensor deploy): + - RuntimeContext가 LA allocator에서 LA 공간 할당 + - PEMemAllocator가 per-channel PA 할당 (1:1) 또는 aggregated PA 할당 (n:1) + - `BaawSegmentInstallMsg`를 PE_DMA로 전송하여 segment table에 등록 + +2. **사용 시점** (kernel 실행): + - 커널이 `tl.load(la_ptr)` → DmaReadCmd(src_addr=LA) + - PE_DMA가 BAAW에서 LA에 해당하는 segment를 lookup + - mode에 따라 PA(들)로 변환 + +3. **해제 시점** (tensor free): + - segment table에서 제거 + - LA 공간 반환, PA 해제 + +--- + +### D4. BAAW (Logical-to-Physical Mapping Unit) + +#### 위치 + +BAAW는 PE_DMA 내부의 front-end stage로 배치된다. +별도의 SimPy 컴포넌트가 아니며, PE_DMA의 `handle_command()` 시작 부분에서 실행되는 +동기적 address resolution 로직이다. + +#### 입력 + +- LA (Logical Address) — DmaReadCmd.src_addr 또는 DmaWriteCmd.dst_addr +- access size (bytes) + +#### 출력 + +- 1:1 mode: `list[PhysicalRequest]` — 각 request는 (PA, nbytes, channel_node_id) +- n:1 mode: `PhysicalRequest` 1개 — (agg_PA, nbytes, agg_node_id) + +```python +@dataclass +class PhysicalRequest: + pa: int # 51-bit Physical Address + nbytes: int # 이 request의 transfer size + dst_node: str # target node_id (channel router or aggregated router) +``` + +#### BAAW Resolve 로직 + +```python +def resolve(self, la: int, nbytes: int) -> list[PhysicalRequest]: + seg = self._find_segment(la) # la_base <= la < la_base + la_size + offset = la - seg.la_base + + if seg.mode == "n_to_one": + pa = seg.agg_pa_base + offset + return [PhysicalRequest(pa=pa, nbytes=nbytes, dst_node=seg.agg_node_id)] + + elif seg.mode == "one_to_one": + requests = [] + per_ch_size = seg.channel_size + for i, (pa_base, ch_id) in enumerate(zip(seg.pa_bases, seg.channel_ids)): + ch_offset = offset % per_ch_size # interleaved or striped + ch_nbytes = nbytes // seg.channel_count + pa = pa_base + ch_offset + dst_node = f"{self._pe_prefix}.ch_r{ch_id}" + requests.append(PhysicalRequest(pa=pa, nbytes=ch_nbytes, dst_node=dst_node)) + return requests +``` + +#### 역할 범위 + +BAAW의 책임: +- logical access를 physical request 단위로 변환 +- mapping mode에 따른 fan-out (1:1) 또는 pass-through (n:1) 수행 +- Physical Address 생성 및 target node 결정 + +BAAW의 책임이 아닌 것: +- 실제 data movement 수행 +- NOC routing 실행 +- bandwidth 소비 시뮬레이션 (downstream component의 역할) + +#### Output Contract + +BAAW의 출력은 추가적인 address decoding 없이 +simulator의 routing 및 resource 모델에서 직접 사용 가능한 request 단위여야 한다. + +--- + +### D5. PE_DMA handle_command() 변경 + +#### 현재 흐름 (VA 기반) + +``` +DmaReadCmd.src_addr (VA) + → MMU.translate(VA) → PA + → PhysAddr.decode(PA) → PhysAddr object + → resolver.resolve(PhysAddr) → dst_node_id (e.g., "sip0.cube0.hbm_ctrl.slice3") + → router.find_path(pe_prefix, dst_node_id) → path + → 1개 sub-Transaction 생성 → fabric inject +``` + +#### 새 흐름 (LA 기반) + +``` +DmaReadCmd.src_addr (LA) + → BAAW.resolve(LA, nbytes) → list[PhysicalRequest] + → 각 PhysicalRequest에 대해: + → router.find_path(pe_prefix, req.dst_node) → path + → compute_drain_ns(path, req.nbytes) → drain + → sub-Transaction 생성 → fabric inject + → 모든 sub-Transaction 완료 대기 + → pe_txn.done.succeed() +``` + +핵심 변경: +- MMU 참조 제거 → BAAW resolve로 대체 +- PhysAddr.decode() + resolver.resolve() → BAAW가 직접 dst_node 반환 +- 1개 request → N개 request 병렬 inject (1:1 mode) + +--- + +### D6. 1:1 Mode 상세 + +- 하나의 logical access → N개(= `channels_per_pe`)의 physical request +- N은 `hbm_pseudo_channels / pes_per_cube`로 결정되는 파라미터 +- 각 request: + - fully resolved 51-bit PA + - 특정 channel router를 target (`{pe_prefix}.ch_r{channel_id}`) +- per-channel link에 의한 BW contention 모델링 +- PE_DMA는 N개 sub-transaction을 동시에 inject + +#### 1:1 Mode 예시 + +구성: `hbm_pseudo_channels=64`, `pes_per_cube=8` +→ `channels_per_pe=8`, PE0이 ch0-7 소유 + +```text +Tensor A (4 KB) → LA 0x1_0000_0000, size=4096 bytes +BAAW segment: { + la_base: 0x1_0000_0000, la_size: 4096, + mode: "one_to_one", channel_count: 8, # = channels_per_pe + pa_bases: [PA_ch0, PA_ch1, ..., PA_ch7], + channel_ids: [0, 1, 2, 3, 4, 5, 6, 7], + channel_size: 512, # = la_size / channel_count +} + +BAAW resolve 결과 (N=8개 request): + → 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: N개 sub-transaction 병렬 inject + 각각 channel router → hbm_ctrl link (channel_bw_gbs)를 통해 HBM 접근 + 총 effective BW = N × channel_bw_gbs +``` + +N이 다른 구성의 예: +- `hbm_pseudo_channels=32`, `pes_per_cube=8` → `channels_per_pe=4`, 4개 request +- `hbm_pseudo_channels=64`, `pes_per_cube=4` → `channels_per_pe=16`, 16개 request + +--- + +### D7. n:1 Mode 상세 + +- 하나의 logical access → 하나의 aggregated request +- target: aggregated router → hbm_ctrl (ADR-0019 참조) +- aggregated link BW = `channels_per_pe` × `channel_bw_gbs` (e.g., 8 × 32 = 256 GB/s) +- single queue / resource로 모델링 +- per-channel PA 분해 없음 + +#### n:1 Mode 예시 + +``` +Tensor A (4 KB) → LA 0x1_0000_0000, size=4096 bytes +BAAW segment: { + la_base: 0x1_0000_0000, la_size: 4096, + mode: "n_to_one", + agg_pa_base: PA_agg, + agg_node_id: "sip0.cube0.pe0.agg_router", +} + +BAAW resolve 결과: + → PhysicalRequest(pa=PA_agg, nbytes=4096, dst_node="sip0.cube0.pe0.agg_router") + +PE_DMA: 1개 sub-transaction inject + aggregated router → hbm_ctrl link (256 GB/s)를 통해 HBM 접근 +``` + +--- + +### D8. Kernel Model 유지 + +- kernel은 여전히 단일 memory op만 발행 (`tl.load`, `tl.store`, `tl.composite`) +- LA가 커널에 전달되는 주소 체계 +- channel 분해/집계는 PE_DMA 내부 BAAW에서 수행 +- kernel 코드에 physical channel 정보가 노출되지 않음 + +--- + +## Consequences + +### Positive + +- 1:1 vs n:1 semantics가 BAAW라는 단일 지점에서 명확히 분리됨 +- kernel abstraction 유지 — 커널 코드 변경 불필요 +- topology 기반 정책 제어 가능 (yaml에서 mode 전환) +- simulation 모델 일관성 및 디버깅 용이성 향상 +- segment-based mapping은 page table 대비 단순하고 overhead가 낮음 + +### Negative + +- VA/MMU 기반 코드 전체 리팩토링 필요 +- request 생성 경로 복잡도 증가 (1:1 mode에서 N개 request 관리) +- n:1 mode에서 per-channel visibility 감소 +- 기존 VA 관련 테스트 재작성 필요 + +--- + +## Alternatives + +### A1. VA 유지 + MMU에서 fan-out + +- MMU가 per-channel PA를 반환하도록 확장 +- 문제: MMU의 역할이 address translation을 넘어 request 분해까지 확장됨 +- 문제: n:1 mode에서 aggregation 표현이 어려움 + +### A2. Kernel이 channel-aware request 생성 + +- 커널이 직접 채널별 load/store를 호출 +- 문제: abstraction leakage, portability 저하 +- 문제: 모든 벤치마크 코드 수정 필요 + +### A3. 항상 PA 사용 (LA 없이) + +- runtime이 직접 per-channel PA를 커널에 전달 +- 문제: aggregation 모델과 충돌 +- 문제: 변환 시점이 불명확, 커널에 channel 정보 노출 + +--- + +## Implementation Notes + +### 구현 순서 + +1. LA 타입 도입 (`policy/address/la_allocator.py`) +2. BAAW segment table 구현 (`policy/address/baaw.py`) +3. `BaawSegmentInstallMsg` 메시지 타입 추가 (`runtime_api/kernel.py`) +4. PE_DMA에 BAAW 통합 (`components/builtin/pe_dma.py` handle_command 변경) +5. RuntimeContext 변경: LA alloc + segment install (`runtime_api/context.py`) +6. Tensor.va_base → la_base 변경 (`runtime_api/tensor.py`) +7. VA/MMU 코드 제거 +8. topology.yaml에서 pe_mmu 제거, mapping mode 설정 추가 +9. 테스트 마이그레이션 + +### 영향받는 기존 테스트 + +| 테스트 파일 | 영향 | +|------------|------| +| `tests/test_mmu_component.py` | 제거 → BAAW segment install 테스트로 대체 | +| `tests/test_mmu_fabric.py` | 제거 → BAAW + fabric 통합 테스트로 대체 | +| `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 테스트로 대체 | + +--- + +## Test Requirements + +- 동일 logical access에 대해: + - 1:1 → N개 request 생성 확인 + - n:1 → 1개 aggregated request 생성 확인 +- 두 모드에서 effective bandwidth 일관성 검증 +- 1:1 → per-channel contention 모델링 확인 +- n:1 → aggregated bandwidth 반영 확인 +- kernel 코드 변경 없이 동작 확인 +- BAAW segment install/uninstall 정상 동작 +- 여러 tensor가 서로 다른 segment에 할당될 때 충돌 없음 + +--- + +## Links + +- ADR-0011 (Memory Addressing Simplification — PA-first, VA/MMU 도입) → 본 ADR이 대체 +- ADR-0019 (NOC Per-Channel HBM 연결 모델) → topology 측 연동 +- ADR-0014 (PE Internal Execution Model) → PE_DMA 변경 영향 diff --git a/docs/adr/ADR-0019-NOC-Local HBM.md b/docs/adr/ADR-0019-NOC-Local HBM.md new file mode 100644 index 0000000..238a618 --- /dev/null +++ b/docs/adr/ADR-0019-NOC-Local HBM.md @@ -0,0 +1,513 @@ +# ADR-0019: CUBE NOC 내 Per-Channel 및 Aggregated HBM 연결 모델 + +## Status + +Proposed + +## Context + +ADR-0018에서는 LA 기반 주소 추상화와 BAAW를 도입하여, +logical memory access가 다음 두 형태의 request로 변환되도록 정의하였다. + +- 1:1 mode: 하나의 logical access → N개의 per-channel request +- n:1 mode: 하나의 logical access → 하나의 aggregated request + +여기서 N = `hbm_pseudo_channels / pes_per_cube` (= `channels_per_pe`)이며, +topology 파라미터로 결정된다. + +### 기존 구조의 문제 + +현재 구현(`topology/builder.py`)에서는: + +- PE_DMA → NOC → xbar_top/xbar_bot → HBM_CTRL.slice{0-7} 경로를 사용 +- HBM은 8개 slice(= PE 수) 노드로 모델링됨 +- local/remote access가 서로 다른 경로를 사용: + - local: NOC → xbar → HBM slice + - cross-half: NOC → xbar_top → bridge → xbar_bot → HBM slice + - remote cube: NOC → UCIe → remote NOC → remote xbar → remote HBM slice + +이 구조의 한계: + +- pseudo-channel 단위 모델링 불가 (slice = PE 단위, channel 단위 아님) +- xbar/bridge가 local/remote 경로를 이원화 +- 1:1 / n:1 mode를 일관되게 표현할 수 없음 + +--- + +## Decision + +### D1. HBM controller는 CUBE당 단일 endpoint로 정의한다 + +현재의 `hbm_ctrl.slice{0-7}` (8개 노드)를 **`hbm_ctrl` 단일 노드**로 통합한다. + +- pseudo channel은 HBM controller 노드 자체가 아니라, + controller에 연결되는 **link의 단위**로 표현한다 +- HBM controller 내부의 read/write resource 모델은 유지하되, + mode에 따라 contention 단위가 달라진다: + - 1:1 mode: per-channel link가 BW contention point (controller는 terminal) + - n:1 mode: aggregated link가 BW contention point (controller는 terminal) + +노드 네이밍 변경: + +| 현재 | 변경 후 | +| ---- | ------- | +| `sip0.cube0.hbm_ctrl.slice0` ~ `slice7` | `sip0.cube0.hbm_ctrl` (단일) | + +--- + +### D2. xbar, bridge 완전 제거 + +기존 다음 노드 및 관련 edge를 모두 제거한다: + +- `{cube}.xbar_top`, `{cube}.xbar_bot` +- `{cube}.bridge.left`, `{cube}.bridge.right` +- `noc_to_xbar`, `xbar_to_noc`, `xbar_to_hbm`, `hbm_to_xbar` 종류의 edge +- `xbar_to_bridge`, `bridge_to_xbar` 종류의 edge + +이들의 역할(PE→HBM 라우팅, cross-half 연결)은 +channel router 및 horizontal line 연결이 대체한다 (D3, D4 참조). + +--- + +### D3. 1:1 mode: per-channel router 기반 연결 + +#### channel router 정의 + +1:1 mode에서 graph compiler는 pseudo-channel 수만큼의 **channel router** 노드를 +생성한다. channel router는 NOC의 일부이다. + +```text +파라미터 예: hbm_pseudo_channels=64, pes_per_cube=8 +→ channels_per_pe = 8, 총 64개 channel router 생성 +``` + +노드 네이밍: `{cube}.ch_r{global_channel_id}` + +| PE | 소유 channel routers | +| -- | -------------------- | +| PE0 | ch_r0, ch_r1, ..., ch_r7 | +| PE1 | ch_r8, ch_r9, ..., ch_r15 | +| ... | ... | +| PE7 | ch_r56, ch_r57, ..., ch_r63 | + +일반화: PE `p`는 channel `p * channels_per_pe` ~ `(p+1) * channels_per_pe - 1`을 소유. + +#### PE_DMA ↔ channel router 연결 + +각 PE_DMA는 자신의 local channel router N개와 양방향 link로 연결된다: + +```text +sip0.cube0.pe0.pe_dma ←→ sip0.cube0.ch_r0 (bw: channel_bw_gbs) +sip0.cube0.pe0.pe_dma ←→ sip0.cube0.ch_r1 (bw: channel_bw_gbs) +... +sip0.cube0.pe0.pe_dma ←→ sip0.cube0.ch_r7 (bw: channel_bw_gbs) +``` + +- edge kind: `pe_to_ch_router` / `ch_router_to_pe` +- BW: `hbm_channel_bw_gbs` (e.g., 32 GB/s) +- distance: PE에서 channel router까지의 물리적 거리 (layout 기반) + +#### channel router ↔ HBM controller 연결 + +각 channel router는 cube의 hbm_ctrl과 양방향 link로 연결된다: + +```text +sip0.cube0.ch_r0 ←→ sip0.cube0.hbm_ctrl (bw: channel_bw_gbs) +sip0.cube0.ch_r1 ←→ sip0.cube0.hbm_ctrl (bw: channel_bw_gbs) +... +sip0.cube0.ch_r63 ←→ sip0.cube0.hbm_ctrl (bw: channel_bw_gbs) +``` + +- edge kind: `ch_router_to_hbm` / `hbm_to_ch_router` +- BW: `hbm_channel_bw_gbs` (e.g., 32 GB/s) + +#### 1:1 mode 전체 데이터 경로 + +```text +PE0.pe_dma + ├→ ch_r0 → hbm_ctrl (32 GB/s) + ├→ ch_r1 → hbm_ctrl (32 GB/s) + ├→ ... + └→ ch_r7 → hbm_ctrl (32 GB/s) + 총 PE0 local BW = N × channel_bw_gbs +``` + +--- + +### D4. 1:1 mode: horizontal line 연결 (cross-PE channel 접근) + +#### 배치 규칙 + +같은 **logical index**를 가지는 channel router들을 동일한 horizontal row에 배치한다. + +logical index 정의: `logical_idx = global_channel_id % channels_per_pe` + +```text +파라미터 예: channels_per_pe=8, pes_per_cube=8 + +Row 0: ch_r0 (PE0) ↔ ch_r8 (PE1) ↔ ch_r16 (PE2) ↔ ... ↔ ch_r56 (PE7) +Row 1: ch_r1 (PE0) ↔ ch_r9 (PE1) ↔ ch_r17 (PE2) ↔ ... ↔ ch_r57 (PE7) +Row 2: ch_r2 (PE0) ↔ ch_r10 (PE1) ↔ ch_r18 (PE2) ↔ ... ↔ ch_r58 (PE7) +... +Row 7: ch_r7 (PE0) ↔ ch_r15 (PE1) ↔ ch_r23 (PE2) ↔ ... ↔ ch_r63 (PE7) +``` + +일반화: Row `r`에는 `{ch_r(p * N + r) | p ∈ 0..pes_per_cube-1}`이 위치. +여기서 `N = channels_per_pe`. + +#### horizontal line edge + +같은 row에서 인접한 channel router끼리 양방향 edge로 연결: + +```text +ch_r0 ↔ ch_r8 ↔ ch_r16 ↔ ... ↔ ch_r56 +``` + +- edge kind: `ch_horizontal` +- BW: `hbm_channel_bw_gbs` (or configurable inter-PE channel BW) +- distance: PE 간 물리적 거리 + +#### cross-PE HBM 접근 경로 (1:1 mode) + +PE0이 PE1의 local channel (ch_r8)에 접근하는 경우: + +```text +PE0.pe_dma → ch_r0 → ch_r8 (horizontal hop) → hbm_ctrl +``` + +Dijkstra router가 horizontal line을 통해 최단 경로를 탐색한다. + +#### 설계 의도 + +이 배치 규칙은: + +- routing 규칙 단순화: horizontal = cross-PE, vertical = PE-local +- 거리 계산 단순화: row 내 hop 수 = |src_pe - dst_pe| +- 구조적 반복성 확보: 모든 row가 동일한 구조 + +--- + +### D5. n:1 mode: aggregated router 기반 연결 + +#### aggregated router 정의 + +n:1 mode에서 graph compiler는 PE당 1개의 **aggregated router** 노드를 생성한다. +aggregated router는 NOC의 일부이다. + +노드 네이밍: `{cube}.pe{p}.agg_router` + +#### 연결 구조 + +```text +sip0.cube0.pe0.pe_dma ←→ sip0.cube0.pe0.agg_router (bw: N × channel_bw_gbs) +sip0.cube0.pe0.agg_router ←→ sip0.cube0.hbm_ctrl (bw: N × channel_bw_gbs) +``` + +- edge kind: `pe_to_agg_router` / `agg_router_to_pe`, `agg_to_hbm` / `hbm_to_agg` +- BW: `channels_per_pe × hbm_channel_bw_gbs` (e.g., 8 × 32 = 256 GB/s) + +#### cross-PE 접근 (n:1 mode) + +PE0이 PE1의 local HBM에 접근하는 경우: + +```text +PE0.pe_dma → PE0.agg_router → PE1.agg_router → hbm_ctrl +``` + +aggregated router 간 연결: + +```text +pe0.agg_router ↔ pe1.agg_router ↔ pe2.agg_router ↔ ... ↔ pe7.agg_router +``` + +- edge kind: `agg_horizontal` +- BW: configurable (inter-PE aggregated BW) + +#### n:1 mode 전체 데이터 경로 + +```text +PE0.pe_dma → PE0.agg_router → hbm_ctrl + (BW = N × channel_bw_gbs = 256 GB/s) +``` + +--- + +### D6. local / remote access를 NOC로 통일한다 + +- 모든 memory access는 NOC(channel router 또는 aggregated router)를 통해 전달된다 +- local access도 별도의 fast path(xbar)를 사용하지 않는다 +- cross-cube (remote) access 경로: + +```text +1:1 mode: PE_DMA → ch_r{local} → ch_r{...} → UCIe → remote_ch_r → remote_hbm_ctrl +n:1 mode: PE_DMA → agg_router → UCIe → remote_agg_router → remote_hbm_ctrl +``` + +UCIe 연결은 기존 구조를 유지하되, +양쪽 endpoint가 xbar 대신 channel router 또는 aggregated router가 된다. + +--- + +### D7. AddressResolver 변경 + +현재 `AddressResolver.resolve()`: + +```python +# 현재: HBM offset → pe_slice → "sip{s}.cube{c}.hbm_ctrl.slice{pe_slice}" +pe_slice = PhysAddr.hbm_pe_id(addr.hbm_offset, self._slice_size_bytes) +return f"sip{s}.cube{c}.hbm_ctrl.slice{pe_slice}" +``` + +변경 후: + +```python +# 변경: HBM → 단일 endpoint +return f"sip{s}.cube{c}.hbm_ctrl" +``` + +pe_slice 계산이 제거된다. +BAAW가 이미 dst_node를 결정하므로, PE_DMA의 1:1 mode에서는 +resolver를 거치지 않고 BAAW가 직접 channel router node_id를 반환한다. +n:1 mode에서도 BAAW가 aggregated router node_id를 반환한다. + +resolver.resolve()는 외부 접근(M_CPU DMA 등) 및 backward compatibility용으로 유지한다. + +--- + +### D8. topology.yaml 설정 변경 + +#### 추가 설정 + +```yaml +cube: + memory_map: + hbm_mapping_mode: n_to_one # one_to_one | n_to_one + hbm_pseudo_channels: 64 # 전체 pseudo channel 수 + hbm_channels_per_pe: 8 # PE당 local channel 수 (= pseudo_channels / pes_per_cube) + hbm_channel_bw_gbs: 32.0 # per-channel bandwidth (GB/s) + hbm_total_gb_per_cube: 48 # 유지 +``` + +#### 제거 설정 + +```yaml +# 제거 대상 +links: + xbar_to_hbm_bw_gbs: 256.0 # → channel_bw_gbs × channels_per_pe로 대체 + xbar_to_hbm_mm: 2.5 # → ch_router_to_hbm_mm으로 대체 + xbar_to_bridge_bw_gbs: 128.0 # → 제거 (bridge 없음) + xbar_to_bridge_mm: 3.0 # → 제거 + noc_to_xbar_bw_gbs: ... # → 제거 + noc_to_xbar_mm: ... # → 제거 +``` + +#### 추가 link 설정 + +```yaml +links: + pe_to_ch_router_bw_gbs: 32.0 # PE_DMA ↔ channel router + pe_to_ch_router_mm: 1.0 # 물리적 거리 + ch_router_to_hbm_bw_gbs: 32.0 # channel router ↔ hbm_ctrl + ch_router_to_hbm_mm: 2.0 # 물리적 거리 + ch_horizontal_bw_gbs: 32.0 # channel router 간 horizontal link + ch_horizontal_mm: 1.5 # PE 간 horizontal 거리 + # n:1 mode용 + pe_to_agg_router_bw_gbs: 256.0 # PE_DMA ↔ aggregated router + agg_to_hbm_bw_gbs: 256.0 # aggregated router ↔ hbm_ctrl + agg_horizontal_bw_gbs: 256.0 # aggregated router 간 link +``` + +--- + +### D9. 대역폭 수치 정합 + +| 구성 | 값 | +| ---- | --- | +| pseudo channels per cube | 64 (파라미터) | +| PEs per cube | 8 (파라미터) | +| channels per PE (N) | `pseudo_channels / pes_per_cube` = 8 | +| per-channel BW | 32 GB/s (파라미터) | +| per-PE local BW | N × 32 = 256 GB/s | +| cube total HBM BW | 64 × 32 = 2048 GB/s | + +두 모드에서 PE당 effective BW는 동일: + +- 1:1 mode: N개 channel link × channel_bw_gbs = N × 32 = 256 GB/s +- n:1 mode: 1개 aggregated link = N × channel_bw_gbs = 256 GB/s + +--- + +## Consequences + +### Positive + +- 1:1 mode에서 pseudo-channel 단위 BW contention 모델링이 자연스럽다 +- n:1 mode에서 aggregated bandwidth 모델이 단순하다 +- local / remote access 경로가 NOC로 통일된다 +- graph compiler 기반 topology 생성과 잘 맞는다 +- channel 수, PE 수가 모두 파라미터이므로 다양한 구성을 테스트할 수 있다 + +### Negative + +- 1:1 mode에서 router 및 link 수가 크게 증가한다 + (64 channel routers + 64 edges to HBM + 56 horizontal edges per cube) +- local access도 NOC 경로를 사용하므로 모델이 더 일반화된다 +- 기존 xbar 기반 테스트 전면 재작성 필요 +- SimPy 노드 수 증가에 따른 시뮬레이션 성능 영향 가능 + +--- + +## Alternatives + +### A1. 기존 xbar + HBM slice 유지 + +- local/remote 경로가 이원화됨 +- pseudo-channel 단위 모델링 불가 +- 1:1/n:1 mode 전환 불가 + +### A2. per-channel link를 항상 생성하고 n:1에서만 집계 + +- topology 구조가 항상 1:1 크기 +- n:1 semantics를 link aggregation으로 표현하기 복잡 +- router 노드 수 감소 효과 없음 + +### A3. 단계적 전환 (xbar 유지 + NOC 경로 추가) + +- 호환성은 높으나 두 경로 공존으로 복잡도 증가 +- 최종적으로 xbar 제거가 필요하므로 중간 단계의 가치가 낮음 + +--- + +## Implementation Notes + +### topology/builder.py 변경 상세 + +#### 제거할 코드 (현재 `_instantiate_cube()` 내) + +- xbar_top, xbar_bot 노드 생성 (~line 495-508) +- bridge.left, bridge.right 노드 생성 +- noc ↔ xbar edge 생성 (~line 540-555) +- xbar ↔ hbm_ctrl.slice edge 생성 (~line 510-538) +- xbar ↔ bridge edge 생성 (~line 557-572) + +#### 추가할 코드 + +1:1 mode: + +```python +N = hbm_channels_per_pe # from topology config +total_ch = hbm_pseudo_channels + +# channel router 노드 생성 +for ch_id in range(total_ch): + pe_id = ch_id // N + nodes[f"{cp}.ch_r{ch_id}"] = Node( + id=f"{cp}.ch_r{ch_id}", kind="noc_router", impl="noc_v1", + attrs={}, pos_mm=(...), # horizontal row = ch_id % N + ) + +# PE_DMA ↔ local channel router edges +for pe_id in range(pes_per_cube): + for local_ch in range(N): + ch_id = pe_id * N + local_ch + edges.append(Edge( + src=f"{cp}.pe{pe_id}.pe_dma", dst=f"{cp}.ch_r{ch_id}", + bw_gbs=channel_bw, kind="pe_to_ch_router", ...)) + edges.append(Edge( + src=f"{cp}.ch_r{ch_id}", dst=f"{cp}.pe{pe_id}.pe_dma", + bw_gbs=channel_bw, kind="ch_router_to_pe", ...)) + +# channel router ↔ hbm_ctrl edges +for ch_id in range(total_ch): + edges.append(Edge( + src=f"{cp}.ch_r{ch_id}", dst=f"{cp}.hbm_ctrl", + bw_gbs=channel_bw, kind="ch_router_to_hbm", ...)) + edges.append(Edge( + src=f"{cp}.hbm_ctrl", dst=f"{cp}.ch_r{ch_id}", + bw_gbs=channel_bw, kind="hbm_to_ch_router", ...)) + +# horizontal line edges (same logical index) +for row in range(N): + for p in range(pes_per_cube - 1): + ch_a = p * N + row + ch_b = (p + 1) * N + row + edges.append(Edge( + src=f"{cp}.ch_r{ch_a}", dst=f"{cp}.ch_r{ch_b}", + bw_gbs=ch_horizontal_bw, kind="ch_horizontal", ...)) + edges.append(Edge( + src=f"{cp}.ch_r{ch_b}", dst=f"{cp}.ch_r{ch_a}", + bw_gbs=ch_horizontal_bw, kind="ch_horizontal", ...)) +``` + +n:1 mode: + +```python +# aggregated router 노드 생성 +for pe_id in range(pes_per_cube): + nodes[f"{cp}.pe{pe_id}.agg_router"] = Node( + id=f"{cp}.pe{pe_id}.agg_router", kind="noc_router", impl="noc_v1", + attrs={}, pos_mm=(...), + ) + +agg_bw = N * channel_bw # aggregated BW + +# PE_DMA ↔ aggregated router +for pe_id in range(pes_per_cube): + edges.append(Edge( + src=f"{cp}.pe{pe_id}.pe_dma", dst=f"{cp}.pe{pe_id}.agg_router", + bw_gbs=agg_bw, kind="pe_to_agg_router", ...)) + edges.append(Edge( + src=f"{cp}.pe{pe_id}.agg_router", dst=f"{cp}.pe{pe_id}.pe_dma", + bw_gbs=agg_bw, kind="agg_router_to_pe", ...)) + +# aggregated router ↔ hbm_ctrl +for pe_id in range(pes_per_cube): + edges.append(Edge( + src=f"{cp}.pe{pe_id}.agg_router", dst=f"{cp}.hbm_ctrl", + bw_gbs=agg_bw, kind="agg_to_hbm", ...)) + edges.append(Edge( + src=f"{cp}.hbm_ctrl", dst=f"{cp}.pe{pe_id}.agg_router", + bw_gbs=agg_bw, kind="hbm_to_agg", ...)) + +# aggregated router 간 horizontal link +for p in range(pes_per_cube - 1): + edges.append(Edge( + src=f"{cp}.pe{p}.agg_router", dst=f"{cp}.pe{p+1}.agg_router", + bw_gbs=agg_horizontal_bw, kind="agg_horizontal", ...)) + edges.append(Edge( + src=f"{cp}.pe{p+1}.agg_router", dst=f"{cp}.pe{p}.agg_router", + bw_gbs=agg_horizontal_bw, kind="agg_horizontal", ...)) +``` + +### 영향받는 기존 테스트 + +| 테스트 파일 | 영향 | +| ---------- | ---- | +| `tests/test_topology_compile.py` | xbar/bridge 노드 참조 제거, channel router 검증 추가 | +| `tests/test_topology_load.py` | topology.yaml 설정 변경 반영 | +| `tests/test_pe_components.py` | PE_DMA 라우팅 경로 변경 | +| `tests/test_sip_parallel.py` | cross-PE 접근 경로 변경 | +| xbar/bridge를 직접 테스트하는 케이스 | 제거 | + +--- + +## Test Requirements + +- 1:1 mode에서 channel별 link로 request가 전달되는지 확인 +- n:1 mode에서 aggregated link로 request가 전달되는지 확인 +- 두 mode에서 topology가 올바르게 생성되는지 검증: + - 1:1: `total_ch`개 channel router + per-PE link + horizontal link + - n:1: `pes_per_cube`개 aggregated router + per-PE link +- 동일 workload에서 effective BW가 두 모드에서 일관적인지 확인 +- cross-PE 접근 시 horizontal line routing이 동작하는지 확인 +- cross-cube 접근 시 UCIe를 통한 routing이 동작하는지 확인 +- 파라미터 변경 (channels_per_pe = 4, 8, 16 등)에서 topology 생성이 정상인지 확인 + +--- + +## Links + +- ADR-0018 (LA + BAAW) → addressing 측 연동 +- ADR-0017 (Cube NOC 2D Mesh) → 본 ADR이 xbar/bridge 부분을 대체 +- ADR-0004 (Memory Semantics) → BW 모델 재정의 +- ADR-0014 (PE Internal Execution Model) → PE_DMA 경로 변경 영향