kernbench2/docs/adr/ADR-0020-data-execution-two-pass.md

# ADR-0020: 2-Pass 데이터 실행 모델 (타이밍 / 데이터 분리)

## Status

Proposed

## Context

현재 시뮬레이션은 **타이밍만** 모델링한다.
`tl.load()`, `tl.composite(op="gemm")` 등은 SimPy latency를 생성하지만,
실제 텐서 데이터를 읽거나 연산하지 않는다.

### 필요한 기능

1. HBM/TCM/SRAM에 실제 데이터를 저장하고 읽을 수 있어야 한다
2. PE_GEMM, PE_MATH가 실제 행렬 연산을 수행하고 결과를 검증할 수 있어야 한다
3. 시뮬레이션 성능 저하를 최소화해야 한다

### 기존 커널 실행 구조의 한계

현재 커널 실행은 3단계로 분리되어 있다:

```
Phase 0: TLContext에서 커널 함수 실행 → PeCommand 리스트 생성 (SimPy 밖, 데이터 없음)
Phase 1: PE_CPU가 PeCommand 리스트를 SimPy로 replay (타이밍만)
```

Phase 0에서 커널이 **전부 실행 완료**된 후에야 SimPy가 시작된다.
`tl.load()`는 TensorHandle(placeholder)을 반환하므로 실제 데이터에 접근할 수 없다.
따라서 데이터 값에 따른 분기(dynamic control flow)가 불가능하다.

본 ADR은 이 한계를 **메모리 연산에 한해** 해소한다 (D1, D3 참조).

### 제약 조건

- 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)

---

## 영향받는 파일

| 파일 | 변경 |
|------|------|
| `src/kernbench/components/base.py` | `_on_process_start/end` hook 추가 |
| `src/kernbench/common/pe_commands.py` | `data_op = True` 추가, metadata 필드 확장 |
| `src/kernbench/sim_engine/op_log.py` | 신규: OpRecord, OpLogger |
| `src/kernbench/sim_engine/data_executor.py` | 신규: DataExecutor, MemoryStore |
| `src/kernbench/sim_engine/engine.py` | op_logger 주입 (optional) |
| `src/kernbench/triton_emu/tl_context.py` | `tl.load()` 등 내부에서 greenlet switch 호출 |
| `src/kernbench/triton_emu/kernel_runner.py` | 신규: KernelRunner (greenlet ↔ SimPy 연결) |
| `src/kernbench/components/builtin/pe_cpu.py` | Phase 0 제거, KernelRunner 호출로 변경 |
| `pyproject.toml` | greenlet 의존성 추가 |

컴포넌트 구현 파일 (pe_gemm.py, pe_dma.py, hbm_ctrl.py 등): **변경 없음**
벤치마크 커널 (benches/*.py): **사용자 API 변경 없음**