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

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. 매핑 모드 설정

토폴로지 수준(큐브) 설정:

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 → 물리 요청(들)로 해석한다.

@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.

@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을 소유.

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 분해 없음.

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 (메모리 주소 지정)