diff --git a/docs/adr-ko/ADR-0064-perf-cpu-issue-cost-model.md b/docs/adr-ko/ADR-0064-perf-cpu-issue-cost-model.md index fee0146..c0bb05e 100644 --- a/docs/adr-ko/ADR-0064-perf-cpu-issue-cost-model.md +++ b/docs/adr-ko/ADR-0064-perf-cpu-issue-cost-model.md @@ -252,6 +252,74 @@ composite (가장 큰 것이 decode opt2 의 `#2` ~322 bytes) 보다 훨씬 위 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 표 유지 diff --git a/docs/adr/ADR-0064-perf-cpu-issue-cost-model.md b/docs/adr/ADR-0064-perf-cpu-issue-cost-model.md index 069284a..8d749e1 100644 --- a/docs/adr/ADR-0064-perf-cpu-issue-cost-model.md +++ b/docs/adr/ADR-0064-perf-cpu-issue-cost-model.md @@ -270,6 +270,80 @@ kernel author, who is best placed to decide how to split the work. Decode opt2's `#2` composite (10 ops, ~322 bytes) sits comfortably inside the 1024 cap — no error for the GQA workload. +### D8. Single-op-cmd fast-path (per-cmd-type FIXED) + +Every **single-op** command — the DMA descriptors (`DmaReadCmd`, +`DmaWriteCmd`) *and* the single-op compute dispatches (`GemmCmd`, +`MathCmd`, `CopyCmd`) — carries a separate, lighter FIXED than the +general control-path cost: **`fixed_per_single_op_cmd_cycles`** (default +**8 cycles**) replaces `fixed_per_cmd_cycles` for all of them. +**`CompositeCmd` is the only command that stays on the general +control-path cost** `fixed_per_cmd_cycles` (= 40 cycles). + +Concretely, the D1 formula becomes +``` +dispatch_cycles(cmd) = FIXED(type(cmd)) + cmd.logical_bytes × R +FIXED(CompositeCmd) = fixed_per_cmd_cycles (= 40) +FIXED(everything else, i.e. every single-op cmd: + DmaReadCmd / DmaWriteCmd / GemmCmd / MathCmd / CopyCmd) + = fixed_per_single_op_cmd_cycles (= 8) +``` + +**Why a separate, lighter FIXED for single-op cmds.** D1 treated every +emitted command uniformly because that is the simplest defensible +model. In practice an single-op command — whether a DMA descriptor or a +single GEMM / elementwise / copy dispatch to an engine — is *much* +cheaper than the generic control-path cost the 40-cycle FIXED was +modeling. It is a single descriptor or instruction pushed to an engine +queue, with no scheduler-side plan to build: + +- NVIDIA Hopper TMA: a single PTX instruction (`cp.async.bulk`) initiates + a bulk async DMA — ~1 ISA cycle on the SM side, with the TMA engine + doing fan-out and chunking itself. +- NVIDIA Ampere `cp.async`: ~1–2 cycles per warp-level async copy. +- A single MMA / tensor-core issue (`wgmma`, `mma.sync`): one + instruction to the math pipe, not a control-path round trip. +- AMD AQL packet write to a HSA queue: ~5–15 cycles. +- Generic descriptor-ring-push designs (Synopsys/Xilinx-style DMA IP): + ~5–20 cycles for the MMIO writes. +- Tenstorrent tile descriptor emission, Habana Gaudi TPC descriptor RAM + + start register: single-digit cycles per descriptor. + +The 40-cycle FIXED is justified for `CompositeCmd` specifically, because +a composite is *not* a single dispatch: the scheduler must build a +tile-feeder plan, track per-tile read/write hazards across its internal +`DMA_READ → FETCH → GEMM → STORE → DMA_WRITE` stages, and wire a +completion handle. An single-op `GemmCmd` is one engine issue with none of +that machinery; charging it the same 40 cycles as a composite conflates +a single instruction with a whole scheduled plan. + +**Why this matters for kernel evaluation.** A uniform 40-cycle FIXED +inflates the dispatch overhead of user-orchestrated single-op-primitive +kernels. A chunked-prefetching async GEMM at `K = N_K · TILE_K` emits, +per K-tile, a `DmaReadCmd` (B-chunk load) **and** a `GemmCmd` (the +per-chunk `tl.dot`) **and** a `MathCmd` (the running `tl.add` +accumulate) — so the per-command FIXED dominates its dispatch budget. +The single-op fast-path keeps the *structural* command-count signal that +D1 was designed to capture (a recipe that emits `N` single-op commands +*does* pay more than a recipe that emits one composite covering `N` +tiles internally) without conflating it with a modeling overcharge that +would be specific to a 40-cycle generic control path. + +**Why 8 cycles specifically.** 8 sits at the mid-range of the survey +above (TMA ~1 to descriptor-ring ~15). It is intentionally a *modeling +default*, not a measured HW number — the role of this default is to +distinguish the single-op dispatch path from the composite control path in +qualitative behaviour. The value is overridable per topology (D4). + +**Effect on composite numbers.** Composite tile-internal stages +(`DMA_READ`, `FETCH`, `GEMM`, `STORE`, `DMA_WRITE` per HW tile) are +emitted by the scheduler's tile-feeder loop, not by the host-side +TLContext, so they do **not** go through `_charge_dispatch`. The single +`CompositeCmd` itself still pays the 40-cycle control-path FIXED. D8 +only changes the FIXED charged to user-side single-op primitives +(`tl.load` / `tl.store` / `tl.dot` / `tl.add` / …). Composite +measurements remain stable. + ## Alternatives ### A1. Keep Revision 1's op-type calibration table diff --git a/src/kernbench/common/pe_cost_model.py b/src/kernbench/common/pe_cost_model.py index 433e2e5..e166bb6 100644 --- a/src/kernbench/common/pe_cost_model.py +++ b/src/kernbench/common/pe_cost_model.py @@ -1,20 +1,34 @@ -"""Structural PE_CPU dispatch cost model (ADR-0064 Revision 2). +"""Structural PE_CPU dispatch cost model (ADR-0064 Revision 2, with D8 +single-op-cmd fast-path amendment). Replaces the Rev1 per-op-type calibration table (``cpu_issue_cost.py``, removed) with a structural formula derived from each command's ``logical_bytes`` (ADR-0064 D2): - dispatch_cycles(cmd) = FIXED_PER_CMD + cmd.logical_bytes * R + dispatch_cycles(cmd) = FIXED(cmd_type) + cmd.logical_bytes * R -The per-command ``FIXED_PER_CMD`` term models the queue-tail update / -MMIO-class RTT / completion-event registration; the byte term ``R`` models -queue-write bandwidth. The primary signal the model exposes is -**command-count reduction** (FIXED-dominated); the byte term is a secondary -refinement (ADR-0064 Context). +The per-command ``FIXED`` term models the queue-tail update / MMIO-class +RTT / completion-event registration; the byte term ``R`` models queue-write +bandwidth. The primary signal the model exposes is **command-count +reduction** (FIXED-dominated); the byte term is a secondary refinement +(ADR-0064 Context). + +ADR-0064 D8 (single-op-cmd fast-path): every single-op command — the DMA +descriptors (``DmaReadCmd`` / ``DmaWriteCmd``) and the single-op compute +dispatches (``GemmCmd`` / ``MathCmd`` / ``CopyCmd``) — carries a +separate, lighter FIXED than the generic control-path cost. The default +(``fixed_per_single_op_cmd_cycles = 8``) reflects descriptor-ring-push / +single-instruction-dispatch patterns common in modern accelerators +(NVIDIA TMA ~1 ISA cycle, AMD AQL packet write ~5–15 cycles, generic +Synopsys/Xilinx DMA IP ~5–20 cycles). The general control-path cost +(``fixed_per_cmd_cycles = 40``) applies only to ``CompositeCmd`` — the +one command that needs scheduler-side plan generation, per-tile RW-hazard +tracking, and completion-handle wiring. Knobs are cycle-domain only; cycle→ns uses the PE node's ``clock_freq_ghz`` (ADR-0064 D3). Defaults anchor a typical 1-OpSpec GEMM composite (≈54 bytes) -at ≈43 ns on a 16 B/cycle on-die descriptor queue. +at ≈43 ns on a 16 B/cycle on-die descriptor queue, while a 64 B single-op +descriptor lands at ≈12 ns. A composite whose ``logical_bytes`` exceeds ``max_composite_logical_bytes`` is rejected at emit time with a ``ValueError`` (ADR-0064 D7, revised: hard @@ -27,15 +41,30 @@ from dataclasses import dataclass @dataclass(frozen=True) class PeCostModel: - """Cycle-domain dispatch cost knobs (ADR-0064 D3/D4).""" + """Cycle-domain dispatch cost knobs (ADR-0064 D3/D4/D8).""" fixed_per_cmd_cycles: float = 40.0 byte_cycles_recip: float = 0.0625 # = 16 bytes/cycle max_composite_logical_bytes: int = 1024 # D7 hard cap + fixed_per_single_op_cmd_cycles: float = 8.0 # D8 single-op-cmd fast-path - def dispatch_cycles(self, logical_bytes: int) -> float: - """FIXED + logical_bytes × R (ADR-0064 D1).""" - return self.fixed_per_cmd_cycles + logical_bytes * self.byte_cycles_recip + def dispatch_cycles(self, cmd_type: type, logical_bytes: int) -> float: + """FIXED(cmd_type) + logical_bytes × R (ADR-0064 D1/D8). + + ``CompositeCmd`` uses the general control-path + ``fixed_per_cmd_cycles``; every single-op command (DmaRead/ + DmaWrite/Gemm/Math/Copy) uses the lighter + ``fixed_per_single_op_cmd_cycles``. + """ + # Local import avoids a hard module-load cycle between common.pe_cost_model + # and common.pe_commands; only CompositeCmd is needed to discriminate + # the control path from the single-op fast path. + from kernbench.common.pe_commands import CompositeCmd + if cmd_type is CompositeCmd: + fixed = self.fixed_per_cmd_cycles + else: + fixed = self.fixed_per_single_op_cmd_cycles + return fixed + logical_bytes * self.byte_cycles_recip DEFAULT_PE_COST_MODEL = PeCostModel() @@ -54,4 +83,7 @@ def from_node_attrs(attrs: dict) -> PeCostModel: max_composite_logical_bytes=int( block.get("max_composite_logical_bytes", d.max_composite_logical_bytes)), + fixed_per_single_op_cmd_cycles=float( + block.get("fixed_per_single_op_cmd_cycles", + d.fixed_per_single_op_cmd_cycles)), ) diff --git a/src/kernbench/triton_emu/tl_context.py b/src/kernbench/triton_emu/tl_context.py index 7bde3ff..e702c96 100644 --- a/src/kernbench/triton_emu/tl_context.py +++ b/src/kernbench/triton_emu/tl_context.py @@ -229,7 +229,8 @@ class TLContext: f"CompositeCmd logical_bytes {lb} exceeds " f"max_composite_logical_bytes {cap} (ADR-0064 D7)" ) - cycles = self._cost_model.dispatch_cycles(cmd.logical_bytes) + cycles = self._cost_model.dispatch_cycles( + type(cmd), cmd.logical_bytes) ns = round(cycles / self._clock_freq_ghz) if ns > 0: self._emit(PeCpuOverheadCmd(cycles=ns)) diff --git a/tests/attention/test_gqa_decode_opt2.py b/tests/attention/test_gqa_decode_opt2.py index b7efee4..6386a7e 100644 --- a/tests/attention/test_gqa_decode_opt2.py +++ b/tests/attention/test_gqa_decode_opt2.py @@ -89,11 +89,23 @@ def _dispatch_cycles(emitter, cost_model) -> float: # ── dispatch ratio (ADR-0064 Test #9 / ADR-0065 Test #7) ───────────── -def test_opt3_dispatch_exceeds_opt2_by_2x(): +def test_opt3_dispatch_exceeds_opt2(): + """ADR-0064 Test #9 / ADR-0065 Test #7 — fusing opt3's per-tile + primitives into opt2's two composites lowers dispatch cost (the + CPU-offload win). + + ADR-0064 D8 (single-op fast-path) narrowed this win: single-op + commands now pay FIXED=8 while a composite pays FIXED=40, so opt2's + two composites no longer dominate opt3's many (now cheap) single-ops + by the pre-D8 2x. At the default model opt2 is still ~46% cheaper + (opt3 ≈ 1.87x opt2); the invariant that survives recalibration is the + direction plus a >1.5x margin. The sibling R-sweep test gates the + formula-level (direction + monotonicity) claim. + """ opt3 = _dispatch_cycles(_opt3_tile, DEFAULT_PE_COST_MODEL) opt2 = _dispatch_cycles(_opt2_tile, DEFAULT_PE_COST_MODEL) assert opt2 > 0 and opt3 > 0 - assert opt3 > 2 * opt2, f"opt3={opt3} must exceed 2x opt2={opt2}" + assert opt3 > 1.5 * opt2, f"opt3={opt3} must exceed 1.5x opt2={opt2}" def test_dispatch_ratio_R_sensitivity(): diff --git a/tests/test_pe_cost_model.py b/tests/test_pe_cost_model.py index f66bdc8..76a2a52 100644 --- a/tests/test_pe_cost_model.py +++ b/tests/test_pe_cost_model.py @@ -3,7 +3,7 @@ Rev2 replaces the Rev1 per-op-type cost table (``cpu_issue_cost.py``, ``DEFAULT_CPU_ISSUE_COST``) with a structural formula - dispatch_cycles(cmd) = FIXED_PER_CMD + cmd.logical_bytes * R + dispatch_cycles(cmd) = FIXED(cmd_type) + cmd.logical_bytes * R where ``logical_bytes`` is a structural property of each Pe command (ADR-0064 D2). Tests are written against the *formula* (D1) and the byte @@ -12,8 +12,9 @@ rule (D2), not specific absolute latencies, so they survive recalibration. Assumed post-Phase-2 surface: - ``kernbench.common.pe_cost_model`` exports ``PeCostModel`` (fields: ``fixed_per_cmd_cycles``, ``byte_cycles_recip``, - ``max_composite_logical_bytes``; method ``dispatch_cycles(lb)``) and - ``DEFAULT_PE_COST_MODEL`` (FIXED=40, R=0.0625, cap=1024). + ``max_composite_logical_bytes``, ``fixed_per_single_op_cmd_cycles``; + method ``dispatch_cycles(cmd_type, lb)``) and + ``DEFAULT_PE_COST_MODEL`` (FIXED=40, FIXED_SINGLE_OP=8, R=0.0625, cap=1024). - ``OpSpec``/``CompositeCmd``/``DmaReadCmd``/… expose ``logical_bytes``. - ``TLContext(cost_model=...)`` charges ``FIXED + lb*R`` (÷ clock) as a ``PeCpuOverheadCmd`` before each dispatched command; ``tl.cycles(n)`` @@ -67,15 +68,59 @@ def test_cost_model_defaults(): assert m.fixed_per_cmd_cycles == 40 assert m.byte_cycles_recip == 0.0625 # = 16 bytes/cycle assert m.max_composite_logical_bytes == 1024 + assert m.fixed_per_single_op_cmd_cycles == 8 # D8 single-op-cmd fast-path -def test_dispatch_cycles_formula(): - """dispatch_cycles(lb) == FIXED + lb*R (ADR-0064 D1).""" +def test_dispatch_cycles_composite_general_path(): + """CompositeCmd is the only command that keeps the general control-path + FIXED (= fixed_per_cmd_cycles); dispatch_cycles == FIXED + lb*R + (ADR-0064 D1/D8).""" from kernbench.common.pe_cost_model import PeCostModel + from kernbench.common.pe_commands import CompositeCmd m = PeCostModel(fixed_per_cmd_cycles=40, byte_cycles_recip=0.0625) - assert m.dispatch_cycles(54) == pytest.approx(43.375) # D3 anchor - assert m.dispatch_cycles(0) == 40 + assert m.dispatch_cycles(CompositeCmd, 54) == pytest.approx(43.375) + assert m.dispatch_cycles(CompositeCmd, 0) == 40 + + +def test_dispatch_cycles_single_op_fast_path(): + """ADR-0064 D8: every single-op command (DmaRead/DmaWrite/Gemm/Math/ + Copy) uses fixed_per_single_op_cmd_cycles instead of the general + fixed_per_cmd_cycles. Only CompositeCmd is excluded.""" + from kernbench.common.pe_cost_model import PeCostModel + from kernbench.common.pe_commands import ( + DmaReadCmd, DmaWriteCmd, GemmCmd, MathCmd, CopyCmd, + ) + + m = PeCostModel( + fixed_per_cmd_cycles=40, + fixed_per_single_op_cmd_cycles=8, + byte_cycles_recip=0.0625, + ) + for T in (DmaReadCmd, DmaWriteCmd, GemmCmd, MathCmd, CopyCmd): + assert m.dispatch_cycles(T, 64) == pytest.approx(12.0) + assert m.dispatch_cycles(T, 0) == 8 + + +def test_dispatch_cycles_yaml_override(): + """ADR-0064 D4 + D8: pe_cost_model.fixed_per_single_op_cmd_cycles in + topology attrs overrides the default; single-op cmds use it while + CompositeCmd uses the general fixed_per_cmd_cycles.""" + from kernbench.common.pe_cost_model import from_node_attrs + from kernbench.common.pe_commands import GemmCmd, CompositeCmd + + attrs = { + "pe_cost_model": { + "fixed_per_cmd_cycles": 50, + "fixed_per_single_op_cmd_cycles": 5, + "byte_cycles_recip": 0.125, + }, + } + m = from_node_attrs(attrs) + assert m.fixed_per_cmd_cycles == 50 + assert m.fixed_per_single_op_cmd_cycles == 5 + assert m.dispatch_cycles(GemmCmd, 0) == 5 + assert m.dispatch_cycles(CompositeCmd, 0) == 50 # ── logical_bytes byte rule (D2 / D3 worked example) ───────────────── @@ -146,7 +191,10 @@ def test_tlcontext_charges_formula_per_command(): tl = TLContext( pe_id=0, num_programs=1, - cost_model=PeCostModel(fixed_per_cmd_cycles=100, byte_cycles_recip=0.0), + cost_model=PeCostModel( + fixed_per_cmd_cycles=100, fixed_per_single_op_cmd_cycles=100, + byte_cycles_recip=0.0, + ), ) tl.load(0x1000, shape=(4, 4), dtype="f16") overheads = [c for c in tl.commands if isinstance(c, PeCpuOverheadCmd)]