# ADR-0020: 2-Pass 데이터 실행 모델 (타이밍 / 데이터 분리) ## Status Accepted ## Context 현재 시뮬레이션은 **타이밍만** 모델링한다. `tl.load()`, `tl.composite(op="gemm")` 등은 SimPy latency를 생성하지만, 실제 텐서 데이터를 읽거나 연산하지 않는다. ### 필요한 기능 1. HBM/TCM/SRAM에 실제 데이터를 저장하고 읽을 수 있어야 한다 2. PE_GEMM, PE_MATH가 실제 행렬 연산을 수행하고 결과를 검증할 수 있어야 한다 3. 시뮬레이션 성능 저하를 최소화해야 한다 ### 제약 조건 - SimPy는 single-thread 이벤트 루프 — numpy matmul을 안에서 하면 전체가 block - 컴포넌트는 교체 가능해야 한다 (ADR-0015) — 프레임워크 요구사항이 구현에 침투하면 안 됨 - 벤치마크 커널은 명령형 코드(tl.load → tl.composite → tl.wait) — 같은 코드를 재사용해야 함 - 커널 함수는 plain Python function으로 유지해야 한다 (generator/async 변환 불가) ### 설계 탐색 결과 | Option | 방식 | 판정 | |--------|------|------| | SimPy 내 직접 실행 | GEMM을 SimPy 안에서 numpy 호출 | 탈락: single-thread block | | SimPy + ThreadPool | future.submit → timeout → result() | 탈락: back-to-back 요청 시 result()에서 block | | Symbolic + lazy | 메타데이터만 추적, 나중에 실행 | 탈락: control-flow dependent 읽기 처리 곤란 | | **2-pass (채택)** | Phase 1: 타이밍, Phase 2: 데이터 | 완전 분리, 성능 영향 없음 | --- ## Decision ### D1. 2-Pass 실행 모델 — Phase 0 제거 기존의 3단계(Phase 0 → Phase 1 → Phase 2)를 **2단계로 통합**한다. 기존: ``` Phase 0: 커널 → PeCommand 리스트 (데이터 없음, 분기 불가) Phase 1: PeCommand 리스트를 SimPy replay (타이밍만) ``` 변경: ``` Phase 1 (타이밍): 커널 + SimPy 통합 실행 — greenlet 기반 - 메모리 읽기/쓰기: SimPy 타이밍 + MemoryStore 실제 데이터 - 연산 (GEMM/Math): SimPy 타이밍 + op_log 기록 (실제 연산은 Phase 2) - dynamic control flow 가능 (tl.load가 실제 데이터 반환) Phase 2 (데이터): op_log 기반 실제 연산 실행 — SimPy 외부, 병렬 가능 ``` 본 ADR은 **메모리 연산에 한해 Phase 1을 data-aware로 확장**한다. Phase 1은 latency/BW 병목 분석 + 메모리 데이터 추적, Phase 2는 GEMM/Math 연산 정합성 검증. Phase 2는 optional — 타이밍만 필요하면 Phase 1만 실행. ### D2. Op Log 기록 — ComponentBase hook op_log 기록은 **컴포넌트 베이스 클래스의 hook**으로 수행한다. 개별 컴포넌트 구현을 수정하지 않는다. ```python class ComponentBase: def _on_process_start(self, env, msg): if self._op_logger and getattr(msg, 'data_op', False): self._op_logger.record_start(env.now, self.node.id, msg) def _on_process_end(self, env, msg): if self._op_logger and getattr(msg, 'data_op', False): self._op_logger.record_end(env.now, self.node.id, msg) ``` `_forward_txn()` 에서 `run()` 전후로 hook을 호출한다. `_op_logger`는 optional — 없으면 오버헤드 제로. **hook 시점 정의**: | 시점 | 의미 | |------|------| | `t_start` | 컴포넌트가 해당 msg의 **service를 시작**한 시점 (`run()` 진입 직전) | | `t_end` | 컴포넌트의 **내부 service가 완료**된 시점 (`run()` 반환 직후) | link traversal latency는 t_start/t_end에 포함되지 않는다. link latency는 발신 컴포넌트의 t_end와 수신 컴포넌트의 t_start 차이로 관측된다. ### D3. Greenlet 기반 커널 실행 — Phase 0 제거 기존 Phase 0 (커널 → PeCommand 리스트)를 제거하고, **greenlet**을 사용하여 커널과 SimPy를 협력적으로 interleave 실행한다. #### 동작 원리 greenlet은 협력적 context switch를 제공하는 C 확장이다. 커널(child greenlet)이 `tl.load()` 등을 호출하면 SimPy 루프(parent greenlet)로 switch하여 타이밍 시뮬레이션을 수행하고, 완료 후 실제 데이터와 함께 커널로 돌아온다. ``` SimPy 루프 (parent greenlet) 커널 (child greenlet) ───────────────────────── ────────────────────── g.switch() ─────────────────────────→ 커널 시작 a = tl.load(ptr, ...) 내부: parent.switch(DmaReadCmd) cmd = DmaReadCmd ←────────────────── (커널 일시정지) yield DmaReadMsg(...) yield env.timeout(dma_latency) data = memory_store.read(...) g.switch(data) ─────────────────────→ (커널 재개) a = data ← 실제 numpy array if a[0][0] > 0.5: ← 분기 가능 ... ``` 커널은 **plain Python function**으로 유지된다. greenlet switch는 `tl.load()`, `tl.store()` 등의 **내부 구현에만** 존재한다. #### KernelRunner — 프레임워크 레이어 greenlet 루프는 PE_CPU 컴포넌트가 아니라 프레임워크 레이어인 **KernelRunner**에 위치한다. ```python # KernelRunner (프레임워크 — greenlet ↔ SimPy 연결) class KernelRunner: def run(self, env, kernel_fn, args, store): g = greenlet(self._run_kernel) cmd = g.switch(kernel_fn, args) while cmd is not None: if isinstance(cmd, DmaReadCmd): yield from self._dispatch_dma(env, cmd) data = store.read(cmd.src_addr, cmd.shape, cmd.dtype) cmd = g.switch(data) # 실제 데이터와 함께 재개 elif isinstance(cmd, GemmCmd): yield from self._dispatch_gemm(env, cmd) cmd = g.switch() # 재개 (데이터 없음) elif isinstance(cmd, DmaWriteCmd): store.write(cmd.dst_addr, cmd.data) # visibility = issue 시점 yield from self._dispatch_dma(env, cmd) # timing만 반영 cmd = g.switch() # PE_CPU (컴포넌트 — 간단하게 유지, greenlet을 모름) def _execute_kernel(self, env): runner = KernelRunner(self.ctx) yield from runner.run(env, kernel_fn, args, store) ``` **Op logging single source of truth**: KernelRunner는 op_log에 직접 기록하지 않는다. 모든 op logging은 **ComponentBase hook (_on_process_start/end)만** 담당한다. KernelRunner가 `_dispatch_gemm()` 등으로 컴포넌트에 메시지를 전달하면, 컴포넌트 베이스 클래스의 hook이 자동으로 기록한다. **레이어 분리**: - **커널 코드**: plain function, greenlet 존재를 모름 - **TLContext**: `tl.load()` 내부에서 `parent.switch(cmd)` 호출 - **KernelRunner**: greenlet ↔ SimPy 연결, MemoryStore 읽기/쓰기 처리. **logging 안 함**. - **ComponentBase hook**: op_log 기록의 유일한 경로 - **PE_CPU**: KernelRunner를 호출만 함, 컴포넌트로서 교체 가능 #### 메모리 읽기/쓰기 vs 연산의 처리 차이 | 연산 | Phase 1에서 | Phase 2에서 | |------|------------|------------| | `tl.load()` | SimPy 타이밍 + MemoryStore read → **실제 데이터 반환** | — | | `tl.store()` | SimPy 타이밍 + MemoryStore write → **실제 기록** | — | | `tl.composite(gemm)` | SimPy 타이밍 + **op_log 기록만** | numpy 실제 연산 | | `tl.dot()` / math ops | SimPy 타이밍 + **op_log 기록만** | numpy 실제 연산 | 메모리 읽기/쓰기는 Phase 1에서 즉시 처리 (numpy slice, 빠름). GEMM/Math 연산은 Phase 2에서 batch 실행 (성능 분리). #### Store Visibility Rule `tl.store()`는 **issue 시점에 MemoryStore에 즉시 반영**된다 (visibility = issue). SimPy DMA 타이밍은 이후 별도로 시뮬레이션된다. 이는 timing과 visibility를 의도적으로 분리한 것이다: - **visibility**: MemoryStore에 반영되는 시점 = `store.write()` 호출 시 - **timing**: SimPy에서 DMA latency가 완료되는 시점 이 분리로 dynamic control flow에서 store 직후 load가 최신 데이터를 볼 수 있다. #### Result Handle Semantics `tl.composite()`(sync/async)는 결과 tensor를 참조하는 **handle**을 반환한다. Phase 1에서의 핵심 계약: 1. **모든 compute handle은 Phase 1에서 항상 pending 상태로 간주한다.** 2. `tl.wait(handle)`은 **timing synchronization만 표현**하며, handle을 ready로 만들지 않는다. 3. handle의 실제 결과 데이터 접근(`handle.data`, element access, numpy conversion 등)은 **Phase 2에서만 가능**하다. 4. 따라서 Phase 1에서 **compute-result 기반 control flow는 지원하지 않는다.** 5. 반면 `tl.load()`는 Phase 1에서 실제 데이터를 반환하므로, **memory-read 기반 control flow는 지원 가능**하다. | handle 상태 | Phase | 허용 동작 | |------------|-------|----------| | pending | Phase 1 | `tl.wait(handle)` — timing 동기화만 | | pending | Phase 1 | handle을 `tl.store()`의 대상으로 전달 (logical destination 연결만, payload는 Phase 2) | | pending | Phase 1 | **데이터 접근 불가** — 값 기반 분기 불가 | | ready | Phase 2 | 실제 numpy 데이터 접근, 검증 | 이 제약은 의도적이다. Phase 1에서 연산을 실행하면 SimPy single-thread가 block되어 2-pass 분리의 존재 이유가 사라진다. #### Phase 1 Materialization — Future Extension 향후 소형 연산(scalar, 작은 reduction)에 대해 Phase 1 eager execution이 필요한 경우, `materialized_in_phase1: bool` 플래그를 op record에 추가하여 선택적 materialization을 지원할 수 있다. 현재 범위에서는 구현하지 않는다. ### D4. data_op 플래그 — 메시지 자기 선언 로깅 대상은 메시지 타입이 아니라 메시지 인스턴스의 `data_op` 속성으로 결정한다. 프레임워크가 메시지 타입을 하드코딩하지 않는다. ```python class MsgBase: data_op: bool = False # 기본: 로깅 안 함 class DmaReadCmd(MsgBase): data_op = True # 메모리 이동 → 로깅 class GemmCmd(MsgBase): data_op = True # 연산 → 로깅 class MathCmd(MsgBase): data_op = True # 연산 → 로깅 ``` 새 메시지 타입(예: IpcqMsg) 추가 시 `data_op = True`만 설정하면 프레임워크 코드 수정 없이 자동 로깅된다. ### D5. Op Log 구조 #### op 분류 체계 2단계로 분류한다: | 레벨 | 필드 | 역할 | |------|------|------| | `op_kind` | `memory` \| `gemm` \| `math` | executor dispatch 기준 | | `op_name` | `dma_read` \| `dma_write` \| `gemm_f16` \| `exp` \| `add` \| `sum` 등 | 구체 연산 식별 | #### OpRecord 정의 ```python @dataclass class OpRecord: t_start: float # SimPy 시각 (ns) — service 시작 t_end: float # SimPy 시각 (ns) — service 완료 component_id: str # e.g. "sip0.cube0.pe0.pe_gemm" op_kind: str # "memory" | "gemm" | "math" op_name: str # 구체 연산명 params: dict # 연산별 파라미터 (아래 참조) dependency_ids: list[int] # 현재는 in-memory record index 기반, 향후 stable op_id로 대체 가능 ``` #### dependency_ids 생성 규칙 `dependency_ids`는 **optional**이며, 기본적으로 executor는 주소 기반 dependency 추론을 수행한다 (D6 참조). 정확한 실행 순서가 필요한 경우에만 명시적으로 설정한다: - **기본 (address-based inference)**: executor가 read/write set을 분석하여 RAW/WAW/WAR 의존성을 자동 추론. 대부분의 경우 이것으로 충분. - **명시적 설정**: TLContext 또는 command 생성 단계에서 logical dependency가 주소로 표현되지 않는 경우에 설정. 예: completion handle 기반 동기화 — handle dependency는 메모리 주소가 아니라 논리적 완료 순서에 의존하므로 address inference로 잡히지 않는다. #### op_log ordering op_log는 `t_start` 기준으로 **stable ordering**을 유지한다. 동일 `t_start`의 record들은 insertion order를 보존한다. #### params 상세 **memory (dma_read / dma_write)**: ```python { "src_addr": int, # source 주소 (byte) "dst_addr": int, # destination 주소 (byte) "nbytes": int, # 전송 크기 "src_space": str, # "hbm" | "tcm" | "sram" "dst_space": str, # "hbm" | "tcm" | "sram" } ``` **gemm**: ```python { "src_a_addr": int, # operand A 주소 "src_b_addr": int, # operand B 주소 "dst_addr": int, # output 주소 "shape_a": tuple, # e.g. (128, 256) "shape_b": tuple, # e.g. (256, 128) "shape_out": tuple, # e.g. (128, 128) "dtype_in": str, # e.g. "f16" "dtype_acc": str, # accumulation dtype, e.g. "f32" "dtype_out": str, # output dtype, e.g. "f16" "transpose_a": bool, "transpose_b": bool, "layout_a": str, # "row_major" | "col_major" "layout_b": str, "layout_out": str, "addr_space": str, # "tcm" (GEMM operand는 항상 TCM) } ``` **math**: ```python { "op": str, # "exp" | "add" | "sum" | "where" | ... "input_addrs": list[int], # operand 주소 목록 "input_shapes": list[tuple], "dst_addr": int, "shape_out": tuple, "dtype": str, "axis": int | None, # reduction axis "addr_space": str, # "tcm" } ``` ### D6. Phase 2 Executor Phase 2는 SimPy 밖에서 op_log를 실행한다. ```python class DataExecutor: def __init__(self, op_log: list[OpRecord], initial_store: MemoryStore): self.store = initial_store # Phase 1의 MemoryStore snapshot을 입력으로 받는다 def run(self): for t, ops in groupby(op_log, key=lambda o: o.t_start): batch = list(ops) independent, sequential = self._classify(batch) self._execute_parallel(independent) self._execute_sequential(sequential) ``` **병렬 실행 판정**: 같은 `t_start`의 op들은 **병렬 후보**로 간주한다. 실제 병렬 실행 여부는 executor가 다음 기준으로 판정한다: - read/write 주소 범위 겹침 여부 (WAW, RAW, WAR 충돌 검사) - `dependency_ids`에 명시된 선행 op 완료 여부 주소 범위가 겹치지 않고 명시적 의존성이 없는 op들만 병렬 실행한다. **배치 최적화**: 동일 op_name이며 **shape, dtype, layout, transpose flag가 모두 동일한** 독립 op들만 batching 대상이 된다. 예: 여러 PE의 동일 shape GEMM → `np.matmul(a_batch, b_batch)` 한 번으로 묶음. CPU에서도 BLAS 효율 향상, GPU에서는 launch overhead 절감. **Phase 2 실행 순서 보장**: Phase 2는 데이터 도착 시점을 고려하지 않으며, dependency (주소 기반 추론 + 명시적 dependency_ids)를 통해서만 실행 순서를 보장한다. ### D7. Memory Store `MemoryStore`는 논리적으로 byte-addressable semantics를 따르며, 현재 구현은 **tensor-granular storage** (addr → numpy ndarray 매핑)를 사용한다. ```python class MemoryStore: def write(self, space: str, addr: int, data: np.ndarray) -> None: ... def read(self, space: str, addr: int, shape: tuple, dtype: str) -> np.ndarray: ... ``` **내부 저장 포맷: numpy ndarray** MemoryStore는 텐서를 **numpy ndarray**로 저장한다. | 후보 | store/load 속도 | Phase 2 연산 | 판정 | |------|----------------|-------------|------| | **numpy ndarray** | 즉시 (참조 전달, 복사 없음) | `np.matmul` 바로 사용 | **채택** | | bytearray | memcpy 필요 | `np.frombuffer` 변환 필요 | 탈락 | | torch tensor | 즉시 | torch 연산 가능 | GPU 최적화 시만 사용 | - write: numpy array를 **참조 저장** (복사 없음) → Phase 1 오버헤드 = dict lookup 1회 - read: numpy array를 **참조 반환** (복사 없음) - 동일 addr에 재 write 시 기존 array를 **tensor 단위로 덮어쓴다** (partial overwrite 미지원) - dtype은 numpy native 사용 (`np.float16`, `np.float32`, `np.bfloat16` 등) - byte-level access가 필요한 경우 `.view(np.uint8)` 로 변환 - Phase 2에서 GPU batch 최적화 시 numpy → torch tensor 변환은 executor가 담당 **read/write contract**: - read/write는 **contiguous tensor** 기준이다. non-contiguous stride view가 필요한 경우 별도 copy op으로 표현한다. - 일반 benchmark path에서는 producer/consumer dtype 일치를 기대한다. reinterpret cast는 low-level memory validation 또는 특수 테스트 케이스를 위한 permissive behavior이다. - addr은 byte-aligned이며, 최소 alignment = dtype 크기. - dtype mismatch (write와 다른 dtype으로 read)는 reinterpret cast로 처리한다. shape 불일치 시 nbytes 기준으로 검증하고, 불일치하면 error. - 정합성 기준은 주소 범위 기반 read/write semantics를 따른다. - 구현 최적화로 tensor object cache를 둘 수 있지만, canonical state는 byte-addressable storage이다. - deploy 시점에 호스트가 초기 텐서 데이터를 주입한다. ### D8. 벤치마크 커널 코드 벤치마크의 **사용자 코드 API는 변경하지 않는다**. `tl.load()`, `tl.composite()`, `tl.store()` 등의 호출 인터페이스는 유지. 단, 내부 command/message schema는 Phase 2 실행에 필요한 metadata를 포함하도록 확장될 수 있다 (예: dtype_acc, transpose 등 추가 필드). ### D9. 컴포넌트 변경 없음 개별 컴포넌트 구현(PE_GEMM, PE_DMA, HBM_CTRL 등)은 수정하지 않는다. op_log 기록은 ComponentBase hook의 책임이다. 커스텀 컴포넌트 교체 시 타이밍 모델만 교체되며, Phase 2 데이터 실행은 영향받지 않는다. ### D10. Phase 2는 Optional ```python engine = GraphEngine(graph) engine.run(benchmark) # Phase 1: 타이밍만 result = engine.get_timing_result() if verify_data: executor = DataExecutor(engine.op_log) # Phase 2: 데이터 executor.run() executor.verify(expected_output) ``` 타이밍 분석만 필요하면 Phase 2를 건너뛴다. op_logger를 비활성화하면 Phase 1 성능도 기존과 동일. ### D11. Verification Contract 기본 검증은 **최종 output tensor**를 reference backend(numpy)와 비교한다. dtype별 tolerance 정책: | dtype | 비교 방식 | tolerance | |-------|----------|-----------| | f32 | `np.allclose` | rtol=1e-5, atol=1e-5 | | f16 | `np.allclose` | rtol=1e-3, atol=1e-3 | | bf16 | `np.allclose` | rtol=1e-2, atol=1e-2 | | int 계열 | `np.array_equal` | exact | - 기본 모드: 최종 output만 비교 (end-to-end correctness) - 디버그 모드: intermediate tensor도 op 단위로 비교 가능 (MemoryStore snapshot at each op boundary) --- ## Non-goals - **Compute-result-based control flow**: 지원하지 않는다. 모든 compute handle은 Phase 1에서 pending 상태이며, `wait()`는 timing synchronization만 표현하고 data readiness를 의미하지 않는다. Phase 1에서 `handle.data` 접근, element access, truth-value evaluation은 **error로 처리**한다. 메모리 데이터 기반 분기(`tl.load()` 결과)는 greenlet으로 지원된다. Phase 1 materialization은 future extension (D3 참조). - **Cycle-accurate overlap reconstruction**: Phase 2에서 Phase 1의 실행 시간 overlap을 정확히 재현하지 않는다. Phase 2는 데이터 정합성만 검증한다. - **GPU kernel compilation**: Phase 2의 GEMM/Math는 numpy/torch 호출이며, 실제 하드웨어 PE의 마이크로아키텍처를 재현하지 않는다. ## Open Questions - **Aliasing / slice view**: 동일 backing storage를 참조하는 slice/view를 MemoryStore에서 어떻게 표현할지 (stride-based view vs copy semantics) - **IPCQ/descriptor read 일반화**: PE-to-PE 통신을 memory op으로 완전히 일반화할지, 별도 op_kind를 둘지 - **Op log streaming**: 대규모 시뮬레이션에서 op_log 메모리 사용량 관리 (in-memory list vs disk-backed streaming) - **Fused operation**: tl.composite의 tiled pipeline (READ→COMPUTE→WRITE)을 하나의 fused op record로 기록할지, 개별 op으로 분리할지 - **Math op schema 일반화**: 현재 math params는 단순 구조이나, broadcasting rule, input별 dtype, keepdims, scalar/immediate operand, where/mask 표현 등 일반화가 필요할 수 있음 - **Op record 식별자**: 현재 dependency_ids는 in-memory list index 기반이며, streaming/disk-backed mode 도입 시 stable op_id로 대체 필요 - **Phase 1 materialization policy**: D3의 Future Extension 참조. 허용 시 해당 op의 Phase 2 처리 방식 (skip / verify / recompute) 정의 필요 --- ## Consequences ### 긍정적 - SimPy 시뮬레이션 성능 영향 최소 (op_log append만 추가) - Phase 2에서 멀티스레드/GPU 자유롭게 사용 가능 - 컴포넌트 교체 자유도 유지 (ADR-0015 설계 철학 보존) - 벤치마크 사용자 코드 API 변경 불필요 - 새 메시지 타입 추가 시 data_op 플래그만 설정 - greenlet으로 Phase 0 제거 — 메모리 데이터 기반 dynamic control flow 지원 - `tl.load()`가 실제 데이터를 반환하므로 커널 디버깅 용이 ### 부정적 - op_log 메모리 사용량 (대규모 시뮬레이션 시) - Phase 2 실행 시간은 텐서 크기에 비례 (대형 GEMM) - pending handle (연산 미완료) 기반 동적 분기 불가 (연산은 Phase 2에서 실행, Phase 1에서 결과 값 미확정). 메모리 데이터 기반 분기는 greenlet으로 지원된다. - greenlet C 확장 의존성 추가 (pip install greenlet)