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>
This commit is contained in:
@@ -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 표 유지
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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)),
|
||||
)
|
||||
|
||||
@@ -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))
|
||||
|
||||
@@ -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():
|
||||
|
||||
@@ -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)]
|
||||
|
||||
Reference in New Issue
Block a user