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:
@@ -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
|
||||
|
||||
Reference in New Issue
Block a user