Files
kernbench2/docs/adr-ko/ADR-0064-perf-cpu-issue-cost-model.md
ywkang 2d8271c981 perf(cost-model): D8 single-op-cmd fast-path (FIXED=8 single-op / 40 composite)
ADR-0064 D8 broadened from "DMA fast-path" to "single-op-cmd fast-path":
every single-op command (DmaRead/DmaWrite/Gemm/Math/Copy) now pays the
lighter FIXED=8; only CompositeCmd keeps the 40-cycle control-path FIXED
(it alone needs scheduler plan generation + per-tile RW-hazard tracking +
completion wiring). Renamed knob fixed_per_dma_cmd_cycles ->
fixed_per_single_op_cmd_cycles. dispatch_cycles now branches on
"is CompositeCmd" rather than enumerating DMA types.

Term choice: "single-op" (not "atomic", which read as sync/async) — the
axis is composition (one engine op vs fused multi-op plan), orthogonal to
timing. single-op <-> composite.

Tests: test_pe_cost_model.py updated to the single-op surface (defaults,
fast-path over all 5 single-op cmd types, composite general path, yaml
override). All green.

Recalibrated tests/attention/test_gqa_decode_opt2.py
::test_opt3_dispatch_exceeds_opt2 — NOT a regression: D8 makes single-op
cmds 5x cheaper, so opt2's two-composite fusion win over opt3's many
single-ops narrowed from pre-D8 ~3.7x to ~1.87x (opt3=224 > opt2=120).
The CPU-offload invariant (opt2 cheaper) still holds; only the model-
dependent ">2x" constant was over-fit to the old uniform-40 model. Gate
now: direction + >1.5x margin (matches sibling R-sweep test's stated
"absolute ratio informative-only" philosophy).
NOTE for review: ADR-0065's "2x CPU-offload win" headline may want a
refresh to reflect the post-D8 ~1.87x — left to user (architectural doc).

Full regression: 826 passed, 1 skipped (tests/ excl. tests/gemm).

--- Remaining work (resume here if interrupted) ---
5. Re-run scripts/paper/paper_plot_gemm_async_vs_composite.py with new
   cost model; verify async-tiled dispatch overhead drops (~4576ns ->
   ~1536ns expected) and the composite-vs-async-tiled gap narrows from
   the prior ~6.3x at K=3072.
6. Copy regenerated gemm_composite_vs_async_tflops.png to
   docs/report/1H-codesign-paper/figures/.
