# ADR-0023: PE-level IPCQ — Inter-PE Collective Communication ## Status Accepted ## Context ### 목표 CCL (Collective Communication Library) 커널을 PE 안에서 실행할 수 있도록 PE 간 데이터 교환 인프라를 추가한다. 호스트는 그저 각 SIP에 커널을 launch만 하고, 실제 동기화와 데이터 이동은 **PE 커널 안에서 IPCQ(Inter-Process Communication Queue)를 통해** 일어난다. 이는 NCCL이 GPU 커널 안에서 NVLink 통신을 수행하는 모델, 또는 Cerebras/Tenstorrent의 core-local 통신 큐와 유사하다. 호스트 레벨 collective(`dist.all_reduce`)는 **미래 작업**으로 미루고, 본 ADR은 커널 collective 인프라에만 집중한다. ### 풀어야 할 문제 1. PE 간 직접 데이터 이동 (peer's memory에 write) 2. 동기화 — 송신 측이 수신 측 buffer 공간을 확인해야 함 (backpressure) 3. compute traffic과 communication traffic의 자원 경쟁 (Head-of-Line blocking) 4. 호스트가 알고리즘에 따라 (ring/mesh/tree) 논리적 neighbor 토폴로지를 구성할 수 있어야 함 --- ## Decision ### D1. PE_IPCQ 컴포넌트 신규 추가 PE 안에 새 컴포넌트 `PE_IPCQ`를 추가한다. PE_GEMM/PE_MATH가 PE_CPU의 sub-block을 별도 컴포넌트로 모델링하는 것과 동일한 패턴이다. ``` PE ├── PE_CPU ├── PE_SCHEDULER ├── PE_DMA ├── PE_IPCQ ← 신규 ├── PE_FETCH_STORE ├── PE_GEMM ├── PE_MATH ├── PE_TCM ├── PE_MMU ``` **역할 분리** (control plane vs data plane): - **PE_IPCQ (control plane)**: ring buffer 주소 계산, head/tail pointer 관리, peer pointer 캐시, backpressure 결정, 4-방향 neighbor 매핑 - **PE_DMA (data plane)**: 실제 데이터를 cube_noc/UCIe/PCIE 경유로 peer 메모리에 전송 PE_IPCQ는 데이터 이동을 직접 수행하지 않고 PE_DMA에 위임한다. ### D2. Ring Buffer 모델 각 PE는 4-방향(N/S/E/W) × {tx, rx} = 총 8개의 ring buffer를 가진다. ```python @dataclass class IpcqQueuePair: direction: Direction # N/S/E/W peer: IpcqEndpoint # init 시 호스트가 설정 (D2.5) tx_buffer_base: int # 내가 보낼 데이터의 base addr (자기 메모리) rx_buffer_base: int # 내가 받을 데이터의 base addr (자기 메모리) slot_size: int # tile 단위 n_slots: int # ring depth my_head: int # 내 send 위치 (다음에 쓸 tx/peer slot) my_tail: int # 내 recv 위치 (다음에 읽을 rx slot) peer_head_cache: int # 캐시: peer가 마지막으로 보낸 head 위치 (D9 piggyback으로 갱신) peer_tail_cache: int # 캐시: peer가 마지막으로 소비한 tail 위치 (D9 fast path credit으로 갱신) ``` **필드명 규약 (canonical)**: 본 ADR 전체에서 다음 4개 이름을 일관되게 사용한다. | 필드 | 소유자 | 갱신 시점 | |------|--------|----------| | `my_head` | 자기 PE_IPCQ | tl.send 호출 후 즉시 (송신 추적용) | | `my_tail` | 자기 PE_IPCQ | tl.recv 호출 후 즉시 (수신 추적용) | | `peer_head_cache` | 자기 PE_IPCQ | IpcqMetaArrival 도착 시 (D9 piggyback) | | `peer_tail_cache` | 자기 PE_IPCQ | IpcqCreditMetadata 도착 시 (D9 fast path) | 다른 표현(`peer_head_local`, `peer_head`, `peer_tail` 등)은 사용하지 않는다. **Slot 단위**: fixed-size, 한 slot이 한 tile 데이터를 통째로 담는다. descriptor 모델이 아니라 **full data embedding** 모델 (D5에서 상세). ### D2.5. PeAddress / IpcqEndpoint 스키마 `IpcqQueuePair.peer`가 가져야 할 정보를 명시한다. 송신 측 PE_IPCQ가 peer rx slot에 직접 DMA write하려면 다음을 모두 알아야 한다. ```python @dataclass(frozen=True) class IpcqEndpoint: """송신 측이 peer's rx_buffer 주소를 계산하기 위해 필요한 모든 정보.""" sip: int # 목적지 SIP cube: int # 목적지 cube pe: int # 목적지 PE (cube 내 local index) buffer_kind: str # "tcm" | "hbm" | "sram" — 어느 메모리 공간 rx_base_pa: int # peer rx_buffer base의 PA (PhysAddr.encode()) rx_base_va: int # peer rx_buffer base의 VA (선택, MMU 사용 시) n_slots: int # peer ring depth (경계 wrap-around 계산용) slot_size: int # peer slot 크기 (offset 계산용) ``` `IpcqQueuePair`의 `peer` 필드는 이 `IpcqEndpoint` 객체를 들고 있다. 주소 계산은: ```python slot_idx = self.my_head % peer.n_slots dst_pa = peer.rx_base_pa + slot_idx * peer.slot_size ``` PE_IPCQ는 이 dst_pa를 `IpcqDmaToken`의 dst_addr로 PE_DMA에 전달한다. PE_DMA(vc_comm)는 fabric 라우팅(cube_noc/UCIe/PCIE)을 통해 dst_pa로 데이터를 전송한다. **Endpoint 생성 시점**: backend init (D10)에서 모든 PE의 IPCQ buffer를 allocator로 할당받고, 각 rank의 neighbor table을 만들 때 peer rank의 endpoint 정보를 install한다. 즉 install 순서는: 1. **모든 rank의 IPCQ buffer 할당** (각 PE의 buffer_kind 메모리 공간에서) 2. **rank별 endpoint table 구성** (자신의 4-방향 peer가 어느 sip/cube/pe/pa를 갖는지) 3. **PE_IPCQ에 install** (`IpcqInitMsg` via fabric or sideband) 이 순서는 모든 rank가 서로의 PA를 알아야 하므로, 단계 1을 모든 rank에 대해 먼저 끝낸 후 단계 2-3을 진행한다. ### D3. 4-방향 매핑 = 논리적 ProcessGroup PE는 4방향(N/S/E/W)을 logical port로 본다. 실제 peer 주소는 호스트 CCL init이 알고리즘에 따라 설정한다. PE 커널은 토폴로지를 알지 못하고 방향만 사용한다. ```python # 호스트 init 예시 — 1D ring for rank in range(world_size): ipcq_set_neighbor(rank, "E", peer=ranks[(rank + 1) % world_size]) ipcq_set_neighbor(rank, "W", peer=ranks[(rank - 1) % world_size]) # 호스트 init 예시 — 2D mesh for r in range(R): for c in range(C): ipcq_set_neighbor((r, c), "N", peer=((r - 1) % R, c)) ipcq_set_neighbor((r, c), "S", peer=((r + 1) % R, c)) ipcq_set_neighbor((r, c), "E", peer=(r, (c + 1) % C)) ipcq_set_neighbor((r, c), "W", peer=(r, (c - 1) % C)) ``` PE 코드 입장에서 `tl.send(dir="E", ...)`가 어디로 가는지는 알 필요가 없다. ### D4. PE 커널 API ```python # Send (blocking, backpressure 발생 가능) tl.send(dir: str, src_addr: int, nbytes: int) -> None # Recv (blocking) data = tl.recv(dir: str) # 특정 방향에서 수신 data = tl.recv() # 4방향 round-robin, 도착한 첫 tile 반환 # Recv (non-blocking) handle = tl.recv_async(dir: str) data = tl.wait(handle) ``` `tl.recv()` (방향 미지정)는 IPCQ가 last_polled_dir 인덱스를 들고 있다가 다음 호출 시 그 다음 방향부터 검사하면서 데이터 있는 첫 슬롯을 반환한다. 4방향 모두 비어있으면 wait. **Fairness는 weak fairness**: polling 시작 방향을 회전시켜 단순 편향을 완화하지만, 한 방향에 데이터가 항상 먼저 도착하면 다른 방향이 starvation될 수 있다. strict fairness가 필요한 알고리즘은 `tl.recv(dir=...)`로 방향을 명시해야 한다. (Open Questions 참조) ### D5. Single-hop DMA Write + Full-data Slot 모델 데이터는 송신 측 메모리에서 수신 측 ring slot으로 **단일 DMA 전송**으로 이동한다. 핵심 속성: - **Single-hop**: 송신 측 IPCQ가 peer rx slot 주소를 직접 알고 있어 한 번의 fabric DMA로 데이터가 도착한다. - **No CPU memcpy**: CPU가 데이터를 복사하지 않는다. - **No intermediate staging**: 송신/수신 어느 쪽에도 별도 staging buffer가 없다 (송신은 자기 source 주소에서 직접, 수신은 자기 ring slot으로 직접). (엄밀히 말하면 fabric DMA write 자체는 발생하므로 "data movement가 전혀 없다"는 의미는 아니다. NCCL의 "zero-copy"가 가리키는 것 — CPU memcpy / staging copy 부재 — 과 동일한 속성이다.) 데이터 이동 모델: ``` PE A: tl.send(E, src_addr, nbytes) 1. IPCQ가 peer rx slot 주소 계산 dst_addr = peer.rx_base_pa + (my_head % peer.n_slots) * peer.slot_size 2. backpressure: my_head - peer_tail_cache < peer.n_slots ? (꽉 찼으면 sleep/poll) 3. PE_DMA(vc_comm)에 DMA 요청 → src_addr에서 peer의 dst_addr로 nbytes 전송 4. my_head += 1 PE B: data = tl.recv(W) 1. 내 rx_buffer[my_tail % n_slots] 위치 확인 2. 데이터 도착 대기 (D7 backpressure 모드) 3. 그 주소를 PE 커널에 반환 (또는 fetch unit으로 register file에 로드) 4. my_tail += 1 5. credit return fast path 발행 (D9) — bottleneck-BW latency 후 peer A의 peer_tail_cache 갱신 ``` **핵심**: Slot에 데이터가 통째로 들어간다. PE B의 recv는 자기 rx_buffer만 읽으면 되고, A의 메모리를 read하지 않는다. 송신 측 IPCQ가 peer rx slot 주소를 알고 있으므로 직접 그 주소로 DMA write한다 (single-hop). 본인의 PE_TCM read/write는 DMA를 거치지 않는다 (PE에 직접 붙어있음). slot이 본인 TCM에 있으면 직접 접근, 아니면 PE_DMA 경유. ### D6. Buffer 위치 — 3-way benchmark 호스트 CCL init이 IPCQ ring buffer의 메모리 위치를 결정한다: ```python ipcq_init( backend="ahbm", buffer_kind="tcm" | "hbm" | "sram", n_slots=8, slot_size=4096, ) ``` | 위치 | 특징 | trade-off | |------|------|-----------| | **PE_TCM** | PE에 직접 붙음, 빠름 | 작음, PE 내부 자원과 경쟁 | | **PE-local HBM** | 큼, DMA 경유 | latency 큼 | | **Cube SRAM** | 중간 크기, cube-shared | cube 내 PE 간 contention | 세 위치 모두 동일 코드로 동작하며 init만 다르다. 벤치마크로 비교 가능. **규칙**: peer가 read/write할 때는 DMA 경유. 본인이 자기 PE_TCM 읽기/쓰기는 DMA 없음. ### D7. Backpressure — 2-mode benchmark 송신 측이 peer slot full을 감지했을 때, 또는 수신 측이 데이터 미도착을 감지했을 때 어떻게 대기하는가: | 모드 | 동작 | 모델 | |------|------|------| | **poll** | 캐시된 peer pointer를 주기적으로 재확인. cache update event를 폴링 | spin loop | | **sleep** | SimPy event를 yield하고 sleep, peer가 update event를 trigger하면 wake | interrupt-like | ```python ipcq_init(backpressure="poll" | "sleep", ...) ``` 두 모드 모두 구현하여 latency/throughput trade-off를 벤치마크할 수 있다. ### D8. PE_DMA Virtual Channel PE_DMA를 단일 큐에서 **2-channel virtual channel** 모델로 확장한다. ``` PE_DMA ├── vc_compute: GEMM/MATH의 tile load/store/writeback └── vc_comm: IPCQ의 send 데이터 ``` 각 VC는 독립적인 state machine을 가진다: - 한 채널이 stall되어도 다른 채널은 진행 - 동일 link(cube_noc, UCIe 등)는 공유하지만, link BW는 두 채널이 분할 사용 **Chunk 단위 인터리브**: - 큰 GEMM tile DMA가 한 번에 link를 점유하지 않음 - chunk_size 단위로 진행 (예: 256B), 매 chunk마다 다른 VC와 link BW 공유 - chunk_size는 init 파라미터 (작을수록 fair, 클수록 효율) 이로써: - HoL blocking 해소 (compute DMA 진행 중에도 IPCQ send 끼어들 수 있음) - compute/comm overlap 자연스러움 (NVIDIA copy engine + compute SM 패턴) - HW 모델 정합 (NoC virtual channel은 실제 HW 기법) **첫 구현의 정확도 한계 (intentional)**: 본 ADR의 첫 구현은 **deterministic chunk-level interleave + weighted round-robin arbitration** (default 50/50, ccl.yaml에 노출)을 채택한다. 이는 first-order approximation이며, 실제 HW의 dynamic contention/credit-based arbitration보다는 단순화된 모델이다. | 모델링 항목 | 첫 구현 | 향후 확장 가능 | |------------|---------|---------------| | VC 간 BW 분할 | 정적 weight | dynamic contention 기반 | | Chunk 단위 인터리브 | deterministic round-robin | priority/QoS 기반 | | Cross-VC dependency | 없음 (독립) | NoC arbiter component 추가 | 이 단순화는 functional correctness에는 영향이 없으며, latency 측정에서 heavy contention 시나리오는 실제보다 약간 optimistic한 결과를 낼 수 있다. 정밀화가 필요하면 별도 ADR로 NoC arbiter를 도입한다. #### Token routing - compute용 token (TileToken): 기존 PE_FETCH_STORE → PE_DMA 체이닝 그대로 - comm용 token (IpcqDmaToken, 신규): PE_IPCQ → PE_DMA로 self-routing - PE_DMA가 token 종류로 채널 결정 ```python class PeDmaComponent: def _process(self, env, token): if isinstance(token, IpcqDmaToken): yield from self._vc_comm_process(env, token) else: yield from self._vc_compute_process(env, token) ``` ### D9. Pointer 동기화 — DMA payload piggyback 실제 하드웨어(NVLink, UCIe 등)는 DMA 메시지의 payload에 메타데이터를 piggyback하여 송수신과 함께 pointer를 갱신한다. 본 시뮬레이션도 같은 모델을 채택하여 **별도의 control 채널 없이** 메타데이터가 data와 함께 도착하도록 한다. 이 모델의 핵심 이점: - **자동 ordering**: 메타데이터가 data와 동일 token으로 이동하므로 data가 먼저 visible해진 다음에야 head_cache가 갱신된다. 별도 ordering invariant 없이 race condition이 원천 차단된다. - **HW 정합**: 실제 NVLink/UCIe의 piggybacked header 모델과 일치 - **컴포넌트 단순화**: 별도 IpcqPtrUpdate event 종류가 필요 없음 #### Send 흐름 (head 측 piggyback) ``` PE A: tl.send(E, src_addr, nbytes) 1. PE_IPCQ가 backpressure 체크 (peer_tail_cache 기준) 2. PE_IPCQ가 IpcqDmaToken 생성: - data 본체 (src_addr → peer dst_addr) - piggyback metadata: (sender_seq, src_sip/cube/pe, src_direction) 3. PE_DMA(vc_comm)에 token put 4. PE A는 자기 my_head++ (송신 추적용) [fabric DMA: latency 만큼 진행] PE B의 PE_DMA가 token 수신 5. data를 dst_addr (B의 rx slot)에 MemoryStore.write 6. token의 metadata를 PE B의 PE_IPCQ로 forward (PE 내부 wire, ~1 cycle) PE B의 PE_IPCQ가 metadata 수신 7. peer_head_cache 갱신 (= A의 head 위치) 8. 대기 중인 recv (해당 direction)가 있으면 wake ``` 여기서 핵심은 **5와 6은 같은 SimPy step**이라는 것이다 — DMA 완료와 동시에 data와 metadata가 atomic하게 visible해진다. #### Recv 흐름 (credit return — fast path with bottleneck-BW latency) 수신측이 slot을 비우면 송신측은 그 사실을 알아야 한다 (backpressure 해제). data 경로의 piggyback 모델과 달리, credit return은 일반 vc_comm fabric을 거치지 않고 **별도 fast path**로 처리한다. 이는 실제 HW의 NVLink/UCIe credit return fast path를 추상화한 것이다. **Latency 계산**: magic constant가 아니라 **라우팅 경로의 full path latency** (per-node overhead + edge propagation + drain) 기준으로 산출한다. ``` credit_size_bytes = 16 (ccl.yaml: ipcq_credit_size_bytes) path = router.find_path(self_pe, peer_pe.pe_dma) latency = compute_path_latency_ns(path, credit_size_bytes) = sum(edge.distance_mm * ns_per_mm) + sum(node_overhead_ns[n] for n in path) + credit_size_bytes / bottleneck_bw_on_path ``` router는 source에만 `.pe_dma`를 자동 부여하므로 destination에는 반드시 `.pe_dma` suffix를 명시해야 한다. 그렇지 않으면 `find_path`가 raise하고 credit이 0 cost로 silently teleport되는 latent bug가 발생한다 (이번 업데이트에서 수정됨). `tl.recv`는 credit-emit 완료를 yield-from으로 기다린다 (이전에는 `env.process`로 fork). 이로써 credit-return cost가 receiver의 `pe_exec_ns`에 반영되어, IPCQ control-plane이 consume-acknowledgement를 완료한 뒤에야 recv가 kernel에 반환된다 — RAW DMA의 non-posted `tl.store`가 HBM ack-trip을 기다리는 것의 protocol-level 등가물이다. 이로써: - **토폴로지 비례 approximation**: cube 내 credit return과 cross-SIP credit이 자동으로 다른 latency를 가짐 - **Magic constant 없음**: 모든 ns 값이 데이터 트래픽과 동일한 edge_map 및 `node_overhead_ns`에서 산출되는 `compute_path_latency_ns`로부터 옴 - **Deadlock 위험 없음**: `peer_credit_store.put`은 unbounded, B가 A에게 보낼 데이터가 없어도 credit이 자동 발행됨 - **`IPCQ ≥ raw DMA`** 보장: matched physical move에 대해 credit-emit이 RAW의 ack-trip cost와 균형을 이룸 ``` PE B: tl.recv(W) → 데이터 가져감 → my_tail++ PE B의 PE_IPCQ: 1. router로 PE A까지 path 계산 2. compute_drain_ns(path, credit_size_bytes) = latency_ns 3. env.process(self._delayed_credit_send(latency_ns, peer_credit_store, my_tail)) [fast path: latency_ns 만큼 timeout, fabric vc 미사용] PE A의 PE_IPCQ가 자기 credit_store에서 IpcqCreditMetadata 수신: 4. peer_tail_cache 갱신 5. 대기 중인 send (해당 direction)가 있으면 wake ``` #### Component 결합도 — SimPy Store 채널 PE B의 PE_IPCQ가 PE A의 PE_IPCQ를 직접 호출하지 않는다. 대신 **init 시점에 양쪽 PE_IPCQ 사이에 SimPy Store를 한 번 wire**해두고 (양방향 fast path 채널), credit metadata는 그 store로 put한다. ```python class PeIpcqComponent: def _delayed_credit_send(self, env, peer_credit_store, my_tail, latency_ns): yield env.timeout(latency_ns) yield peer_credit_store.put(IpcqCreditMetadata(seq=my_tail, ...)) ``` backend init에서 IpcqInitMsg fan-out 시 양방향 fast path channel을 함께 설치한다 (D12 IpcqInitMsg에 명시). #### Credit return fast path의 한계 - `credit_size_bytes`는 estimate. 보통 16-64 bytes로 충분하며, 실제 HW의 credit return wire 크기를 모방한 값. - fast path는 일반 vc_comm BW contention 모델에서 **제외**된다 (별도 채널). 실제 HW의 credit return wire는 매우 lightweight이므로 1차 근사로 합리적. - 정밀화가 필요하면 후속 ADR에서: - credit fast path를 별도 link로 모델링 (BW limit + contention) - 또는 piggyback 모드로 변경 가능 (`credit_return_mode: piggyback`) #### PE_DMA의 책임 추가 PE_DMA(vc_comm)는 token 수신 시 다음 시퀀스로 처리한다: Transaction terminal의 BW drain을 먼저 지불하고, 이어서 atomic하게 data write + metadata forward 수행. **data write와 metadata forward 사이에는 SimPy yield를 두어서는 안 된다** (I6 MUST 규칙 참조). drain yield는 atomic 구간 안이 아니라 그 앞에 위치해야 한다: ```python def _on_vc_comm_recv(self, env, txn): # Sender PE_DMA가 찍어 둔 drain_ns (= nbytes / bottleneck_bw) 를 # 여기서 지불. atomic 구간보다 앞이어야 한다 — recv는 bytes가 # "도착"한 이후에만 깨어나야 하므로. drain = getattr(txn, "drain_ns", 0.0) if drain > 0: yield env.timeout(drain) token = txn.request # ── ATOMIC: 두 동작 사이에 yield 금지 ── # 1. data를 dst_addr에 write (dst의 메모리 공간은 token.dst_endpoint.buffer_kind) data = self._memory_store.read(token.src_space, token.src_addr, shape=..., dtype=...) self._memory_store.write(token.dst_endpoint.buffer_kind, token.dst_addr, data) # 2. token의 metadata를 자기 PE의 IPCQ로 forward yield self.out_ports[self._ipcq_id].put(IpcqMetaArrival(token=token)) # ───────────────────────────────────── ``` `out_ports[ipcq_id].put`은 SimPy Store의 yield-able 호출이지만, PE 내부 wire로 capacity가 unbounded인 store를 사용하므로 즉시 완료된다 (실질적으로 single-step). 이 최종 put이 atomic 구간의 끝이며, 그 이전에 다른 yield가 삽입되면 안 된다. #### Drain-at-inbound semantics (D9 timing model) Transaction은 sender PE_DMA가 `drain_ns = nbytes / bottleneck_bw_on_path` 를 찍어 둔 상태로 fabric에 들어간다. 이 simulator에서 per-hop `overhead_ns` 는 각 forwarding component의 `run()` 에서 지불되고, 남은 BW drain은 Transaction의 terminal node에서 한 번 지불된다. IPCQ가 아닌 모든 Transaction (raw DMA, kernel-launch fanout 등) 은 `ComponentBase._forward_txn` 이 terminal에서 이 drain을 지불한다. IPCQ의 경우 목적지 PE_DMA가 `_handle_ipcq_inbound` 핸들러로 Transaction을 가로채서 (IPCQ 전용 data write + metadata forward를 해야 하므로) **이 핸들러 최상단에서 drain을 명시적으로 지불해야 한다** — 그래야 IPCQ의 timing model이 다른 모든 fabric Transaction과 동일선상에 놓인다. 여기서 drain을 지불할 때의 side-effect: - **SRC `tl.send`**: 동작 불변. sender PE_DMA가 `sub_done` 을 `yield` 하지 않으므로 fire-and-forget 의미가 보존된다. metadata forward 이후 호출되는 `sub_done.succeed()` 는 sender 입장에서 listener가 없는 이벤트. - **DST `tl.recv`**: `drain_ns` 만큼 늦게 깨어난다. recv는 local PE_IPCQ 의 `IpcqMetaArrival` 수신 시에만 wake되며, metadata forward가 drain 이후로 이동했으므로 recv는 bandwidth까지 포함한 전체 fabric transfer 시간을 관측하게 된다. 물리적 그림과 일치: send는 dispatch하고 바로 반환; recv는 bytes가 실제로 자신의 inbox로 drain될 때까지 대기. #### Backpressure latency 정확도 backpressure 해제까지 걸리는 시간: - **데이터 send 측 latency** = full fabric DMA (data + piggyback metadata 함께) - **Credit return 측 latency** = fast path with bottleneck-BW (`credit_size_bytes / bottleneck_bw_on_path`) | 시나리오 | 모델링된 latency | 실제 HW와의 관계 | |---------|----------------|----------------| | Cube 내 (fast link) | 작음 (bottleneck = cube_noc BW) | topology-aware approximation | | Cross-cube (UCIe) | 중간 (bottleneck = UCIe BW) | topology-aware approximation | | Cross-SIP (PCIE) | 큼 (bottleneck = PCIE BW) | topology-aware approximation | 별도 magic latency 파라미터 없이 토폴로지에 비례한 first-order approximation이 자동으로 반영된다. 실제 HW와 정확히 일치하지는 않지만 (credit fast path는 contention 모델에서 제외, credit_size_bytes는 estimate), magic constant 모델보다 훨씬 의미 있는 비교 가능. 정밀화는 후속 ADR로 넘긴다. ### D9.5. ADR-0020 (2-Pass) 통합 `tl.send/recv`는 ADR-0020의 2-pass 모델과 통합되어야 한다. Phase 1은 타이밍과 실제 데이터 이동(MemoryStore) 모두 모델링하고, Phase 2는 op_log 기반 정합성 검증을 가능케 한다. #### Phase 1 (타이밍 + 데이터 이동) D9는 head 갱신과 tail 갱신을 다른 메커니즘으로 모델링한다: - **Send-side (head update)** — DMA payload piggyback. data write와 metadata forward가 동일 SimPy step에 일어나므로 자동으로 atomic visibility 보장. - **Recv-side (tail credit return)** — fast path SimPy Store 채널. bottleneck-BW 기반 latency 후 peer_tail_cache 갱신. 두 메커니즘을 합쳐서 전체 ring buffer pointer 일관성을 유지한다. **send 시**: 1. PE_IPCQ가 backpressure 체크 (peer_tail_cache 기준) 2. PE_IPCQ가 IpcqDmaToken 생성 (data + piggyback metadata) → PE_DMA(vc_comm)에 put 3. PE_DMA가 fabric DMA 시뮬레이션 (latency 진행) 4. **DMA 완료와 동일한 SimPy step에 atomic 시퀀스**: - **MemoryStore.write(buffer_kind, dst_pa, data)** — single-hop DMA write - 수신측 PE_IPCQ에 metadata forward → peer_head_cache 갱신 → 대기 recv wake 5. **op_log 기록**: `OpRecord(op_kind="ipcq", op_name="send", params={src_space, src_addr, dst_space, dst_addr, nbytes, dir, dtype, shape, sender_seq})` - `dst_space`는 `token.dst_endpoint.buffer_kind`에서 derive된 값이다 (별도 token 필드가 아니다). dst_addr은 `token.dst_addr`. **recv 시**: 1. PE_IPCQ가 (peer_head_cache > my_tail) AND (MemoryStore.has(slot_addr)) 조건 대기 (D9 piggyback 모델에서는 두 조건이 같은 step에 truthy가 되지만, defensive check) 2. 조건 만족 시: `slot_addr = my_rx_base + slot_idx * slot_size` 3. **두 가지 모드** (`recv_mode`로 op_log에 기록): - **`return_slot`** (default): slot_addr을 그대로 PE 커널에 반환. 데이터 복사 없음. 커널이 slot 메모리를 직접 사용한다. - **`copy_to_dst`**: 호출 시 dst_addr이 지정된 경우. slot 데이터를 읽어서 dst_addr에 write. `data = memory_store.read(...)`; `memory_store.write(dst_space, dst_addr, data)` 4. PE_IPCQ가 my_tail++, fast path credit return을 발행 (D9 — vc_comm fabric을 거치지 않고 별도 SimPy Store 채널로 bottleneck-BW latency 후 peer 측 peer_tail_cache 갱신) 5. **op_log 기록**: `OpRecord(op_kind="ipcq", op_name="recv", params={recv_mode, src_space, src_addr, dst_space, dst_addr, nbytes, dir, dtype, shape, consumer_seq})` - `recv_mode="return_slot"`: src_space/src_addr가 slot 위치, dst_addr=None - `recv_mode="copy_to_dst"`: src_space/src_addr가 slot 위치, dst_space/dst_addr가 사용자 지정 위치 #### Phase 2 (op_log replay) DataExecutor가 `op_kind="ipcq"` 레코드를 만나면: - **send**: src → dst (peer rx slot)로 ndarray를 idempotent하게 write - **recv (`recv_mode="return_slot"`)**: no-op. slot 데이터는 Phase 1에서 이미 적절한 위치에 있으며, 커널이 해당 slot 메모리를 직접 사용함. - **recv (`recv_mode="copy_to_dst"`)**: slot → dst_addr로 ndarray를 idempotent 하게 copy 본질적으로 IPCQ는 **데이터 이동**만 하므로 Phase 2가 추가로 계산할 것은 없다. DataExecutor의 GEMM/Math가 그 데이터를 사용하면 자동으로 정합성이 검증된다. ```python class DataExecutor: def _execute_op(self, op): if op.op_kind == "ipcq": self._execute_ipcq(op) elif op.op_kind == "memory": ... elif op.op_kind == "gemm": ... def _execute_ipcq(self, op): """IPCQ ops are data movement; Phase 1 already wrote to MemoryStore.""" p = op.params if op.op_name == "send": data = self.store.read(p["src_space"], p["src_addr"], shape=p["shape"], dtype=p["dtype"]) self.store.write(p["dst_space"], p["dst_addr"], data) elif op.op_name == "recv": if p.get("recv_mode") == "copy_to_dst": data = self.store.read(p["src_space"], p["src_addr"], shape=p["shape"], dtype=p["dtype"]) self.store.write(p["dst_space"], p["dst_addr"], data) # recv_mode == "return_slot": no-op (data already in slot) ``` #### `--verify-data` 흐름 (CCL 커널) ``` 1. kernbench run --bench ccl_allreduce --verify-data 2. backend init → IPCQ buffers 할당, neighbor table install 3. 모든 rank greenlet 동시 실행 4. 각 PE 커널이 tl.send/recv → MemoryStore에 데이터 누적 5. 시뮬레이션 완료 후 DataExecutor.run() → ipcq op 멱등 replay (no-op) 6. 벤치마크가 print(out) 또는 out.data 비교 → 정합성 확인 ``` 벤치 작성자는 `out.data`로 결과를 읽고 expected와 비교하면 된다 (ADR-0020 D7 Tensor.data 패턴). ### D10. 호스트 CCL Init은 PyTorch 패턴 그대로 호스트 코드는 실제 PyTorch distributed 코드와 동일하게 유지한다. `init_process_group`은 backend 객체만 만들고, IPCQ 설정 (neighbor topology, buffer_kind, backpressure 등)은 받지 않는다. ```python # benches/ccl_allreduce.py — 실제 PyTorch와 동일한 호스트 코드 def run_rank(rank, world_size, torch): dist = torch.distributed dist.init_process_group(backend="ahbm", world_size=world_size, rank=rank) tensor = torch.zeros((M, K), dtype="f16", dp=...) from kernbench.ccl.algorithms import ring_allreduce torch.launch("ring_allreduce", ring_allreduce.kernel, tensor, rank, world_size) ``` IPCQ 설정은 backend가 **init_process_group 시점에** `ccl.yaml`을 읽고 즉시 PE_IPCQ neighbor table을 install한다. 호스트 코드는 IPCQ를 인지할 필요가 없다. 벤치마크 하나는 하나의 알고리즘을 사용하는 것을 가정하며, 사용할 알고리즘은 `ccl.yaml`의 `defaults.algorithm` 으로 지정한다 (D11). 호스트 코드 변경 없이 ccl.yaml만 수정하여 다른 알고리즘으로 교체할 수 있다. #### Init 흐름 (eager) 1. `init_process_group(backend="ahbm")` 호출 2. backend가 `ccl.yaml` 로드 → `defaults.algorithm` 결정 3. `algorithms[]`에서 topology + buffer_kind + backpressure + slot/size 결정 4. **즉시** 모든 PE의 PE_IPCQ에 neighbor table을 install (sideband 또는 fabric `IpcqInitMsg`) 5. 이후 `torch.launch(kernel_name, ...)`는 일반 launch와 동일하게 처리 (CCL kernel이든 아니든 PE_IPCQ는 이미 준비됨) ### D11. CCL 설정 파일 (`ccl.yaml`) IPCQ 설정과 알고리즘 metadata는 별도 YAML 파일에 둔다. `components.yaml`/`topology.yaml`과 같은 패턴을 유지하며, 변경 이력이 코드처럼 추적 가능하다. 벤치마크 한 번 실행은 한 알고리즘만 사용한다 (`defaults.algorithm`). 다른 알고리즘으로 교체하려면 `ccl.yaml`의 `defaults.algorithm` 만 바꾸면 된다. ```yaml # ccl.yaml — CCL backend (ahbm) configuration # # 이 파일은 init_process_group(backend="ahbm") 시점에 로드되며, # defaults.algorithm 으로 지정된 알고리즘에 따라 PE_IPCQ neighbor table을 # install한다. 호스트 코드는 IPCQ 설정을 인지하지 않는다. defaults: # 이번 벤치 실행에서 사용할 알고리즘. algorithms 섹션에 정의된 것 중 하나. algorithm: ring_allreduce # IPCQ ring buffer가 위치할 메모리. # tcm — PE-local TCM (작지만 빠름, PE 내부 자원과 경쟁) # hbm — PE-local HBM (큼, DMA latency 큼) # sram — Cube-shared SRAM (중간 크기, cube 내 PE 간 contention) buffer_kind: tcm # send/recv가 peer slot full / data 미도착을 만났을 때의 대기 방식. # poll — peer pointer 캐시를 spin loop로 재확인 # sleep — SimPy event yield 후 wakeup 대기 (interrupt-like) backpressure: sleep # Ring buffer depth (한 방향당 slot 개수). 클수록 in-flight 가능, 메모리 ↑ n_slots: 8 # Slot 하나의 크기 (bytes). 한 tile을 통째로 담을 수 있는 크기여야 함. slot_size: 4096 # PE_DMA virtual channel chunk 크기 (bytes). 작을수록 fair, 클수록 효율. # IPCQ traffic과 compute traffic 사이의 인터리브 granularity (D8 참조). vc_chunk_size: 256 # Credit return fast path 메시지 크기 (bytes). 실제 HW의 credit return wire # 크기를 모방. backend가 라우팅 경로의 bottleneck BW를 보고 latency를 # 계산한다 (D9 참조). 보통 16-64로 충분. ipcq_credit_size_bytes: 16 algorithms: # ── 알고리즘 정의 ───────────────────────────────────────────────── # 각 entry는 알고리즘 모듈과 그 알고리즘이 요구하는 topology를 명시한다. # 알고리즘별 default override 가능 (buffer_kind, backpressure 등). ring_allreduce: # PE 커널이 정의된 모듈. `kernel(t_ptr, rank, world_size, tl)` 함수를 export. module: kernbench.ccl.algorithms.ring_allreduce # 이 알고리즘이 요구하는 neighbor topology. builtin 이름 또는 "custom". # ring_1d — 1D 양방향 ring (E/W) # ring_1d_unidir — 1D 단방향 ring (E only) # mesh_2d — 2D mesh (N/S/E/W) # tree_binary — binary tree (parent/children direction) # custom — 모듈의 neighbors(rank, world_size) 함수 사용 topology: ring_1d tree_allreduce: module: kernbench.ccl.algorithms.tree_allreduce topology: tree_binary # 알고리즘별 override (이 알고리즘만 hbm 사용) buffer_kind: hbm custom_mesh: module: kernbench.ccl.algorithms.custom_mesh topology: custom # 모듈이 직접 neighbors() 함수 제공 ``` #### 알고리즘 모듈 구조 알고리즘 모듈은 두 개의 hook을 export한다 — `kernel`은 필수, `neighbors`는 선택. ```python # src/kernbench/ccl/algorithms/ring_allreduce.py def kernel(t_ptr, rank, world_size, tl): """필수 — PE 커널. IPCQ 설정은 backend가 ccl.yaml + neighbors() 결과로 install한 상태이다. 커널은 그저 4-방향 send/recv API만 사용하면 된다. """ for step in range(world_size - 1): ... tl.send(dir="E", ...) data = tl.recv(dir="W") def neighbors(rank, world_size, neighbor_map): """선택 — neighbor table override hook. backend는 ccl.yaml의 topology 필드에 따라 builtin neighbor_map을 생성한 뒤, 이 함수가 정의되어 있으면 호출하여 결과를 override 한다. Args: rank: 이 rank의 인덱스 world_size: 전체 rank 수 neighbor_map: ccl.yaml의 topology 필드가 만든 builtin 매핑 예: ring_1d → {"E": (rank+1)%ws, "W": (rank-1)%ws} mutable dict — 직접 수정 가능 Returns: dict | None: dict — neighbor_map을 override한 결과 None — override 안 함, neighbor_map 그대로 사용 """ return None # 또는 수정 후 반환 ``` #### `neighbors` override 패턴 대부분의 알고리즘은 builtin topology만으로 충분하므로 `neighbors` 정의가 필요 없다. 정의가 필요한 경우의 패턴: **Pattern A — builtin을 base로 일부만 수정**: ```python def neighbors(rank, world_size, neighbor_map): # 짝수 rank만 W 사용 if rank % 2 == 1: neighbor_map.pop("W", None) return neighbor_map ``` **Pattern B — 완전히 새로 만들기 (skip-connection ring 등)**: ```python def neighbors(rank, world_size, neighbor_map): # neighbor_map은 무시하고 새로 작성 return {"E": (rank + 2) % world_size} ``` #### Builtin topology generators `ccl.yaml`의 `topology` 필드가 다음 builtin 이름이면 backend가 알아서 처리: | topology | 설명 | direction set | |----------|------|---------------| | `ring_1d` | 1D 양방향 ring | E, W | | `ring_1d_unidir` | 1D 단방향 ring | E only | | `mesh_2d` | 2D mesh | N, S, E, W | | `tree_binary` | binary tree (root = rank 0) | parent, child_left, child_right | | `none` | 빈 매핑 — 알고리즘이 `neighbors()`로 처음부터 작성 | (없음) | `topology: none`은 builtin이 빈 dict를 반환하므로 알고리즘의 `neighbors()`가 처음부터 매핑을 만들어야 한다. #### 알고리즘 추가 절차 1. `src/kernbench/ccl/algorithms/.py`에 `kernel` 함수 작성 2. `ccl.yaml`의 `algorithms` 섹션에 entry 추가 (`module`, `topology`) 3. (선택) 같은 모듈에 `neighbors()` 함수 추가하여 builtin override 4. `defaults.algorithm`을 새 알고리즘으로 설정하면 적용 호스트 코드는 손대지 않는다. ### D12. 메시지 / 토큰 스키마 본 ADR이 추가하는 모든 메시지/토큰의 필드를 명시한다. 구현 시 이 정의를 `src/kernbench/common/pe_commands.py`와 `src/kernbench/runtime_api/kernel.py`에 그대로 추가한다. #### `IpcqInitMsg` (sideband, init 시 fan-out) backend가 모든 PE의 PE_IPCQ에 neighbor table을 install하기 위해 사용한다. 구조는 `MmuMapMsg`와 유사 (target_sips, target_cubes, target_pe + entries). ```python @dataclass(frozen=True) class IpcqInitEntry: direction: str # "N" | "S" | "E" | "W" peer: IpcqEndpoint # D2.5 참조 my_rx_base_pa: int # 자신의 rx_buffer base my_rx_base_va: int # 선택 n_slots: int slot_size: int # Credit fast path 채널 (D9). # 계약: 이 필드는 반드시 simpy.Store 인스턴스이며, IpcqCreditMetadata # 객체만을 받는 receive endpoint이다 (peer's PE_IPCQ가 자기 입력 큐로 # 사용). 송신 측 PE_IPCQ는 _delayed_credit_send에서 이 store에 직접 # IpcqCreditMetadata를 put한다. 다른 객체 type을 put해서는 안 된다. # backend init 시 양방향 SimPy Store가 한 번 wire되며 이후 변경 불가. peer_credit_store: "simpy.Store[IpcqCreditMetadata]" @dataclass(frozen=True) class IpcqInitMsg: correlation_id: str request_id: str target_sips: tuple[int, ...] target_cubes: tuple[int, ...] target_pe: int | tuple[int, ...] | str entries: tuple[IpcqInitEntry, ...] # 이 PE의 4-방향 entry backpressure_mode: str # "poll" | "sleep" buffer_kind: str # "tcm" | "hbm" | "sram" credit_size_bytes: int # D9 fast path latency 계산용 (default 16) ``` **Credit fast path channel wiring**: backend init이 모든 PE의 PE_IPCQ에 양방향 fast path 채널을 한 번 설치한다. PE A의 IpcqInitEntry(direction=E)에 PE B의 credit-receive Store reference를 넣어 송신 측이 직접 put할 수 있게 한다 (별도 fabric routing 없음). #### `IpcqSendCmd` (PE_CPU → PE_IPCQ) ```python @dataclass(frozen=True) class IpcqSendCmd: direction: str # 어느 방향으로 보낼지 src_addr: int # 보낼 데이터의 원본 주소 (TCM/HBM) src_space: str # "tcm" | "hbm" | "sram" nbytes: int shape: tuple[int, ...] # data shape (op_log/MemoryStore용) dtype: str handle_id: str # completion 추적용 data_op: bool = True # ADR-0020 op_log 기록 대상 ``` #### `IpcqRecvCmd` (PE_CPU → PE_IPCQ) ```python @dataclass(frozen=True) class IpcqRecvCmd: direction: str | None # None이면 round-robin (weak fairness, D4) # recv_mode: 두 가지 동작 모드 # "return_slot" — slot 주소를 그대로 PE 커널에 반환 (default, zero-copy) # "copy_to_dst" — slot 데이터를 dst_addr에 copy 후 반환 recv_mode: str = "return_slot" # dst_addr / dst_space는 recv_mode="copy_to_dst"일 때만 사용됨 dst_addr: int = 0 dst_space: str = "" shape: tuple[int, ...] = () # data shape (op_log/MemoryStore용) dtype: str = "" handle_id: str = "" blocking: bool = True # blocking vs non-blocking data_op: bool = True ``` #### `IpcqDmaToken` (PE_IPCQ → PE_DMA, vc_comm 채널) D9의 piggyback 모델에 따라 token이 data + head metadata를 함께 담아 fabric을 따라 이동한다. 수신 측 PE_DMA가 도착 시점에 data를 dst_addr에 write하고 metadata를 PE_IPCQ로 forward한다 (atomic). ```python @dataclass class IpcqDmaToken: # ── Data movement (single-hop DMA write) ── src_addr: int # 자기 메모리 주소 src_space: str dst_addr: int # peer rx slot 주소 (이미 계산됨) dst_endpoint: IpcqEndpoint # 라우팅용 (sip/cube/pe) nbytes: int # data 크기 handle_id: str # 완료 시 송신 측 PE_IPCQ로 알림 # ── Piggyback metadata (수신측 PE_IPCQ가 자동 갱신할 정보) ── sender_seq: int # 단조 증가 sequence number # peer가 자기 head_cache로 사용 src_sip: int # 송신 측 (수신측이 어느 peer인지 식별) src_cube: int src_pe: int src_direction: str # 송신측 기준 방향 (수신측은 reverse 매핑으로 자기 direction 결정) data_op: bool = True # ADR-0020 op_log 기록 대상 ``` PE_DMA는 token type으로 채널 결정 (D8): TileToken → vc_compute, IpcqDmaToken → vc_comm. **수신 측 PE_DMA의 처리** (vc_comm 도착 시): ```python def _vc_comm_arrival(self, env, token: IpcqDmaToken): # 1. data를 dst_addr에 write (data와 metadata atomic visibility) if self._memory_store is not None: data = self._memory_store.read(token.src_space, token.src_addr, shape=..., dtype=...) self._memory_store.write(token.dst_endpoint.buffer_kind, token.dst_addr, data) # 2. metadata를 자기 PE의 IPCQ로 forward (PE 내부 wire, 같은 step) yield self.out_ports[self._ipcq_id].put(IpcqMetaArrival(token=token)) ``` PE_IPCQ는 `IpcqMetaArrival`을 받아 sender_seq를 보고 peer_head_cache를 갱신한다. #### `IpcqCreditMetadata` (PE_IPCQ → peer PE_IPCQ, fast path 채널) Credit return은 D9의 fast path 모델에 따라 vc_comm fabric을 거치지 않고 **별도의 SimPy Store 채널**로 전달된다. backend init 시 양방향 channel이 미리 wire되며, latency는 bottleneck-BW 기반으로 계산된다. ```python @dataclass(frozen=True) class IpcqCreditMetadata: """Credit return — recv 측 → send 측 fast path.""" consumer_seq: int # my_tail (recv 측의 새 tail) src_sip: int # 누가 보냈는지 (수신 측이 어느 peer credit인지 식별) src_cube: int src_pe: int src_direction: str # 송신 측 기준 방향 (수신 측은 reverse 매핑) ``` **전송 흐름**: ```python class PeIpcqComponent: def _delayed_credit_send(self, env, peer_credit_store, my_tail, latency_ns): yield env.timeout(latency_ns) yield peer_credit_store.put(IpcqCreditMetadata( consumer_seq=my_tail, src_sip=..., src_cube=..., src_pe=..., src_direction=..., )) ``` `latency_ns`는 D9에 정의된 대로: ```python path = self.ctx.router.find_path(self_pe_prefix, peer_pe_prefix) latency_ns = self.ctx.compute_drain_ns(path, credit_size_bytes) ``` **별도의 IpcqPtrUpdate 이벤트는 없다** — head 갱신은 D9 piggyback 모델로, tail 갱신은 D9 fast path SimPy Store 채널로 처리된다. ### D13. 테스트 전략 단위/통합/regression 테스트를 명시한다. #### T1. 단위 테스트 (component-level) - **PE_IPCQ 단위** (`tests/test_pe_ipcq.py`): - send: backpressure 미발생 시 즉시 PE_DMA로 token forward - send: peer slot full → backpressure (poll/sleep 모드별) - send: peer credit return (IpcqCreditMetadata) 도착 후 backpressure 해제 - recv: 데이터 도착 시 즉시 반환 - recv: 데이터 미도착 → wait → IpcqMetaArrival (D9 piggyback) 수신 시 wake - recv (round-robin): 4-방향 중 도착한 첫 데이터 반환 (weak fairness) - 잘못된 방향 → IpcqInvalidDirection 예외 - **PE_DMA virtual channel** (`tests/test_pe_dma_vc.py`): - vc_compute / vc_comm 독립 진행 (한 채널 stall 시 다른 채널 진행) - chunk-level 인터리브 verification - link BW 분할 (50/50 또는 weighted) - **builtin topology** (`tests/test_ccl_topologies.py`): - ring_1d/mesh_2d/tree_binary 각각 (rank, world_size) → neighbor dict 정합성 - mesh_2d non-square → ValueError - resolve_topology(custom, module) → module.neighbors 반환 #### T2. 통합 테스트 (E2E send/recv) - **`tests/test_ipcq_e2e.py`**: - 2-rank ring: rank 0 send(E) → rank 1 recv(W) → 데이터 정합성 - 4-rank ring: 양방향 send/recv 동시 진행, deadlock 없음 - mesh_2d 4×4: N/S/E/W 4방향 동시 send/recv - **CCL kernel + 2-pass** (`tests/test_ipcq_2pass.py`): - greenlet 모드 + IPCQ → op_log에 ipcq 레코드 생성 검증 - DataExecutor가 ipcq op 처리 후 결과 정합성 (`out.data` 확인) #### T3. Backend init 테스트 (`tests/test_ccl_backend_ipcq.py`) - ccl.yaml 로드 → `defaults.algorithm` 추출 - builtin topology → IpcqInitMsg fan-out - IpcqEndpoint의 PA가 모든 PE에서 일관 (rank A의 peer E의 rx_base_pa = rank A+1의 자기 rx_base_pa) - buffer_kind 별 메모리 할당 (tcm/hbm/sram) #### T4. Regression - 기존 401 tests 전부 PASS - ADR-0020 통합으로 인한 op_log/DataExecutor 영향 없음 (CCL 미사용 벤치) #### T5. 성능 / overhead - 단일 send/recv pair latency = (DMA latency) + (IPCQ overhead) - 비교: 같은 nbytes의 일반 PE_DMA write와 거의 동일해야 함 (IPCQ overhead < 100 ns) ### D14. Invariants & Failure Modes CCL 인프라에서 흔히 발생하는 hang/오류 상황을 명시하고, 대응 방식을 정의한다. #### Invariants (시뮬레이션이 보장해야 하는 것) I1. **Slot lifecycle exactly-once**: 한 send → 정확히 한 recv. 중복 send나 중복 recv는 sequence 오류로 간주. I2. **Pointer monotonicity**: my_head, my_tail은 단조 증가 (감소 없음). sender_seq는 송신 측에서 단조 증가, 수신 측 cache 갱신도 단조 증가. I3. **Endpoint consistency**: rank A의 IpcqEndpoint(direction=E)의 peer가 rank B라면, rank B의 IpcqEndpoint(reverse(E))의 peer는 rank A여야 함. backend init 시 검증. I4. **buffer_kind consistency**: 한 ProcessGroup 내 모든 PE의 buffer_kind는 동일 (mixed kind는 supported 안 함, 첫 구현). 검증 실패 시 init 에러. I5. **op_log ordering**: send → DMA 완료 → recv 가능. op_log의 t_start 순서가 이 인과관계를 위배하지 않음. I6. **Atomic data + metadata visibility (MUST)**: 본 ADR의 correctness 핵심 조건이다. 수신 측에서 data write (MemoryStore.write)와 metadata forward (peer_head_cache 갱신)는 동일한 SimPy step에 일어나야 한다. control이 data를 앞지를 수 없다. **구현 규칙 (MUST)**: - PE_DMA의 vc_comm token 도착 처리(`_vc_comm_arrival`)는 다음 두 동작 사이에 **어떤 SimPy yield도 두어서는 안 된다**: 1. `MemoryStore.write(token.dst_endpoint.buffer_kind, token.dst_addr, data)` 2. PE_IPCQ에 `IpcqMetaArrival` forward - 두 동작은 동일 SimPy event callback 내에서 연속 실행되어야 한다. - 코드 리뷰에서 이 사이에 `yield` (또는 `yield from`)을 추가하는 것은 correctness 위반으로 reject한다. 이 규칙을 위반하면 다른 SimPy process가 끼어들어 head_cache가 data visibility보다 먼저 또는 늦게 보이는 race condition이 발생한다. I7. **MemoryStore slot existence ↔ pointer**: I6의 결과로, `peer_head_cache > my_tail`이 truthy가 되는 step과 `MemoryStore.has(slot_addr)` 이 truthy가 되는 step이 동일하다. recv는 두 조건을 모두 체크하지만 (defensive), 단일 조건만 체크해도 정확하다. #### Failure Modes (런타임 에러) F1. **잘못된 direction**: - PE 커널이 `tl.send(dir="X")` 호출 → install 안 된 direction - PE_IPCQ가 즉시 `IpcqInvalidDirection` 예외 raise - SimPy 시뮬레이션 즉시 abort, 사용자에게 명확한 에러 F2. **타입 mismatch**: - send와 recv의 dtype/shape/nbytes가 일치하지 않음 - 첫 구현은 검증 안 함 (dtype/shape는 hint), 향후 strict mode로 추가 F3. **Deadlock detection (timeout 기반)**: - send: peer_tail_cache가 갱신 안 되고 영원히 wait - recv: peer_head_cache 갱신 안 되고 영원히 wait - 시뮬레이션 timeout (default 10ms simulated time) 초과 시 abort - 디버그를 위해 각 PE의 last send/recv 위치, blocking 상태 dump F4. **Backend init 실패**: - ccl.yaml에 `defaults.algorithm` 누락 - `algorithms[name]` 정의 누락 - 알고리즘 모듈 import 실패 - topology 검증 실패 (I3, I4) → 모두 `init_process_group` 시점에 즉시 에러 F5. **Slot full + 무한 backpressure**: - peer가 영원히 안 받음 - F3과 같이 timeout으로 처리 - 디버그: 막힌 PE의 my_head, peer_tail_cache 출력 #### 진단 도구 (구현 단계에서 추가) - **CCL trace**: 각 send/recv를 (rank, t, dir, nbytes) 형태로 로깅 - **Pointer dump**: 시뮬레이션 종료 시 또는 hang 시 모든 PE의 IPCQ pointer 상태 출력 - **Deadlock graph**: hang 발생 시 wait-for 그래프 출력 (어느 PE가 어떤 PE를 기다리는지) ### D15. 알고리즘 작성자 가이드 (요약) 본 섹션은 알고리즘 작성자가 한 화면으로 시작점을 잡을 수 있도록 한다. 자세한 step-by-step 가이드는 [docs/onboarding/ccl-author-guide.md](../onboarding/ccl-author-guide.md) 참조. #### 만지는 것 / 만지지 않는 것 | 만지는 것 | 만지지 않는 것 | |----------|---------------| | `src/kernbench/ccl/algorithms/.py` (kernel + 선택적 neighbors) | `benches/ccl_allreduce.py` 호스트 코드 | | `ccl.yaml` 의 한 entry 추가 + `defaults.algorithm` | `src/kernbench/ccl/` 프레임워크 | | (선택) `tests/test_.py` 단위 테스트 | `src/kernbench/components/builtin/pe_ipcq.py` 컴포넌트 | | | `src/kernbench/runtime_api/distributed.py` backend | #### 알고리즘 모듈 인터페이스 contract ```python # src/kernbench/ccl/algorithms/.py def kernel(*args, tl) -> None: """필수. PE 커널. Args (positional): tensor pointers, rank, world_size, 알고리즘 파라미터 Args (keyword): tl — TLContext (자동 주입) 사용 가능한 IPCQ API: tl.send(dir, src_addr, nbytes) # blocking, backpressure 시 wait tl.recv(dir) # 특정 방향에서 blocking recv tl.recv() # 4방향 round-robin tl.recv_async(dir) → handle # non-blocking tl.wait(handle) # non-blocking 완료 대기 기존 API도 그대로 사용: tl.load / tl.store / tl.composite / tl.program_id 등 """ ... def neighbors(rank, world_size, neighbor_map) -> dict | None: """선택. ccl.yaml의 builtin topology가 만든 neighbor_map을 override. None 반환 → builtin 그대로 사용 dict 반환 → 그 dict로 override (builtin을 base로 수정 가능) """ return None ``` #### 5-step 흐름 1. **kernel 함수 작성** — `src/kernbench/ccl/algorithms/.py` 신규 파일 2. **ccl.yaml 등록** — `algorithms.` entry + `defaults.algorithm` 변경 3. **(선택) neighbors override** — builtin topology를 base로 수정이 필요할 때 4. **단위 테스트** — `kernbench.ccl.testing.run_kernel_in_mock` (SimPy 없이 빠름) 5. **시뮬 검증** — `kernbench run --bench ccl_allreduce --verify-data` 호스트 코드 (`benches/ccl_allreduce.py`)는 손대지 않는다. #### 사용 가능한 헬퍼 (`kernbench.ccl.helpers`) | Helper | 설명 | |--------|------| | `chunked(addr, n_chunks, ...)` | 텐서를 n개 chunk view로 슬라이싱 | | `ring_step(rank, step, ws)` | ring algorithm의 step별 (send_idx, recv_idx) | | `tree_step(rank, level)` | binary tree의 level별 parent/child 인덱스 | #### 디버깅 도구 - `KERNBENCH_CCL_TRACE=1` — send/recv trace 출력 - 시뮬 종료 시 자동 IPCQ pointer dump - Deadlock 시 (10ms 시뮬 시간 초과) wait-for graph dump #### 흔한 실수 1. **install 안 된 direction 사용** — ccl.yaml의 topology가 ring_1d면 N/S 사용 불가 2. **send/recv 짝 맞지 않음** — peer 측 recv 없으면 hang (slot full backpressure) 3. **dtype/shape 불일치** — 첫 구현은 검증 안 함, 작성자 책임 자세한 step-by-step과 hello-world 예제는 `docs/onboarding/ccl-author-guide.md` 참조. --- ## HW Realization Notes (Informative) **Status of this section**: Forward-looking. Describes how the simulator contract (D1–D15) would map to silicon. Not currently implemented; subject to revision before tapeout. The simulator implements the contract via Python/SimPy equivalents in [pe_ipcq.py](../../src/kernbench/components/builtin/pe_ipcq.py) and [pe_dma.py](../../src/kernbench/components/builtin/pe_dma.py). ### D16. Proposed HW Block Diagram and End-to-End Dataflow ![PE Baseline Architecture](../diagrams/pe_baseline.png) > Source: [`../diagrams/pe_baseline.d2`](../diagrams/pe_baseline.d2) — `d2 --layout=elk --scale 1.5`. ![PE Proposed Architecture](../diagrams/pe_proposed.png) > Source: [`../diagrams/pe_proposed.d2`](../diagrams/pe_proposed.d2) — `d2 --layout=elk`. **Baseline → Proposed 핵심 변경**: - 단일 FIFO inbox → **compute port / IPCQ port 분리 + WRR Arbiter** (NEW) - PE_IPCQ (SimPy component) → **IPCQ Controller** (HW register + combinational logic) - TCM 내 **IPCQ Slot Region 예약 영역** 명시 - Credit Injector / Receiver가 Fabric Port를 통해 NoC에 직접 연결 #### End-to-End Sequence (HW view) ```mermaid sequenceDiagram participant CPU_A as PE_A: PE_CPU participant IPCQ_A as PE_A: IPCQ Ctrl participant DMA_A as PE_A: DMA participant NOC as NoC Fabric participant DMA_B as PE_B: DMA participant IPCQ_B as PE_B: IPCQ Ctrl participant TCM_B as PE_B: TCM participant CPU_B as PE_B: PE_CPU Note over CPU_A: tl.send(dir="E", src=0x1000) CPU_A->>IPCQ_A: MMIO: send request Note over IPCQ_A: Backpressure check:
(head - peer_tail_cache) < n_slots → PASS
Slot addr gen:
dst = peer_rx_base + (head%n) × slot_size IPCQ_A->>DMA_A: IpcqDmaToken {src, dst, sender_seq=head} Note over IPCQ_A: my_head++ IPCQ_A-->>CPU_A: send returns (fire-and-forget) Note over DMA_A: TCM read → snapshot in read buffer
Flit pack: data + {sender_seq, dst_addr} DMA_A->>NOC: IPCQ data flit(s) Note over NOC: hop latency + BW drain NOC->>DMA_B: IPCQ data flit(s) Note over DMA_B: Terminal BW drain
Slot write latency rect rgb(255, 240, 220) Note over DMA_B,IPCQ_B: ATOMIC (I6): same cycle, no stall DMA_B->>TCM_B: write data → slot address DMA_B->>IPCQ_B: Meta Extractor: {sender_seq, dst_addr} end Note over IPCQ_B: Range match dst_addr → direction "W"
peer_head_cache["W"] = sender_seq + 1 IPCQ_B-->>CPU_B: recv_wake signal Note over CPU_B: tl.recv(dir="W") wakes up CPU_B->>IPCQ_B: recv request Note over IPCQ_B: peer_head_cache > my_tail → YES
slot_addr = rx_base + (tail%n) × slot_size IPCQ_B-->>CPU_B: return slot_addr CPU_B->>TCM_B: read data from slot Note over IPCQ_B: my_tail++ IPCQ_B->>NOC: Credit (16B): {consumer_seq, dst_rx_base_pa} Note over NOC: credit traversal (NoC latency) NOC->>IPCQ_A: Credit arrival Note over IPCQ_A: Match dst_rx_base_pa → direction "E"
peer_tail_cache["E"] = consumer_seq
Backpressure deassert (if stalled) ``` ### D17. IPCQ Controller HW Module (신규) PE_CPU와 DMA Engine 사이에 위치하는 하드웨어 제어 블록. 시뮬레이터의 `PeIpcqComponent`에 대응한다. #### QPair Register File 방향별 queue pair 상태를 flip-flop으로 유지. PE_CPU가 MMIO(CSR)로 읽기/쓰기 가능하며, init 시점에 소프트웨어가 채워넣는다. ``` Per-direction registers (each 64-bit): my_head — sender write position (monotonic) my_tail — receiver read position (monotonic) peer_head_cache — last known peer head (updated by Meta Extractor) peer_tail_cache — last known peer tail (updated by Credit Receiver) rx_base_pa — this PE's rx buffer base physical address peer_rx_base_pa — peer's rx buffer base physical address n_slots — ring depth (power-of-2 제약, D21 참조) slot_size — bytes per slot peer_credit_tgt — peer PE의 credit receive 주소 Directions: 최대 8 (N/S/E/W/parent/child_left/child_right + spare) Total: 8 dirs × 9 regs × 8B = 576B flip-flops ``` #### Slot Address Generator (combinational) ``` Input: pointer (my_head or my_tail), n_slots, slot_size, base_pa Output: slot_addr = base_pa + (pointer % n_slots) * slot_size Implementation: n_slots power-of-2 → pointer & (n_slots - 1) (AND mask, 1 gate) slot_size power-of-2 → barrel shift (1 cycle) 64-bit add → ripple/kogge-stone adder (1 cycle) Latency: 1-2 cycles combinational ``` #### Backpressure Comparator (combinational) ``` full = (my_head - peer_tail_cache) >= n_slots Implementation: 64-bit subtract + unsigned compare Output: stall signal → PE_CPU (IPCQ send blocked) or DMA issue hold Latency: 1 cycle ``` #### Meta Extractor (inbound datapath sideband) DMA Engine의 inbound vc_comm path에 wired. 도착하는 IPCQ flit의 header에서 metadata를 추출하여 queue pair 상태를 갱신한다. ``` Trigger: DMA inbound write completion (same cycle) Extract: {sender_seq, dst_addr} from flit header Direction matching (ADR-0025 D2): for each dir: match = (base_pa[dir] <= dst_addr) && (dst_addr < base_pa[dir] + n_slots[dir] * slot_size[dir]) 8× parallel range comparators + priority encoder Update: peer_head_cache[matched_dir] = max(peer_head_cache, sender_seq + 1) Output: recv_wake signal → PE_CPU interrupt/flag Latency: 1 cycle (pipelined with DMA write — I6 atomicity 자연 보장) ``` #### Credit Injector (outbound) ``` Trigger: recv completion (my_tail 증가 후) Action: pack 16B credit packet → DMA vc_comm (또는 dedicated credit VC) Packet: {consumer_seq = my_tail, dst_rx_base_pa = my_rx_base_pa} Latency: 1 cycle to generate, then NoC traversal ``` #### Credit Receiver (inbound sideband) ``` Trigger: 16B credit packet arrival (from NoC) Extract: {consumer_seq, dst_rx_base_pa} Direction matching (ADR-0025 D3): for each dir: match = (peer_rx_base_pa[dir] == credit.dst_rx_base_pa) Update: peer_tail_cache[matched_dir] = max(peer_tail_cache, consumer_seq) Output: send_wake signal → deassert backpressure stall Latency: 1 cycle ``` ### D18. DMA Engine vc_comm IPCQ-aware Mode 기존 vc_comm 채널(D8)에 IPCQ flit 처리 모드를 추가한다. **Outbound**: 1. IPCQ Controller로부터 command 수신: `{src_addr, dst_addr, nbytes, sender_seq}` 2. TCM에서 src_addr read → DMA read buffer에 snapshot (standard DMA behavior) 3. Flit pack: data + piggyback metadata (sender_seq, dst_addr) 4. NoC fabric port에 inject 5. Fire-and-forget (completion 미대기) **Inbound**: 1. NoC로부터 IPCQ flit 수신 2. Terminal BW drain charge (`drain_ns = nbytes / bottleneck_bw`) 3. Slot write latency charge (backing memory tier) 4. **ATOMIC** (same pipeline stage, no stall insertion): - TCM write: data → slot address - Meta Extractor trigger: sender_seq + dst_addr → IPCQ Controller 5. Done **I6 atomicity 하드웨어 보장**: TCM write completion과 Meta Extractor trigger가 동일 pipeline stage에서 발생하므로 별도 synchronization이 불필요. 시뮬레이터의 "no SimPy yield between MemoryStore.write and IpcqMetaArrival put" (D9, I6)이 자연스럽게 보장된다. #### Data Snapshot Semantics DMA read buffer에 latch된 데이터는 src memory의 이후 수정에 영향받지 않는다. 이는 DMA standard read-then-write behavior이므로 추가 HW 불필요. #### Credit Virtual Channel (선택적) - **옵션 A**: vc_comm에 credit을 multiplexing (16B header-only flit으로 구분). - **옵션 B**: 3rd dedicated credit VC 추가 (strict priority > data). 옵션 B가 deadlock prevention에 유리하나, 16B credit의 BW 영향이 무시 가능하므로 옵션 A로도 충분. ### D19. Fabric Flit Format Extension ``` 일반 data flit (예: 512-bit): ┌──────────────────────────────────────────┐ │ [511:480] routing header (32b) │ │ [479:0] payload (480b = 60B) │ └──────────────────────────────────────────┘ IPCQ data flit (첫 flit에만 metadata 포함): ┌──────────────────────────────────────────┐ │ [511:480] routing header (32b) │ │ [511] ipcq_flag (1b) │ ← IPCQ vs normal DMA 식별 │ [510:509] vc_id (2b) │ │ [508:480] route + hop count │ │ [479:416] ipcq_metadata (64b) │ ← piggyback │ [479:448] sender_seq (32b) │ │ [447:416] dst_addr[31:0] (32b) │ ← direction matching용 │ [415:0] payload (416b = 52B) │ └──────────────────────────────────────────┘ 후속 flits: full 60B payload (metadata 없음) Credit-only flit (128-bit, header-only): ┌──────────────────────────────────────────┐ │ [127:96] routing header (32b) │ │ [127] credit_flag (1b) │ │ [95:64] consumer_seq (32b) │ │ [63:0] dst_rx_base_pa (64b) │ └──────────────────────────────────────────┘ ``` 첫 flit의 payload가 60B → 52B로 감소 (13% overhead). Multi-flit transfer에서는 후속 flit이 full payload이므로 대형 전송에서 overhead < 1%. ### D20. TCM IPCQ Slot Region Layout ``` TCM Memory Map (16MB): ┌─────────────────────────────┐ 0x000000 │ Kernel Working Memory │ │ (compute tensors) │ │ ~14MB │ ├─────────────────────────────┤ 0xE00000 │ IPCQ RX Buffers │ │ Dir N: slots × slot_size │ │ Dir S: slots × slot_size │ │ Dir E: slots × slot_size │ │ Dir W: slots × slot_size │ │ ~1MB │ ├─────────────────────────────┤ 0xF00000 │ IPCQ Metadata / Scratch │ │ ~1MB │ └─────────────────────────────┘ 0xFFFFFF ``` IPCQ region을 TCM의 상위 bank에 배치하여 compute access와의 bank conflict를 최소화한다 (Risk D22 참조). ### D21. 2nm Implementation Analysis #### Area Estimate | Module | Gate Count | Area (2nm est.) | Notes | |---|---|---|---| | QPair Register File | ~4.6K FF | 0.002 mm² | 576B flip-flops | | Slot Addr Gen + Backpressure | ~5K gates | 0.001 mm² | Combinational | | Meta Extractor + Credit Logic | ~3K gates | 0.001 mm² | 8× parallel comparators | | **IPCQ Controller subtotal** | **~12.6K** | **~0.004 mm²** | **PE 전체 대비 < 0.1%** | | DMA vc_comm 확장 | ~2K gates | 0.002 mm² | Flit pack/unpack | | **Total 변경분** | **~14.6K** | **~0.006 mm²** | | #### Timing | Path | Delay (2nm est.) | Target Clock | Margin | |---|---|---|---| | Backpressure (sub + cmp) | ~0.3 ns | 1 GHz (1 ns) | 3× | | Slot Addr Gen (mask + shift + add) | ~0.5 ns | 1 GHz | 2× | | Meta Extractor (8× range match) | ~0.4 ns | 1 GHz | 2.5× | | Credit Receiver (8× equality) | ~0.3 ns | 1 GHz | 3× | 모든 critical path가 1 cycle 이내. Timing closure 문제 없음. #### Power - Active: ~1 mW (register R/W + comparators, send/recv 동작 시) - Idle: leakage only - PE 전체 전력 대비 무시 가능 #### Constraints | 항목 | 제약 | 근거 | |---|---|---| | `n_slots` | **반드시 power-of-2** | mod → AND mask (1 gate). 임의 값은 divider 필요 (~10 cycles) | | `slot_size` | **power-of-2 권장** | mul → barrel shift. 임의 값은 multiplier 필요 | | TCM IPCQ region | **전용 bank 배치** | Compute access와 bank conflict 방지 | ### D22. Risk Assessment #### TCM Bank Conflict - **Risk**: IPCQ slot write와 compute read가 동일 bank 접근 시 stall - **Mitigation**: IPCQ region을 TCM 상위 address의 전용 bank에 배치 (D20) - **Cost**: TCM banking flexibility 소폭 감소 - **Severity**: Medium (성능 영향), Low (correctness 문제 아님) #### Credit Return Latency under Congestion - **Risk**: NoC 혼잡 시 credit return 지연 → sender backpressure stall - **Mitigation**: - Credit을 별도 VC로 분리 + strict priority (16B로 BW impact 미미) - 또는 n_slots를 넉넉히(8+) 설정하여 credit 지연을 buffer로 흡수 - **Severity**: Low (credit 16B는 congestion에 거의 기여하지 않음) #### Inter-Direction Ordering - **Risk**: 같은 PE에서 여러 방향으로 동시 send 시 순서 - **Mitigation**: Per-direction monotonic seq으로 충분. Inter-direction ordering은 kernel(소프트웨어) 책임 — 현재 시뮬레이터 모델과 동일 (D2 + D4) - **Severity**: Low (아키텍처 설계에 의해 해소) ### D23. HW Alternatives Considered #### Doorbell + Polling (전통적 방식) ``` Send: DMA write data → DMA write doorbell register at peer → peer polls doorbell Recv: Polling loop on doorbell, or interrupt-driven ``` | 장점 | 단점 | |---|---| | 단순한 HW (IPCQ controller 불필요) | 2번의 DMA transaction (data + doorbell) | | 기존 DMA 재사용 | Data/doorbell 사이 ordering 보장 필요 (fence) | | | Polling은 전력 낭비, interrupt는 latency overhead | **평가**: Piggyback 대비 latency 2-3× 증가. **불채택.** #### Hardware Message Queue (NVIDIA NVLink 스타일) ``` Send: CPU → HMQ에 descriptor push → HW가 peer HMQ로 자동 전달 Recv: HMQ에서 descriptor pop → data pointer 확인 ``` | 장점 | 단점 | |---|---| | CPU는 descriptor만 작성 | 별도 HMQ engine 필요 (~0.05 mm²) | | Descriptor/data 분리 → 유연 | DMA와 별개 datapath → area/power 중복 | | | Large tensor에는 결국 DMA 필요 | **평가**: CCL의 large tensor 패턴에서 DMA 필수이므로 HMQ + DMA 이중 구조는 면적 낭비. **불채택.** #### RDMA-style Completion Queue (CQ) ``` Send: DMA write → peer에 CQE 자동 생성 Recv: CQ poll/interrupt → data 위치 확인 ``` | 장점 | 단점 | |---|---| | InfiniBand/RoCE 성숙 모델 | CQ 관리 logic + CQE memory overhead | | Multi-tenant/isolation 용이 | CQE/data ordering 보장 추가 필요 | | | PE-to-PE CCL에는 over-engineered | **평가**: RDMA CQ는 host-facing NIC의 multi-tenant 격리에 적합. PE 간 단일 owner 환경에서는 불필요한 복잡성. **불채택.** #### Credit-in-Data Piggyback (v2 최적화 후보) 현재 설계에서 credit return은 별도 16B packet이다. Bidirectional 통신 패턴에서는 **reverse 방향 data flit에 credit을 합칠 수 있다.** ``` PE_A →E→ PE_B: data + sender_seq=3 PE_B →W→ PE_A: data + sender_seq=5 + credit_ack=4 ← credit이 data에 합쳐짐 ``` | 장점 | 단점 | |---|---| | Credit 전용 packet 제거 → NoC BW 절약 | Unidirectional 패턴에서는 fallback 필요 | | Bidirectional allreduce에서 credit latency → 0 | Flit header에 8B 추가 (overhead 미미) | | | Logic 복잡도 소폭 증가 | **평가**: 현재 설계의 우수한 최적화. Bidirectional allreduce에서 credit packet을 완전 제거 가능. Standalone credit fallback도 유지. **v2로 채택 권고.** ### Open HW Questions - IPCQ slot region size를 TCM의 몇 %까지 허용할 것인가? (현재 가정: ~1MB / 16MB = 6.25%) - Credit VC를 별도로 둘 것인가, vc_comm에 multiplexing할 것인가? (D18 참조) - Inter-SIP link에서의 flit format 호환성 검증 필요 - n_slots 최대값 제한? (8 directions × 8 slots × 64KB = 4MB → TCM의 25%) --- ## Non-goals - **호스트 collective**: `dist.all_reduce`가 데이터 이동을 직접 수행하는 모델은 본 ADR 범위 외. 본 ADR은 PE 커널 안에서 일어나는 통신만 다룬다. - **All-reduce 알고리즘**: ring/tree 등 알고리즘 자체는 별도 ADR (또는 커널 코드)에서 다룬다. 본 ADR은 인프라(IPCQ + VC)만 정의. - **Reliability/error handling**: send/recv 실패, link 장애 등은 다루지 않음. - **NoC arbiter 정밀 모델**: VC 간 dynamic contention은 첫 구현 범위 외 (D8). --- ## Open Questions - **VC arbitration 정확도**: 첫 구현은 deterministic chunk interleave + weighted round-robin. heavy contention 시나리오에서 실제보다 optimistic한 결과가 나올 수 있음. 정밀화 필요 시 별도 NoC arbiter component 도입을 검토. - **Credit return fast path BW 모델**: 첫 구현은 fast path가 fabric BW contention 모델에서 제외 (별도 lightweight wire 가정). 정밀화 필요 시 credit fast path를 별도 link로 모델링하거나, `credit_return_mode: piggyback` 옵션 추가. - **Ring buffer slot의 메모리 할당**: TCM/HBM/SRAM 어디에 두든 IPCQ가 알아야 할 metadata (base addr, slot_size, n_slots). init 시 호스트가 사이드밴드로 넣을지, fabric MmuMapMsg와 유사한 메시지로 넣을지 결정 필요. - **VC 간 BW 분할 default**: 균등 분할(50/50)인지, weighted(예: 80% compute, 20% comm)인지. ccl.yaml에 노출하되 default 값 결정 필요. - **Direction 개수**: 4방향(N/S/E/W) 고정인지, 6방향(+ Up/Down for 3D), 또는 가변 N개로 확장할지. 첫 구현은 4방향 고정. - **다중 channel 데이터 구조 (multi-tile aggregation)**: 한 collective에서 여러 tile을 fan-out 받는 경우 기존 round-robin recv로 충분한지, 별도 primitive(`tl.recv_all`)가 필요한지. - **Round-robin recv fairness**: 첫 구현은 last_polled_dir 인덱스 기반 weak fairness. 한 방향에 데이터가 항상 먼저 도착하면 starvation 가능. strict fairness가 필요하면 별도 fairness counter 추가. - **Deadlock detection 정밀화**: 첫 구현은 timeout 기반. 향후 wait-for graph 실시간 추적으로 deterministic deadlock detection 가능. --- ## Consequences ### 긍정적 - PE 간 직접 통신 가능 → CCL 커널 작성 가능 - 호스트는 launch만, 동기화는 PE 안에서 → 단순한 호스트 코드, 강한 compute/comm overlap - VC를 통해 HoL blocking 제거 → collective latency가 compute traffic에 block되지 않음 - Buffer 위치/backpressure 모드를 init 파라미터로 선택 가능 → 벤치마크 가능 - 4-방향 logical neighbor → 호스트가 ring/mesh/tree 등 알고리즘 자유롭게 매핑 ### 부정적 - 컴포넌트 1개 신규 추가 (PE_IPCQ), PE_DMA 재설계 (VC 추가) - IPCQ 메모리 (8 ring × slot_size × n_slots) 만큼 PE-local 메모리 사용 - VC arbitration 모델이 first-order approximation이므로 heavy contention 시나리오에서 실제 HW보다 약간 optimistic한 latency 결과 가능 (D8 한계) - VC chunk-level 인터리브로 PE_DMA 구현이 더 복잡해짐