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