7. Paper §3.4 (03-gemm.tex sec:gemm-vs-async): finish naive->async-full /
   chunked->async-tiled rename AND reframe FIXED_DMA wording to single-op
   vs composite (currently still says "lighter FIXED for DMA descriptors,
   FIXED_DMA=8"). Table 2 (02-platform) + §2 dispatch prose already done.
8. Paper §3.4 K=3072 corner para + mechanism #3: update dispatch breakdown
   to new model (96 DMA*8 + 95 single-op*8 ≈ 1.5us vs old 4.6us); update
   headline ratio if it changed.
9. Rebuild docs/report/1H-codesign-paper/build/main.pdf (tectonic) +
   verify via pdftotext.
10. Then this is the bench-harness + paper commits (Groups 2 & 3).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
2026-06-17 15:04:53 -07:00

435 lines
21 KiB
Markdown
Raw Permalink Blame History

This file contains ambiguous Unicode characters
This file contains Unicode characters that might be confused with other characters. If you think that this is intentional, you can safely ignore this warning. Use the Escape button to reveal them.
# ADR-0064: 구조적 CPU dispatch cost 모델 (`logical_bytes` + FIXED + R)
## Status
Accepted (Revision 2)
> **ADR-0060** (AHBM GQA Fused Attention) 와 **ADR-0065** (flat-ops
> composite + 첫 stateful recipe) 의 보조 ADR. 그 hybrid 의 핵심 — "GEMM 은
> `tl.composite` 로, softmax merge 는 커널에서" — 은 **PE_SCHEDULER 로
> tiling 을 offload 해서 CPU 가 굵은 descriptor 만 issue 하고 앞서가며
> 엔진을 포화시킴** 으로 이깁니다. 그 win 이 현재 시뮬레이터에선 보이지
> 않습니다 (per-op CPU issue cost = 0).
>
> Revision 2 는 원안의 **op-type calibration 표** 를 **logical_bytes 기반
> 구조 공식** 으로 대체합니다 — op-type 별 calibration 불필요, 새 op_kind
> 도 자동 적용.
## Context
### 현재 코드
- 모든 `tl.*` op 가 cmd emit 전에 `_emit_dispatch_overhead()` 호출
(`tl_context.py:196-212`) → `dispatch_cycles > 0` 일 때만
`PeCpuOverheadCmd(cycles=dispatch_cycles)` 발행. knob 은 op kind 무관
**uniform** 이며 두 실행 경로 모두 **0** 하드코딩
(`pe_cpu.py:101` greenlet, `:195` replay).
- ⇒ 명령 issue (descriptor 구성 + scheduler 큐 push) 가 PE_CPU 에서
**0 ns** 비용.
- `PeCpuOverheadCmd` 는 PE_CPU 에서 `yield env.timeout(cmd.cycles)`
소비 (`kernel_runner.py:131-132`).
### uniform-zero 가 hybrid 에 부적합한 이유
ADR-0060 §1 의 핵심: composite 1 개가 `N_tiles` 분량의 GEMM tiling 을
offload → CPU 는 `O(1)` 개의 굵은 cmd 만 issue, primitive 경로는
`O(N_tiles × ops/tile)`. issue cost = 0 이면 모델은:
- primitive 경로의 *CPU 가 push 를 못 따라가 엔진 idle* 을 못 보임
- composite 의 *구성 비용 ≫ primitive 1 개이지만 ≪ 대체된 primitive 들의 합* 을 못 보임
### Revision 1 의 op-type calibration 이 과한 이유
원안은 `cost_table[kind]` (kind = `composite`, `load`, `dot`, `math`, …)
형태였습니다. 비용:
- kind 마다 값 필요 (calibration ≥ |kinds|)
- 새 kind 추가 시마다 entry 추가
- 그런데 잡으려던 *ratio* — "composite ≫ primitive, ≪ 대체된 primitive 들의 합" —
*op kind* 가 아니라 *cmd 가 들고 있는 필드 수* 의 함수.
cmd 의 *byte 발자국* 이 자연 proxy: N OpSpec composite 는 1 OpSpec primitive 의 ~N 배 bytes.
*고정* 부분 (큐 tail 업데이트, completion 등록, MMIO-급 latency) 은 cmd 마다.
둘 합치면: `FIXED + bytes × R`.
### 이 모델이 실제로 노출하는 것
**1차** 신호는 per-cmd FIXED 비용에 의한 **command-count 감소**.
**byte** 항은 *비정상적으로 큰* composite 가 공짜로 보이지 않도록 막는
2차 보정. 현실적 on-die 큐 대역폭 (16 B/cycle, D3) 에서 FIXED 가
decode opt2 의 opt3 (≈10 cmds/tile) vs opt2 (≈2 cmds/tile) 의 dispatch
cost 차이 중 ≥85% 차지.
이 framing 이 모델의 한계도 정의: 단일 composite 가 무한히 커지면 byte
항만으로는 HW 가 받지 못할 정도의 fused command 를 모델이 보상하는 것을
막지 못함 — 그래서 **D7** 의 descriptor 크기 cap 이 필요.
## Decision
### D1. 구조적 dispatch cost 공식
PE_SCHEDULER 로 가는 모든 PE command 가 PE_CPU 에서 dispatch cycles 소모:
```
dispatch_cycles(cmd) = FIXED_PER_CMD + cmd.logical_bytes × R
```
- `FIXED_PER_CMD` (cycles/cmd): 큐 tail 업데이트, MMIO-급 RTT,
completion-event 등록 — cmd 크기와 무관 고정 비용.
- `R` (cycles/byte): scheduler 큐에 cmd 를 직렬화하는 큐-write 대역폭.
- `cmd.logical_bytes` (int): cmd 의 *HW-논리* byte 수 — D2 룰로 계산,
Python `sys.getsizeof` 가 아님.
PE_CPU 는 기존 hook 그대로 `PeCpuOverheadCmd(cycles=dispatch_cycles(cmd))`
를 dispatch 전에 발행 — 사이클 값만 변경.
### D2. `logical_bytes` 룰
각 PE command dataclass 가 `logical_bytes: int` (property) 노출. 룰
(HW 친화적, Python overhead 무시):
| 필드 종류 | bytes |
|---|---|
| cmd framing (cmd 종류 discriminator + completion id 참조) | 4 |
| Opcode (op kind enum) | 1 |
| Enum (scope 등) | 1 |
| `TensorHandle` 참조 (address only — shape/dtype 은 descriptor table 가정) | 8 |
| Scalar (int/float) | 4 |
| Tuple length marker | 1 |
`CompositeCmd``ops``rw_handles` 재귀 합산:
```python
@property
def logical_bytes(self) -> int:
return (
4 # framing
+ 1 + sum(op.logical_bytes for op in self.ops)
+ 1 + 8 * len(self.rw_handles)
)
```
`OpSpec`:
```python
@property
def logical_bytes(self) -> int:
return (
1 + 1 # opcode + scope
+ 1 + 8 * len(self.operands) # named operand handles
+ (8 if self.out is not None else 0) # out handle
+ 1 + sum(_extra_bytes(v) for v in self.extra.values())
)
def _extra_bytes(v) -> int:
"""OpSpec.extra 값의 타입-aware byte 카운트."""
if isinstance(v, bool): return 1
if isinstance(v, (int, float)): return 4
if isinstance(v, (tuple, list)): return 1 + 4 * len(v) # shape, axes, …
if isinstance(v, str): return 1 # opcode-like tag
return 4 # default scalar
# extra 의 예시 타입:
# m, k, n int → 4 each
# reduce_axis int → 4
# shape=(64, 64) tuple → 1 + 8 = 9
# factor=1.0 float → 4
```
(`DmaReadCmd`, `MathCmd` 등도 같은 룰 — dataclass 당 property 1 개,
약 3 줄.)
**카운팅 룰 — per-op 합산, dedup 없음.** `OpSpec` 안의 각 `TensorHandle`
참조는 독립적으로 카운트 — 같은 핸들이 여러 OpSpec 에 등장하거나
`rw_handles` 에도 있어도 마찬가지. `rw_handles` block 은 cross-composite
해저드 추적 (ADR-0065 D6.3) 의 메타데이터로 `ops` 의 operand 참조와
**별도** 카운트. 중복 제거 없음. HW 와 일치: descriptor 가 각 operand
slot 을 독립 address field 로 인코딩하고 dispatcher 가 `rw_handles`
별도 metadata block 으로 추적. 예: `OpSpec(kind="mul_bcast", operands=
{"src_a": O, "src_b": corr}, out=O)` 는 핸들 `O`*두 번* 카운트
(`src_a` 한 번, `out` 한 번); 만약 `O` 가 둘러싼 CompositeCmd 의
`rw_handles` 에도 있으면 세 번째 카운트.
### D3. Default — 일반 composite ≈ 43 ns 기준
Anchor: **DMA-staged GEMM 경로용 단일-OpSpec composite**
`OpSpec(kind="gemm", ...)` 1 개; DMA stage 는 PE_SCHEDULER 가 operand
`space` 에서 자동 삽입 (ADR-0065 D4), **`logical_bytes` 에 등장하지 않음**
(커널이 별도 cmd 로 emit 안 함). 분해:
```
framing 4
ops tuple length 1
GEMM OpSpec 40 (opcode 1 + scope 1 + 1 + 2 handles 16
+ out 8 + 1 + extra m/k/n 12)
rw_handles tuple length 1
rw_handles content 8 (출력 RW 핸들 1 개)
─────────────────────────────
total ~54 bytes
```
Dispatch 목표 ≈43 ns. on-die producer→consumer 큐 16 bytes/cycle
(일반 on-die descriptor queue 폭) 가정.
```
FIXED_PER_CMD = 40 cycles
R = 0.0625 cycles/byte (= 16 bytes/cycle)
```
검증: `40 + 54 × 0.0625 = 43.375 cycles ≈ 43 ns`
Cycle→ns 변환은 PE 노드의 기존 `clock_freq_ghz` attr (PE_MATH `_compute_ns`
가 쓰는 같은 것) 사용. cost-model knob 은 **cycle-domain 만** — clock
설정 중복 안 함.
### D4. topology config override
default 는 `pe_cpu.py` 에 내장. topology yaml 의 PE 노드 attrs 아래
`pe_cost_model:` 절로 override (cycle-domain knob 만; clock 은 PE 의
기존 `clock_freq_ghz` 사용):
```yaml
pe:
attrs:
clock_freq_ghz: 1.0 # 기존, cycle→ns 변환용
pe_cost_model:
fixed_per_cmd_cycles: 40
byte_cycles_recip: 0.0625 # = 16 bytes/cycle
max_composite_logical_bytes: 1024 # D7 — descriptor 크기 cap
```
누락된 키는 default. dispatch 공식은 PE_CPU init 시
`node.attrs["pe_cost_model"]` 에서 읽음.
### D5. Scope — 무엇이 cost 를 내고 무엇이 안 내나
| 경로 | dispatch cost? |
|---|---|
| PE_CPU → PE_SCHEDULER 의 모든 `PeCommand` | **예** |
| `PeCpuOverheadCmd` 자체 (이미 cycles 명시) | **아니오** (공식 우회) |
| PE_SCHEDULER 가 자동 생성한 Stage (DMA_READ/WRITE/FETCH/STORE) | **아니오** (PE_SCHEDULER 내부) |
| 엔진 compute latency (PE_DMA `drain_ns`, GEMM/MATH `_compute_ns`) | **변화 없음** — 엔진에 그대로 (SPEC §0.1) |
"latency 는 모델링된 컴포넌트에서" invariant 유지 — dispatch cost 는
*추가* CPU-side 시간, 엔진 시간에 fold 안 함.
### D6. 설정 가능 값; goldens 재생성
issue cost 를 0 → non-zero 로 바꾸면 **모든** bench 의 latency 변화.
이 ADR land 시 골든 latency **한 번** 재생성 — ADR-0062 D3 lazy-load 와
같은 패턴. 재생성 후 동일 calibration 이 ADR-0065 opt2 측정에 적용.
### D7. Composite 크기 cap (하드 제한 — validation error)
`CompositeCmd``logical_bytes`**`max_composite_logical_bytes`**
(default **1024 bytes**, D4 로 override) 로 제한. `logical_bytes` 가 cap
을 초과하는 composite 는 host-side TLContext 가 emit 시점에 `ValueError`
**거부****자동 segmentation 없음**. 커널 작성자가 각 `CompositeCmd`
가 descriptor capacity 에 맞도록 recipe 를 재구성(예: 여러 개의 더 작은
composite 로 명시적으로 분할)해야 함.
**cap 이 필요한 이유.** 실제 하드웨어는 descriptor queue entry 크기,
scheduler parser buffer, command SRAM, firmware 입력의 하드 제한 보유.
cap 없으면 `FIXED + bytes × R` 모델이 하드웨어가 받지 못할 정도로 큰
fused composite 를 보상 (예: primitive 100 개를 composite 1 개로 fuse,
FIXED 1 회만 지불).
**1024 bytes 라는 *특정* 숫자의 근거.** 이는 측정된 HW 숫자가 아니라
**safe engineering limit** — kernbench 코드베이스의 현재 모든 알려진
composite (가장 큰 것이 decode opt2 의 `#2` ~322 bytes) 보다 훨씬 위이면서
미래 recipe 도 지켜야 할 *유한* descriptor capacity 를 의도적으로 표현.
값은 topology 별 override (D4); 실제 HW reference 등장 시 재캘리브레이션.
이 default 의 역할은 cap 을 *원칙으로 존재시키는 것*, 특정 HW 에 맞추는
것이 아님.
**자동 segmentation 대신 하드 에러를 택한 근거.** 초과 composite 를 자동
분할(원래 Revision 2 제안)하는 것은 inter-segment 순서, `rw_handles`
공유, completion 체이닝 등 emitter 복잡도를 추가하면서, 실질적으로는
하드웨어가 허용하는 것보다 큰 descriptor 를 요청한 커널을 덮어주는
것일 뿐. 명시적 에러로 표면화하면 emitter 가 단순해지고, HW 제약이
작업을 어떻게 분할할지 가장 잘 아는 커널 작성자에게 보임.
Decode opt2 의 `#2` composite (10 ops, ~322 bytes) 는 1024 cap 안에
편안히 — GQA workload 에 에러 없음.
### D8. Single-op-cmd fast-path (cmd 타입별 FIXED)
모든 **single-op** cmd — DMA descriptor (`DmaReadCmd`, `DmaWriteCmd`)
*그리고* single-op compute dispatch (`GemmCmd`, `MathCmd`, `CopyCmd`)
— 는 일반 control-path 비용과 분리된, 가벼운 FIXED 를 지불한다:
**`fixed_per_single_op_cmd_cycles`** (기본 **8 cycles**) 가 이들 전부에서
`fixed_per_cmd_cycles` 를 대체한다. **`CompositeCmd` 만이 일반
control-path 비용** `fixed_per_cmd_cycles` (= 40 cycles) 를 유지하는
유일한 cmd 다.
구체적으로 D1 의 공식이 다음과 같이 확장된다:
```
dispatch_cycles(cmd) = FIXED(type(cmd)) + cmd.logical_bytes × R
FIXED(CompositeCmd) = fixed_per_cmd_cycles (= 40)
FIXED(나머지 전부, 즉 모든 single-op cmd:
DmaReadCmd / DmaWriteCmd / GemmCmd / MathCmd / CopyCmd)
= fixed_per_single_op_cmd_cycles (= 8)
```
**왜 single-op cmd 에 가벼운 FIXED 를 따로 두나.** D1 은 모든 발행 cmd 에
동일한 FIXED 를 매겼다 — 가장 단순한 모델이라서. 그러나 실제로 single-op
cmd — DMA descriptor 든, 엔진에 단일 GEMM / elementwise / copy 를
dispatch 하는 것이든 — 는 40-cycle FIXED 가 모델링하던 일반 control-path
비용보다 훨씬 가볍다. scheduler 측 plan 을 만들 필요 없이, descriptor 나
instruction 하나를 엔진 queue 에 push 하는 것뿐이다:
- NVIDIA Hopper TMA: 단일 PTX 명령 (`cp.async.bulk`) 으로 bulk async DMA
발사 — SM 측에선 ~1 ISA cycle, TMA 엔진이 fan-out / chunking 자체 처리.
- NVIDIA Ampere `cp.async`: warp-level async copy ~12 cycles.
- 단일 MMA / tensor-core 발행 (`wgmma`, `mma.sync`): math pipe 에 명령
하나, control-path round trip 아님.
- AMD AQL packet 을 HSA queue 에 write: ~515 cycles.
- 일반 descriptor-ring-push 디자인 (Synopsys/Xilinx-style DMA IP): MMIO
write 들 합쳐 ~520 cycles.
- Tenstorrent tile descriptor 발행, Habana Gaudi TPC descriptor RAM
+ start register: descriptor 당 한 자릿수 cycles.
40-cycle FIXED 는 `CompositeCmd` 에 한해 정당하다. composite 는 단일
dispatch 가 아니기 때문이다: scheduler 가 tile-feeder plan 을 만들고,
내부 `DMA_READ → FETCH → GEMM → STORE → DMA_WRITE` stage 들에 걸쳐
per-tile read/write hazard 를 추적하고, completion handle 을 배선해야
한다. single-op `GemmCmd` 는 그런 기구 없이 엔진 발행 하나일 뿐이다 —
이를 composite 와 같은 40 cycles 로 매기면 단일 명령을 스케줄된 plan
전체와 혼동하는 것이다.
**왜 이게 커널 평가에 중요한가.** 일정한 40-cycle FIXED 는 user 가 직접
orchestrate 하는 single-op-primitive 커널의 dispatch overhead 를 부풀린다.
`K = N_K · TILE_K` 의 chunked-prefetching async GEMM 은 K-tile 마다
`DmaReadCmd` (B-chunk load) **와** `GemmCmd` (per-chunk `tl.dot`) **와**
`MathCmd` (누적 `tl.add`) 를 발행하므로, per-cmd FIXED 가 dispatch
예산을 지배한다. single-op fast-path 는 D1 이 잡으려던 *구조적 cmd 수
신호* (`N` 개의 single-op cmd 를 발행하는 recipe 가 `N` tile 을 내부로
묶는 composite 하나보다 더 비싸야 한다는 점) 를 유지하면서, 40-cycle
일반 control-path 비용을 덧씌우는 모델링 과잉 청구는 분리한다.
**왜 굳이 8 cycles 인가.** 위 조사 범위 (TMA ~1, descriptor-ring ~15) 의
중간값. 의도적인 *모델링 디폴트* — 측정된 HW 수치가 아님. 이 디폴트의 역할
은 정성적 동작에서 single-op dispatch path 를 composite control path 와
구별하는 것. Topology override 가능 (D4).
**Composite 수치에 미치는 영향.** Composite tile 내부 stage (HW tile 당
`DMA_READ`, `FETCH`, `GEMM`, `STORE`, `DMA_WRITE`) 는 scheduler 의
tile-feeder loop 가 발행하므로 host-side TLContext 의 `_charge_dispatch`
를 안 거친다. 단일 `CompositeCmd` 자체는 여전히 40-cycle control-path
FIXED 를 지불한다. D8 은 user-side single-op primitive (`tl.load` /
`tl.store` / `tl.dot` / `tl.add` / …) 에 매기는 FIXED 만 바꾼다.
Composite 측정치는 안정 유지.
## Alternatives
### A1. Revision 1 의 op-type calibration 표 유지
기각: calibration 비용이 |kinds| 에 비례, 그리고 표가 잡으려던 *ratio*
구조적으로 cmd 크기의 함수. 구조 공식은 같은 정성적 동작을 N 개가 아닌
2 개의 calibratable 숫자로 달성.
### A2. byte-only 공식 (FIXED 항 없음)
기각. FIXED = 0 이면 opt2 (ADR-0065 Option Y) 가 opt3 를 **이기지 못함**
— per-tile dispatch *총 bytes* 가 비슷 (opt3 ≈ 232, opt2 ≈ 380); win 은
전적으로 *per-cmd fixed cost 횟수 감소* 에서 옴. byte-only 는 모델이
보여야 할 신호를 지움.
### A3. PE_SCHEDULER 에 dispatch 비용 부과
기각: saturation 질문은 *"CPU 가 descriptor 를 push 하는 속도가 엔진을
계속 바쁘게 할 수 있는가?"* — **PE_CPU** 의 issue-bandwidth 특성.
scheduler 에 부과하면 CPU back-pressure 모델링 안 됨.
### A4. DMA program/setup 시간을 별도 fixed per-descriptor 비용으로 모델
연기: 처음엔 descriptor-program 비용을 **issuing op 의** dispatch cost 에
fold. calibration 이 분리 필요성을 보이면 PE_DMA fixed setup 으로 분리.
## Consequences
### Positive
- hybrid 의 CPU-offload / saturation win (ADR-0060 §1) 이 **측정 가능**,
구조적으로 정직한 모델 (calibration 표 없음).
- 새 op_kind 추가 (예: ADR-0065 의 `softmax_merge` 8-step recipe) 가
*zero cost* — 같은 공식 자동 적용.
- 하드웨어에 더 충실 (queue-head MMIO RTT + queue-write 대역폭).
### Negative
- **모든** bench goldens 변화 → 일회성 재생성 (D6); CI 골든 fixtures 업데이트.
- 두 calibration knob (FIXED, R) 필요; default 는 문서화된 가정에 기반 —
HW reference 없는 동안엔 절대 latency 를 잠정 취급, **ratio** 만 방어.
- 각 PE command dataclass 에 작은 `logical_bytes` property 추가.
## Open review items
1. **FIXED 와 R 의 calibration 출처.** default 는 "일반 composite ≈ 43 ns
+ on-die 큐 16 bytes/cycle"; on-die descriptor 큐로 합리.
HW reference 등장 시 재방문.
2. **대형 composite 에서 Scheduler plan-gen 비용.** 0 유지 — D5 가
PE_SCHEDULER 의 plan 생성을 dispatch 공식 밖에 둠. D7 cap (1024 bytes
≈ 3035 OpSpec) 이 최악을 묶지만, cap 근처 composite 도 현재
zero-cost 모델에서 1-op composite 와 같은 scheduler 비용. stress test
(대형 composite microbench) 가 scheduler-bound 행동을 보이면
per-op-count `overhead_ns` 노출.
3. **Override 위치.** topology yaml 의 PE 노드 attrs 아래 `pe_cost_model:`
block — 모든 knob 을 한 곳에, 리뷰 가능. clock 은 PE 의 기존
`clock_freq_ghz` 에서 — 여기 중복 안 함.
4. **경로 parity.** greenlet (`_execute_legacy`, `kernel_runner`) 와
replay 둘 다 같은 cost model 읽어야 함. 검증.
5. **R 에 대한 결론의 민감도.** opt2 < opt3 가 합리적 `R` 범위 (큐 대역폭
가정) 에서 유지되어야 함. 민감도 sweep 이 Test Requirements (#9) 의
일부.
## Test Requirements
테스트는 specific 숫자 anchor 가 아닌 **공식 (D1)** 에 대해 작성 —
calibration 이 바뀌거나 OpSpec/CompositeCmd 필드가 추가되어도 유효.
1. **공식 보존.** 임의 `CompositeCmd` `c` 에 대해 PE_CPU 의 기록된
dispatch overhead 가 `FIXED_PER_CMD + c.logical_bytes × R`
(floor/round-off 의 ±1 cycle 내). 여러 composite parametrize:
1-OpSpec GEMM composite, 5-op MATH chain, 10-op recipe composite.
Default-calibration 숫자 (1-OpSpec composite 의 anchor ≈43 ns) 는
*informative* 참고, 테스트 gate 가 아님 — gate 는 공식 등가.
2. **정성적 ratio (robust).** opt3 per-tile PE_CPU dispatch 가 opt2
per-tile dispatch 를 최소 2× 차이로 엄격히 초과 — `opt3 > 2 × opt2`.
Default-calibration 모델은 ≈4× 예측; gate 는 느슨한 2× 한계라
calibration 이동 (예: HW reference 가 default 대체) 시 깨지지 않음.
Informative 숫자 — ADR-0065 §verification 와 DDD-0065 §11 의 모델
기대치 참조.
3. **Override 경로.** topology yaml 의 `pe_cost_model:` block 이 per-PE
dispatch cost 변경; block 누락 시 default 복귀. #1 의 공식 등가가
override 값에서도 성립.
4. **`PeCpuOverheadCmd` 우회.** 수동 `tl.cycles(n)` 는 정확히 `n` cycles,
`n + dispatch_cycles(...)` 아님.
5. **double-count 없음.** PE_DMA `drain_ns`, PE_GEMM/MATH `_compute_ns`
pre-ADR 값과 동일.
6. **결정성.** 동일 입력 → 동일 op_log + latency (SPEC §0.1).
7. **경로 parity.** greenlet 과 replay 가 동일 커널에 대해 동일
dispatch-cycle 회계.
8. **Composite 크기 cap (D7).** `logical_bytes >
max_composite_logical_bytes` 가 될 composite 는 emit 시점에 `ValueError`
(segmentation 없음). cap 이내 composite 는 정상 emit.
9. **민감도 (정성적).** `R ∈ {0.25, 0.0625, 0.03125}` cycles/byte 에서
3 지점 모두 `opt3 > opt2`. 방향 (R 감소 시 ratio 단조 증가) 도 단언,
단 절대 ratio 값은 *불필요*.
## Migration
ADR-0064 Revision 2 는 단일 PR 로 land:
- 각 `PeCommand` dataclass 의 `logical_bytes` property (D2 의 타입-aware
extra 카운트)
- `pe_cpu.py` dispatch 경로의 공식 적용
- PE_CPU init 의 `pe_cost_model:` override read (cycle-domain knob +
`max_composite_logical_bytes`)
- **composite 크기 cap (D7)** — TLContext-side 하드 cap,
`max_composite_logical_bytes` default 1024: cap 초과 composite 는 emit
시점에 `ValueError` (자동 segmentation 없음). 기존 bench 는 트리거
안 됨 (현재 최대 composite 가 ~322 bytes), 하지만 recipe 가 커질 때
descriptor-capacity 제한이 강제되도록 체크 land.
- 일회성 골든 재생성
land 후 ADR-0065 가 위에 쌓임 — 기존 bench 추가 골든 churn 없음
(ADR-0065 는 기존 경로의 `CompositeCmd` 의미 보존 refactor; opt2 만 새 bench).