# 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 변경 영향