2d8271c981
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>
435 lines
21 KiB
Markdown
435 lines
21 KiB
Markdown
# 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 ~1–2 cycles.
|
||
- 단일 MMA / tensor-core 발행 (`wgmma`, `mma.sync`): math pipe 에 명령
|
||
하나, control-path round trip 아님.
|
||
- AMD AQL packet 을 HSA queue 에 write: ~5–15 cycles.
|
||
- 일반 descriptor-ring-push 디자인 (Synopsys/Xilinx-style DMA IP): MMIO
|
||
write 들 합쳐 ~5–20 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
|
||
≈ 30–35 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).
|