ywkang/kernbench2
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Compare commits

base: ywkang/kernbench2:9270d3435a5e66e912500a17a5499a2311c3cbab
Branches Tags
ywkang/kernbench2:master ywkang/kernbench2:feat/gqa-composite-commands ywkang/kernbench2:sccl-distributed-allreduce ywkang/kernbench2:kernbench_Apr16
...
compare: ywkang/kernbench2:feat/gqa-composite-commands
Branches Tags
ywkang/kernbench2:master ywkang/kernbench2:feat/gqa-composite-commands ywkang/kernbench2:sccl-distributed-allreduce ywkang/kernbench2:kernbench_Apr16

47 Commits
9270d3435a ... feat/gqa-composite-commands

Author SHA1 Message Date
ywkang a4a2683aad paper(gqa): single-rank decode-streaming study — composite's masked bandwidth win 
Adds the memory-bound mirror of the compute-bound prefill figure. New single-rank decode sweep (T_q=1, M=8) reuses the prefill_compute_bound kernels, sweeping per-rank S_kv with NO cross-CUBE reduce — isolating the local attention that the 64-way Case-6 decode masks under its reduce tail.

Finding: even memory-bound decode benefits from the composite command. Its scheduler-streamed concurrent per-tile DMAs reach ~233 GB/s (91% of the 256 GB/s per-rank roofline) while the primitive's blocking tl.dot serializes one tile DMA and plateaus at ~166 GB/s — a 25-28% latency win that widens with context. This refines the paper's 'decode is latency-neutral' claim: neutral at 64-way production scale (reduce-dominated), but the composite extracts bandwidth at the local level. The bandwidth roofline here mirrors the MAC roofline in prefill.

Adds: gqa_decode_streaming.py sweep (wired into the milestone umbrella as GQA_1H_SWEEPS=decode_streaming), paper_plot_gqa_decode_streaming.py, the figure, and a new fig:gqa-decode-stream + paragraph in 05-gqa.tex; rebuilt main.pdf. Prefill (Experiment B) re-verified byte-identical (794.1/668.1/646.9us).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
 2026-06-23 19:39:45 -07:00
ywkang faf011dae5 paper(ipcq): regen alternatives figures + comparison plot + allreduce section 
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
 2026-06-23 19:19:55 -07:00
ywkang 1dade267ff experiment(gqa-decode): 16x16x16 MAC-block dispatch model 
Research artifact (NOT wired into the production sweep) for the decode composite-vs-primitive investigation. Models a primitive decode kernel that hand-blocks each tl.dot into 16x16x16 GemmCmds, charging per-MAC-block PE_CPU dispatch -- testing whether faithfully charging the dispatch that the composite form offloads to PE_SCHEDULER flips the 'composite gives no decode latency benefit' result.

Finding: it does NOT. Even with up to 512 serialized blocking GemmCmds per matmul (vs 2 coarse tl.dot), end-to-end decode latency is unchanged (8192: 28.9us, 32768: 114.9us) -- the inter-cube (m,l,O) reduce DMA tail dominates and the local-attention GEMM/dispatch sits in critical-path slack. Confirms the two-regime conclusion (24c7054): composite helps compute-bound prefill, not memory/reduce-bound decode.

Caveats: measured at S_kv 8K/32K (reduce-dominated); large-S_kv streaming regime extrapolated only. The -5.5% tiled<untiled is reduce-tail ordering noise, not fully nailed down.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
 2026-06-23 18:24:26 -07:00
ywkang 24c705419e paper(gqa): two-regime "use of composite commands" — decode + compute-bound prefill 
Complete the composite-command study across both attention regimes and
write up the result, resolving when the composite command helps latency.

Decode (memory-bound, T_q=1, M=8): the command form is latency-neutral —
the kernel is bound by KV streaming and the MAC array is near-idle, so the
composite's only benefit here is host-issue offload (O(n_tiles) -> O(1)
PE_CPU commands). Figure: gqa_decode_long_ctx_composite (latency flat,
command count saturates).

Prefill (compute-bound, large M=G*T_q tile-filling): add three command-form
variants of a single-rank FlashAttention prefill kernel
(_gqa_prefill_compute_bound: primitive / composite / composite_extended).
Here the composite's scheduler-internal per-tile DMA<->compute pipeline
keeps the MAC array fed while the primitive's blocking tl.dot starves it,
so composite wins on MAC utilization (68% flat -> 83%) and wall-clock
(19% faster at ctx=1024), with the margin growing with context (deeper
P.V reduction = more tiles to pipeline) — the compute-bound mirror of the
GEMM result in section 3. Figure: gqa_prefill_compute_bound (latency +
MAC util). Sweep wired as GQA_1H_SWEEPS=prefill_cb; tests cover structure,
e2e, and composite<primitive in the compute-bound regime.

The synthesis: composite has two benefits set by roofline position —
host-issue offload (always) and MAC-array feeding (compute-bound only).
Decode exercises the first, prefill both.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
 2026-06-23 17:04:26 -07:00
ywkang cd6f0ed91d feat(gqa): Case-6 composite-command decode variants + on-chip-operand DMA fix 
Add two command-form variants of the Case-6 (Cube-SP x PE-SP) long-context
decode kernel alongside the primitive baseline:
  - composite: per-tile GEMMs issued as one coarse tl.composite over the
    full S_local (PE_SCHEDULER tiles/streams K,V) -- O(1) PE_CPU commands
    vs the primitive kernel's O(n_tiles).
  - composite_extended: Q.K^T composite + softmax_merge recipe composite
    (ADR-0065), folding the per-tile online merge + P.V.

Extract the shared two-level (m,l,O) reduce into _gqa_mlo_reduce and
refactor the baseline to use it (byte-equal, guarded by a 64-rank
command-stream digest test).

Engine fix: _make_compute_out omitted pinned=True, so an on-chip TCM
compute result consumed as a composite GEMM operand was scheduled as an
HBM DMA_READ of its bit-61 scratch address -> PhysAddrError in data mode.
Mark compute outputs pinned (resident), matching the composite
auto-output and recipe scratch handles. .pinned is read only by the
tiling DMA gate, so this is byte-equal for existing benches.

Add the composite sweep (emit-level PE_CPU dispatch to 1M + data-mode
latency to 128K), its plot, umbrella wiring (GQA_1H_SWEEPS=composite),
and tests (refactor guard, dispatch saturation, engine-fix, e2e).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
 2026-06-23 16:07:41 -07:00
ywkang e9a5c438e3 paper(1H): 6-case GQA long-ctx + IPCQ design alternatives + GEMM terminology 
- §6 GQA: rewrite long-context decode from 4-case to 6-case. Data
  Placement Policy now presents the six placement options and their
  intrinsic per-PE memory / per-token comm costs (no winner predicted);
  the long-context subsection selects the best placement for that regime
  (Case 6 ★, both-axes S_kv shard) from the measured 6-case sweep. Add
  KV-sharding diagram + analytical budget/summary figures.
- Regenerate the 6-row decode sweep (milestone-1h-gqa) and the
  sweep-dependent decode panels (latency/traffic/parallelism/memory) so
  figures, prose, and sweep_decode.json are mutually consistent.
- §5 All-Reduce: add IPCQ design alternatives (architecture + decision
  matrix) as design-rationale schematics (illustrative step-counts, not
  measured) and the bench-generated topology diagram.
- §4 GEMM: rename "user-orchestrated/user-level" -> "kernel-orchestrated/
  kernel-level" (orchestration runs in the kernel program vs the
  scheduler-orchestrated composite); minor accuracy fixes (7.18 TFLOP/s
  ~10% below peak; ~781 ns DMA).
- Recompile build/main.pdf.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
 2026-06-22 00:55:54 -07:00
mukesh bb668a3ac3 paper(gqa): derive (m,l,O) AR per-PE cost from kernel topology 
Replaces the unsourced 32 KB/layer placeholder in the analytical chart
with T*(N-1)/N where T = h_q * S_q * (d_head*2 + 8) ~ 2.1 KB
(per-PE (m, l, O) payload) and N is the participant count of the
reduce-only chain per stage. Cases 2/3/6 analytical bars now match
the simulator-measured numbers within ~10% (was 6-30x over).

Also:
- Bumps the measurement S_kv from 8 K to 64 K to verify Cases 1, 2, 3, 6
  are genuinely S_kv-independent and Cases 4, 5 scale linearly.
- In-bar descriptor on the paired chart is now wrapped to fit the bar
  width, black, no background box, and centered between the analytical
  and measured bars so the label applies to both.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-18 15:38:20 -07:00
mukesh 0662c76fa7 paper(ipcq): de-emphasize flow-diagram outer frame so inner flow stands out 
Drops the per-case panel outline from linewidth 2.2 to 0.9 with 85% alpha.
The case-color frame still anchors each panel but no longer competes
with the swimlanes, flit packets, and annotations inside.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-18 15:01:09 -07:00
mukesh 37dc48fd27 paper(ipcq): add latency-model-style flow + stacked architecture diagrams 
Adds two new IPCQ architecture views in the latency_model.png aesthetic:
a 2x2 flow figure (one per case) and a 4x1 stacked figure for taller
horizontal panels. Refactors per-case panel drawing into shared helpers
and disambiguates the flit labels (HMQ desc -> Hd so it does not
collide with data D). Two-row legend keys.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-18 14:55:59 -07:00
mukesh 7c45047e93 paper(ipcq): swap control-data overlay for latency PNG in figures/ 
Drops the control-data overlay from the paper figures folder and
adds the per-send latency cycle-stack PNG instead.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-18 14:06:32 -07:00
mukesh 73e0b315fe paper(gqa): table cleanup — drop measured-comm column, narrower fig, wrapped title 
Removes the (measured) comm column, widens the Notes column with
shorter wrapped text, scales the table figure down, and wraps the
heading to two lines so it sits flush with the table.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-18 14:02:42 -07:00
mukesh 1552028c25 paper(ipcq): add IPCQ vs HW alternatives comparative study 
6 generated figures comparing IPCQ (chosen) against the 3 HW alternatives
(Doorbell+Polling, HMQ, RDMA-CQ): architecture blocks, latency stack,
op counts, sequence flow, decision matrix, and control/data plane overlay.

Three figures also copied into the paper's figures/ folder for inclusion.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-18 13:27:23 -07:00
mukesh 56c51b184a paper(gqa): refresh 4 decode 6-case PNGs in figures/ with actual 6-case content 
The previous commit (8102ddb) renamed the figures-dir
gqa_decode_long_ctx_4cases_*.png → 6cases_*.png but only the rename
was staged; the cp that brought the regenerated 6-case content
landed after `git add`, so the committed bytes were still the old
4-case ones. This commit syncs the bytes.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-18 11:51:30 -07:00
mukesh 8102ddbe30 paper(gqa): rename long-ctx artifacts to gqa_long_ctx_6cases_* and add 4 decode 6-case charts 
Filename cleanup so every long-ctx GQA artifact has a consistent
"gqa_long_ctx_6cases_*" prefix (or "gqa_decode_long_ctx_6cases_*"
for decode-only charts). Old "4cases" / mixed names retired.

Renames (long_ctx + figures, content unchanged):
  gqa_hbm_budget.png                  -> gqa_long_ctx_6cases_hbm_budget.png
  gqa_4cases_summary.png              -> gqa_long_ctx_6cases_summary.png
  gqa_4cases_memory_comm_analytical   -> gqa_long_ctx_6cases_memory_comm_analytical.png
  gqa_4cases_memory_comm_paired       -> gqa_long_ctx_6cases_memory_comm_paired.png
  gqa_kv_sharding_6cases_diagram      -> gqa_long_ctx_6cases_kv_sharding_diagram.png
  gqa_kv_sharding_6cases_table        -> gqa_long_ctx_6cases_kv_sharding_table.png
  gqa_3cases_measured_comm.json       -> gqa_long_ctx_6cases_measured_comm.json
  gqa_decode_long_ctx_4cases_*.png    -> gqa_decode_long_ctx_6cases_*.png
                                         (figures dir; long_ctx never had old)

New 4-chart 6-case set in long_ctx output dir (regenerated by
paper_plot_gqa_decode_long_ctx_4cases.py, which now reads all 6
sweep_decode.json panels — Cases 1-6 with the same colour scheme
used elsewhere: red = overflow per-PE HBM, grey = neutral, blue
= Pareto-best ★):

  gqa_decode_long_ctx_6cases_latency.png
  gqa_decode_long_ctx_6cases_memory.png
  gqa_decode_long_ctx_6cases_parallelism.png
  gqa_decode_long_ctx_6cases_traffic.png

Generator scripts updated to write the new filenames + handle the
two new d_head-TP variants (Cases 4, 5) in their per-PE memory and
active-PE-count helpers. Figure widths bumped 10 -> 12 in to fit 6
multi-line case labels.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-18 11:49:03 -07:00
mukesh 84bb418e1e gqa(decode): add d_head-TP kernels + measurement runner + figure generators 
Two new long-ctx decode attention kernels for the d_head-TP sharding
variants the 6-case chart predicts:

  · _gqa_attention_decode_long_ctx_cube_sp_pe_tp_dhead.py
        Cube-SP × PE-TP-d_head (Case 4 in chart). Per cube holds
        S_kv/C tokens of full d_head; per PE holds same tokens but
        only d_head/P dims. Partial Q·Kᵀ scores reduced intra-cube
        before softmax; outer (m,ℓ,O) merge two-phase (intra+inter).

  · _gqa_attention_decode_long_ctx_cube_tp_dhead_pe_sp.py
        Cube-TP-d_head × PE-SP (Case 5). Per cube holds full S_kv
        with only d_head/C dims; per PE holds S_kv/P of those dims.
        Partial scores reduced inter-cube (UCIe) before softmax.

Sweep dispatch (gqa_decode_long_ctx_4cases.py) extended with two
new panels so the milestone-1h-gqa sweep covers all 6 cases.

Smoke test scripts/verify_case4_dhead_tp.py runs Cases 4/5/6 at
S_kv=2K to validate the kernels load and execute end-to-end.

Plus the figure-generation toolchain that produced the committed
PNGs in the prior commit (dd3337f):

  · paper_plot_gqa_4cases_summary.py    - 3-panel summary +
        2-panel (analytical / paired-measured) chart generator.
        _plot_comm now takes mode="analytical" | "paired".
  · paper_plot_gqa_kv_sharding_diagram.py  - 6-case 2-D KV-tensor
        diagram + companion comparison-table PNG.
  · measure_gqa_decode_placement_comm.py   - runs all 6 kernels at
        S_kv=8K, sums actual IPCQ-copy bytes from engine.op_log,
        scales partial-score AR ×128 to S_kv=1M, writes
        gqa_3cases_measured_comm.json (committed).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-18 11:26:05 -07:00
mukesh dd3337f2e4 paper(gqa): add 6-case KV-sharding figures (analytical + measured) 
Six PNGs covering the 6-case decode KV-placement story, copied
into both the kernbench 1H_milestone_output staging dir and the
1H-codesign-paper figures/ dir:

  · gqa_hbm_budget.png                 - per-PE HBM budget bar
  · gqa_4cases_summary.png             - 3-panel: HBM + KV mem + comm
  · gqa_4cases_memory_comm_analytical  - 2-panel: KV mem + comm (analytical)
  · gqa_4cases_memory_comm_paired      - 2-panel: + measured overlay
  · gqa_kv_sharding_6cases_diagram     - 2-D KV-tensor view (Y=S_kv, X=d_head)
  · gqa_kv_sharding_6cases_table       - companion comparison table

The original 4-case gqa_4cases_summary.png is updated to the new
6-case 3-panel layout. Generator scripts and Phase-1 d_head-TP
kernels remain untracked - separate commit scope.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-18 11:22:57 -07:00
ywkang b6315c3c90 paper(1H): reflect D8 single-op cost model in §2/§3.4 + figure/diagram regen 
Two strands bundled as the 1H-codesign-paper refresh unit:

(A) This session — single-op cost-model reflection (depends on 2d8271c):
  - §2 Table 2 (tab:hw): split "FIXED per command" into "FIXED per
    single-op command" (8 cycles) and "FIXED per composite command"
    (40 cycles); §2 dispatch-overhead prose updated to the two-class split.
  - §3.4 (sec:gemm-vs-async): rename paragraph headers + prose to
    async-full / async-tiled; "atomic" -> "single-op" throughout; reframe
    mechanism #3 from the old DMA-only fast-path to the single-op
    fast-path. Headline narrative now: even with EVERY single-op cmd
    (96 DMA + 48 dot + 47 add) charged the light 8-cycle FIXED, composite
    still wins ~2.8x at K=3072 purely on command-count structure (1 vs
    192 commands) -- down from the pre-D8 ~6.3x, and explicitly NOT a
    modelling artifact. Numbers refreshed from the regenerated sweep:
    async-full 3.83->3.91, async-tiled 1.14->~2.53, under-tile corner
    1.06->1.21, depth-2 vs depth-inf spread <1%. New figure wired in.
  - build/main.pdf rebuilt (tectonic); pdftotext-verified (no broken
    refs; Table 2 split, single-op terms, 2.8x/2.53/192-host-commands
    all present).

(B) Prior-session paper work riding along uncommitted: §4 all-reduce
    deep-edit, §5 GQA, §6 discussion trims; milestone_1h_ccl.py plot
    label "FSIM" -> "H2 2025 SW queue baseline"; regenerated diagrams
    under docs/diagrams/** and gemm output PNGs under
    1H_milestone_output/gemm/. (Composite-window gemm plots are
    unaffected by D8 — D8 only changes single-op dispatch FIXED, which
    the composite window excludes.)

All TODO items for the D8 single-op extension are now complete and
pushed across 3 commits (2d8271c cost-model+ADR+tests, 821bbf2 bench
harness, this paper refresh). Full regression green (826 passed, 1
skipped). No remaining work.

NOTE for review (carried from 2d8271c): ADR-0065's "2x CPU-offload win"
headline for GQA decode opt2 may want a refresh to the post-D8 ~1.87x.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
 2026-06-17 15:12:31 -07:00
ywkang 821bbf26a2 gemm(bench): composite-vs-async dispatch comparison harness + single-op-model sweep 
Adds the user-orchestrated async GEMM variants used in paper §3.4 to
contrast with the composite (load_ref) path:
  - matmul-async              : naive (tl.load full A+B, one tl.dot)
  - matmul-async-chunked      : depth-inf prefetch (all B-tiles queued)
  - matmul-async-chunked-db   : depth-2 double-buffer (TCM-bounded)
plus scripts/paper/paper_plot_gemm_async_vs_composite.py (4-way per-PE
TFLOPS sweep + figure) and the regenerated outputs under
1H_milestone_output/gemm/.

Sweep regenerated under the ADR-0064 D8 single-op cost model (commit
2d8271c). At K=3072 the composite-vs-async-tiled gap narrows from the
pre-D8 ~6.3x to ~2.8x (composite 7.18 / async-tiled 2.54 TFLOPS): D8
makes the async-tiled kernel's ~191 single-op dispatches 5x cheaper, so
its dispatch overhead drops to ~1.5us. Qualitative order persists:
composite > async-full (3.91) > async-tiled (2.54).

test_bench_registry.py: register matmul-async / -chunked / -chunked-db
(also picks up the earlier milestone-gqa-headline -> milestone-1h-gqa
rename already present in the tree).

--- Remaining work (resume here if interrupted) ---
- Paper §3.4 (03-gemm.tex sec:gemm-vs-async): finish naive->async-full /
  chunked->async-tiled rename AND reframe "FIXED_DMA=8 for DMA descriptors"
  to single-op(8) vs composite(40). Figure already copied to
  docs/report/1H-codesign-paper/figures/gemm_composite_vs_async_tflops.png.
- Paper §3.4 K=3072 para + mechanism #3: update dispatch breakdown to new
  model (191 single-op * 8c ~= 1.5us, was 4.6us) and headline gap
  ~6.3x -> ~2.8x.
- Rebuild build/main.pdf (tectonic) + verify via pdftotext.
- Then commit Group 3 (paper + diagrams + PDF).
- Table 2 / §2 prose / ADR-0064 D8 already done in 2d8271c.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
 2026-06-17 15:06:36 -07:00
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
ywkang 0c6ca0aaed paper(gemm): inline §3.1 pipeline-stage chain; drop display equation 
The five-stage chain was set as a centered display equation, which
overflowed the column and clipped the last 'DMA_WRITE' label. The
inline prose right below it also repeated the same chain in a
shorter form. Consolidate into one inline parenthetical listing the
five stage names, and reuse 'flows through the chain' for the
self-routing description — no clipping, no repetition.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-06-17 10:40:23 -07:00
ywkang b64c43b947 paper(gemm): drop in-chart titles + subtitles, let LaTeX caption carry that 
The matplotlib title and gray subtitle were duplicating what the
figure caption already says in the LaTeX source, and crowding the
top of every chart. Drop both — chart now leaves the title/subtitle
slots blank and the caption is the single source of truth for
chart titling.

Re-render all 5 GEMM figures (mac util measured, mac util theoretical
vs measured, stage breakdown, HBM BW util, per-PE TFLOPS) and resync
into docs/report/1H-codesign-paper/figures/. No data changes.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-06-17 10:29:13 -07:00
ywkang e8b7f4f064 paper(gemm): per-PE HBM BW + TFLOPS evaluation plots; unify peak to 8 TFLOP/s 
Two new evaluation plots for §3 GEMM, both using the composite window
as the denominator (the up-front tl.load of A in load_ref is therefore
excluded — only HBM traffic and compute inside the composite count):

- gemm_hbm_bw_util.png: per-PE HBM bandwidth utilization against the
  256 GB/s per-PE ceiling. Answers "is this workload memory-bound on
  this configuration?" — small/single-tile shapes sit at 22–46%
  (pipeline-fill-limited), large or output-write-dominated shapes
  saturate (≥85%).

- gemm_per_pe_tflops.png: per-PE GEMM throughput against the 8 TFLOP/s
  engine peak. The deep-K K=3072 load_ref case lands at
  7.18 TFLOP/s (~90% of peak) — the configuration's clean win.
  ref_ref at the same shape drops to 3.83 TFLOP/s because the second
  operand doubles HBM pressure and the kernel hits the BW ceiling.
  The variant gap on this chart is the operational cost of NOT
  pre-staging the activation.

Both charts compare two operand-staging variants on every shape:
load_ref ("activation pre-staged", weight-only HBM streaming) and
ref_ref (both A and W streamed from HBM). load_load is omitted —
with both operands pre-staged the composite carries no HBM traffic
and the BW metric collapses to 0.

Analytic / measured peak unified to the single hardware spec
(8.0 TFLOP/s = 8000 flops/ns). T_STAGE is now derived as
tile_flops / peak (= 16.384 ns for the 32×64×32 tile) rather than
the previous hardcoded 16 ns, so MAC efficiency and per-PE TFLOPS%
share the same denominator and never disagree. Side effect: max
analytic-vs-measured gap tightens from 2.2 ppt to 1.4 ppt across
the sweep.

§3 Results restructured: lead with the two new evaluation plots
(BW saturation, achieved throughput), then the MAC-utilization
analytic-vs-measured chart serves as the validation step, then the
per-stage engine wall-clock as the supporting diagnostic. §3
Analysis updated to use the load_ref / ref_ref contrast as the
hardware-software boundary line that motivates the GQA kernel of §5.

§2.4 Accuracy claim updated to match: GEMM analytic-vs-measured
agreement is now "within 1.4 ppt across every swept shape", from
~7.7% at single-tile up to ~90% at K=3072. Subtitle/title vertical
stacking glitch in the matplotlib charts fixed in passing.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-06-17 10:24:42 -07:00
ywkang 23992548f7 gemm(perf): fix back-to-back DMA pipelining + tighten analytic model 
Three coupled fixes that recover small-tile GEMM pipeline efficiency
from 53% to 88% (32x3072x32 load_ref, composite_window basis).

1. PE_DMA channel-hold (ADR-0014 D4 clarified): both the
   _handle_with_hooks (PeInternalTxn) and _pipeline_process
   (TileToken) paths used to hold the cap=1 DMA channel through the
   full HBM round-trip, which double-serialized with the HBM_CTRL's
   own per-PC `available_at` model and prevented back-to-back tile
   DMAs from amortizing their per-request head latency. Channel is
   now released after the request is enqueued onto the next hop;
   HBM serialization is HBM_CTRL's responsibility alone.

   Tests: new test_pe_dma_back_to_back_pipelining as the oracle
   (asserts wall < 75% of strict-serialized N x single_op). Existing
   test_pe_dma_record_start_after_channel_acquire rewritten to assert
   t_start clustering (channel released fast) instead of the old
   round-trip-hold invariant. test_pe_dma_same_channel_serializes
   still passes — HBM_CTRL preserves ordering.

   Probe regression: PE→local-HBM 32 KiB stays at 141 ns
   (single-request, unaffected).

2. milestone_1h_gemm bench: matmul_composite was reading MATMUL_M/K/N
   env vars at module load, so every sweep row replayed the cached
   256³ result; values now read inside run(). Drops the stale
   sys.modules deletion hack.

3. Analytic ideal-pipeline model: dropped the (n_mn-1)·dma_w_per_pair
   penalty (over-pessimistic for under-tile shapes — it pushed
   measured > theoretical) and replaced the D_STAGES-derived head
   with empirical T_PIPELINE_FILL=60 ns / T_PIPELINE_TAIL=30 ns.
   Max analytic-vs-measured gap across all 7 swept shapes now 2.2 ppt
   (was 9-44 ppt under the old constants).

Paper updates:
- §3 (GEMM): 78%→88% measured at 48 tiles, 23%→15% at 1 tile,
  stage breakdown numbers refreshed (DMA in / Fetch / GEMM all
  ~785 ns at K=3072), analytic-vs-measured agreement tightened
  to "within 2.2 ppt".
- §2.4 (Accuracy): GEMM tracking claim refreshed accordingly.
- §5 (GQA): restore long-ctx 4-cases figures into figures/
  (they were dropped from bench output dir as derived artifacts in
  92b9221 / e45626c but §5 still cites them by name).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-06-17 09:56:05 -07:00
mukesh 92b9221533 paper(gqa): drop committed long-ctx 4-cases derived chart PNGs 
These 8 prefill/decode 4-case PNGs are derived artifacts regenerated
on demand from scripts/paper/paper_plot_gqa_{prefill,decode}_long_ctx_4cases.py
and no longer belong in the committed bench output.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-17 09:39:26 -07:00
mukesh 694e0cc9b9 paper(gqa): add 6-case per-PE memory + comm summary chart 
Analytical plot generator + PNG covering the 4 sharding cases from
slide 17 plus the 2 d_head-TP variants (Cases 4, 5). Three panels:
per-PE HBM budget breakdown (Wq + Wk + Wv + Wo + 4.24 GB KV
headroom), per-PE KV memory at S_kv=1M, and per-PE communication
per output token (decode). Numbers match slide 17 (max KV context
~7.1M for the 64-way sharded cases).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-16 16:12:26 -07:00
ywkang 6b6e29968a paper: §2 platform deep-edit + §1 GQA framing + figure regen 
§2 (KernBench Platform):
- Fix Table 2 HBM aggregate BW (1024 → 2048 GB/s); drop stale
  hbm_total_bw_gbs from topology.yaml (never read by sim_engine)
- Split PE_CPU / PE_SCHED fixed-cost row; disambiguate from the
  40-cycle command-dispatch FIXED term
- §2.2 two-pass: expand to describe Pass 1 timing and Pass 2 data
  data-correctness path
- §2.3 dispatch model: add motivation sentence for the
  descriptor-size linear form
- §2.4 Accuracy: reorder GEMM → All-reduce → Probe →
  Simplifications; drop FSIM aside (already covered by §4 Fig 5
  caption); soften 'every ns' → 'every modeled latency
  contribution'
- §2.5 HW config: add 64 TFLOP/s vs 2048 GB/s (~31 FLOP/byte)
  balance-point intuition and forward pointer to §5 GQA decode
- Fig 1 caption: separate illustrative topology from experimental
  configuration

§1: tighten GQA-as-primary-bandwidth-bottleneck framing.

Figures: regenerate SIP / CUBE architecture (SVG sources + PDF +
generator scripts).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-06-16 16:00:23 -07:00
mukesh 77eece81c4 gqa(milestone-1h-gqa): add umbrella bench + first sweep outputs 
Single @bench entry that drives the prefill + decode long-context
4-cases sweeps in one invocation, mirroring milestone-1h-gemm. The
individual gqa_helpers/long_ctx/gqa_{prefill,decode}_long_ctx_4cases
modules are now pure helpers (no @bench), reached only via this
umbrella.

Env var contract:
  GQA_1H_RUN=1                    (required gate)
  GQA_1H_SWEEPS=prefill,decode    (default: both; pick subset to skip)
  GQA_1H_TOPOLOGY=topology.yaml   (override)

Output layout:
  benches/1H_milestone_output/gqa/long_ctx/
    sweep_prefill.json
    sweep_decode.json
    gqa_prefill_long_ctx_4cases_{latency,traffic,memory,parallelism}.png
    gqa_decode_long_ctx_4cases_{latency,traffic,memory,parallelism}.png

The decode bench config drops S_kv from 131_072 to 8_192 so the umbrella
finishes in minutes — the comparative-story bars are the same shape at
8K. The 131K production headline number is recoverable by overriding
the dispatch; the helper docstring notes this.

Outputs (2 sweep JSONs + 8 PNGs) ship in this commit alongside the
umbrella that produced them, following the milestone-1h-gemm pattern
where derived artifacts live with the bench that emits them.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-16 13:06:33 -07:00
mukesh 443ede99c7 gqa(prefill-4cases): add long-context prefill 4-cases comparative study 
Mirrors the existing decode 4-cases study for prefill on the LLaMA-3.1-70B
single-KV-head group (h_q=8, h_kv=1, d_head=128) across 8 cubes × 8 PEs:

  Case 1  Cube-SP × PE-TP  → KV S_kv-Ring across cubes; Q/O T_q-split 64-way
  Case 2  Cube-Repl × PE-TP → full KV per cube; only CUBE 0's P PEs split T_q
  Case 3  Cube-Repl × PE-SP → full KV per cube; PEs SP on S_kv (intra-cube AR)
  Case 4  Cube-SP × PE-SP  → KV split 64-way + Q T_q-split across cubes
                              (Ring KV + per-cube intra-CUBE AR)  ★ optimal

Cases 2 and 3 use Q-axis tiling (TILE_Q) so the per-PE scratch is bounded
by TILE_Q × TILE_S_KV regardless of T_q (the full Q tensor would be
G·T_q·d_head·2 bytes which scales linearly with T_q and blew the 1 MB
budget at T_q≥128 without tiling).

Per-panel try/except in run_sweep tolerates per-panel failures so
sweep.json always lands with whichever cases succeed plus a failures
list. Case 4 uses configure_sfr_intercube_ring with the snake submesh
(2,4) which installs both the Ring E/W lanes and the intra_* lanes
needed for the intra-CUBE reduce — single SFR config for all 4 cases.

The plot script generates 4 PNGs (latency, traffic, memory, parallelism)
into 1H_milestone_output/gqa/long_ctx/. The parallelism chart shows
total compute work (active_PE × T_q_per_PE × S_kv_processed) — exposes
Case 3's 8× redundancy vs. Cases 1/2/4 which all do unique work.

gqa_prefill_long_ctx_4cases.py exposes run_sweep() rather than a @bench
decorator — the umbrella bench in the next commit invokes it.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-16 13:06:08 -07:00
mukesh e45626c036 gqa: reorganize benches into gqa_helpers/ subpackage; drop legacy headline 
Splits the GQA helpers into a dedicated subpackage to make room for the
prefill 4-cases study (next commit) and a single umbrella bench
(milestone-1h-gqa, after that).

Layout:
  benches/gqa_helpers/
    long_ctx/    — decode 4-cases kernels + sweep runner
    short_ctx/   — prefill/decode short-context kernels
    shared/      — _gqa_panel_helpers + decode_opt2 (context-agnostic)

The registry audit now skips subpackages so gqa_helpers/ (without a
leading underscore) doesn't get audited for @bench decorators.

Also drops the legacy milestone-gqa-headline bench, its
_gqa_attention_prefill_long kernel, 6 dependent prefill tests, the
paper_gqa_latency.py report harness, and the 3 stale headline-derived
PNGs the §6 wire-up referenced (paper will re-pull from the new
1H_milestone_output/gqa/long_ctx/ once §6 is updated).

The _ccl_cfg and _summarize_op_log helpers used to live in the
headline bench; extracted them to gqa_helpers/shared/_gqa_panel_helpers.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-16 13:05:41 -07:00
ywkang 359a0eaa44 paper(platform): Table 1 -> table*, shrink Fig 1 row 1, move Fig 1/2 source up 
- Probe table (Table 1, tab:probe-pe-dma) promoted to table* so it
  spans both columns and no longer gets clipped at the column edge.
- Fig 1 (hw-arch) row 1 (SIP, CUBE subfigures) trimmed: widths
  0.495 -> 0.42 textwidth and a height=0.6\linewidth cap added so
  the figure occupies less vertical space and floats can land
  earlier in the document.
- Fig 1 and Fig 2 source blocks relocated to the top of §2 so they
  enter LaTeX's float queue before any §2 body text. As a result
  Fig 1 now lands on page 3 and Fig 2 on page 4 (was 4 and 5).
- Table 2 (tab:hw, modeled hardware config) reverted to plain
  \begin{table}; the previous \begin{table*} reshape was not
  requested.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-06-16 11:58:48 -07:00
ywkang 1cf8dd7868 paper(platform): widen SIP/CUBE subfigures, drop height cap 
- subfig widths 0.46 -> 0.495 textwidth so the \hfill gap between
  (a) SIP and (b) CUBE shrinks to ~2% textwidth instead of ~8%.
- Removed the height=0.85\linewidth cap on both subfigures; with the
  cap, the (roughly 1:1) SIP and CUBE diagrams were being scaled
  down to 85% even when the column had room. Let them render at
  their natural aspect ratio.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-06-16 11:22:01 -07:00
ywkang b715002b5a paper: §2 figure layout + Accuracy subsection + Table 1 → 2-col; §6 trim historical 4-panel content 
§2 platform:
- Accuracy promoted from \paragraph to \subsection (sec:accuracy);
  it sits at the same heading level as Why / Device / Latency / HW
  config and so reads as a first-class component of the platform
  description rather than a tail-end footnote.
- Three architecture diagrams (SIP, CUBE, PE) collapsed into one
  figure* using subcaption: Fig.~\ref{fig:hw-arch}(a) SIP and (b)
  CUBE side-by-side on row 1, (c) PE wide on row 2, total height
  capped at ~40% of the page. Subfigure cross-references rewritten
  to Fig.~ref{fig:hw-arch}\subref{...} in the body text. main.tex
  now pulls in the subcaption package.
- Table 1 (modeled hardware configuration) promoted to table*
  (two-column / full page width) and the row layout rewritten as a
  4-column tabular so Hierarchy + PE + Command-issue sit on the
  left and Memory + Interconnect sit on the right. The previous
  single-column rendering was getting cut at the right edge of the
  printed column.

§6 GQA:
- Removed the historical four-panel headline table (tab:gqa), the
  two figures (fig:gqa-lat, fig:gqa-break), and the prose paragraph
  that cited their per-panel numbers. Reason: the underlying
  milestone-gqa-headline bench has been simplified in collaborator
  commit 65c365f ("drop misleading single_user_/multi_user_
  panels") and no longer reproduces that data, so the section was
  left referencing a dataset the current bench cannot regenerate.
- The previously-added 4-cases long-context decode comparison is
  now §6's only Results subsection, retitled to
  "Results: long-context decode and parallelism strategies".

The headline GQA result that survives into the paper is therefore
the parallelism trade-off study (Case 4 / Cube-SP × PE-SP at 34 us
buying an 8x KV-memory reduction over the latency leaders).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-06-16 11:16:16 -07:00
ywkang 22f3968b51 paper(gqa): wire the 4-cases long-context decode comparison into §6 
The collaborator commit 7c346de added the comparative figures
(gqa_decode_long_ctx_4cases_{latency,memory,traffic}.png) but the
fused-GQA section still only referenced the four headline panels.
This commit closes that loop:

- New §6 subsection "Long-context decode: parallelism strategies"
  added between Results and Analysis. It lays out the four
  parallelism strategies (Cube-SP/Repl x PE-TP/SP) and pulls in the
  three new figures.
- The discussion highlights the central trade: the fastest strategy
  (Case 3, Cube-Repl x PE-SP at 20.2 us) requires replicating the
  full KV cache to every CUBE, while Case 4 (Cube-SP x PE-SP, the
  chosen design marked *) gives back ~14 us in exchange for an 8x
  KV-memory reduction.  Case 4's ~190 IPCQ copies + ~190 DMA reads
  are precisely the on-device collective traffic PE_IPCQ and the
  torus links of §5 are provisioned to absorb -- a direct payoff
  of the communication-side codesign work.
- Connects back to §5 (PE_IPCQ / all-reduce) so the reader sees the
  capstone arc: the GEMM enabler exposes the data-movement bound,
  the communication enabler attacks it, and the long-context
  parallelism study shows how the choice between the two extremes is
  framed by KV memory vs. on-device collective traffic.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-06-16 10:44:32 -07:00
ywkang e3f08972da paper: §2 polish + §5 fresh all-reduce data + kernbench probe table 
§2 platform:
- Add a "memory-centric AHBM" sentence to the device-and-execution
  intro: each PE is paired with dedicated HBM bandwidth and on-PE
  TCM, so the performance question is about feeding compute from
  locally-attached memory + moving the unavoidable inter-PE traffic
  efficiently. Makes the AHBM character of the platform visible
  before §2.3 starts unpacking the latency model.
- Promote "Congestion and contention modeling" from \paragraph to
  \subsubsection: this is the platform's core differentiator over a
  peak-BW roofline, so it deserves its own heading.
- Rename "Control-plane (issue) cost model" to "Command dispatch
  overhead model" -- describes what it actually charges (the PE_CPU
  paying a fixed + per-byte cost to push a command to one of the
  accelerator engines) without the more abstract "control-plane"
  framing.
- Add a third Accuracy cross-check from kernbench probe: a
  PE→HBM DMA-distance sweep that confirms monotonic hop progression,
  bandwidth saturation matching the per-edge model, and the built-in
  invariants (D2H >= H2D, cross-CUBE best < worst). New
  Table~\ref{tab:probe-pe-dma} reports per-traversal latency and
  utilisation at 32KiB / 1MiB across five hop classes.
- Captured probe output as figures/probe_pe_dma_summary.txt for
  reproducibility.

§5 PE_IPCQ / all-reduce:
- Re-ran milestone-1h-ccl on current sim_engine (post-ADR-0064 Rev2
  and the IPCQ slot-wrap Phase-2 race fix). Updated the topology-
  comparison table and the buffer-kind caption to the fresh raw
  latencies. Headline ratios are preserved:
    torus vs mesh saving at 96 KB/PE: 24.8% (was 24.8%) -> ~25%
    torus vs ring saving at 96 KB/PE: 19.3% (was 19.3%) -> ~19%
    TCM vs HBM saving at 64 KB/PE: 13.4% (was 13.9%) -> ~13%
    TCM vs SRAM saving at 64 KB/PE: 37.2% (was 38.3%) -> ~37%
  The Executive Summary's "up to ~14%" / "up to ~38%" framing stays
  consistent with these post values.

Artifacts refreshed in src/kernbench/benches/1H_milestone_output/:
- ccl/summary.csv + per-topology PNGs + buffer-kind CSV/PNG
- ccl/comparison_mesh_vs_ring_vs_2DTorus_vs_theoretical_vs_fsim.png
- gemm/* (re-run yields identical pe_window structure; PNGs
  refreshed)

Paper figures synced to the fresh artifacts:
- figures/allreduce_comparison.png
- figures/allreduce_buffer_kind.png

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-06-16 10:38:47 -07:00
mukesh 7c346dec1b paper(figures): GQA long-context decode 4-cases comparative figures 
Generated by scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py
from the bench's sweep.json (S_kv=8192, the headline 128K run was
sidestepped — Case 2's single-PE 128-tile sweep is the long pole
and shape/ordering is preserved across scales).

  gqa_decode_long_ctx_4cases_latency.png   end-to-end latency / case
  gqa_decode_long_ctx_4cases_traffic.png   ipcq / dma op-count breakdown
  gqa_decode_long_ctx_4cases_memory.png    KV bytes per cube (slide-11)

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 17:48:34 -07:00
mukesh 1fbe833992 gqa(decode-4cases): wire latency + engine occupancy into sweep.json; add comparative plot script (5C.F) 
Bench changes:
  - new _end_to_end_ns(op_log) and _engine_occupancy_ns(op_log) helpers
    in milestone_gqa_decode_long_ctx_4cases.py (mirror paper_gqa_latency.py)
  - _run_panel return dict now carries latency_ns + engine_occupancy_ns
    alongside op_log_summary, so sweep.json is the single source of truth
    for the comparative figures

Plot script:
  - new scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py reads
    sweep.json and emits 3 PNGs to docs/report/1H-codesign-paper/figures/:
      gqa_decode_long_ctx_4cases_latency.png   (end-to-end latency / case)
      gqa_decode_long_ctx_4cases_traffic.png   (ipcq/dma op counts / case)
      gqa_decode_long_ctx_4cases_memory.png    (KV bytes per cube / case)

Test changes:
  - 2 new tests verifying the helpers + _run_panel dict shape
  - lower smoke S_kv from 8192 -> 2048 (4x faster Case 2; assertions
    are S_kv-independent; one fold-loop iteration preserved)

18/18 tests pass in ~4 min.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 17:45:44 -07:00
ywkang 7f437a20bd paper(platform): edit-pass through §2 KernBench Platform done 
- Reordered §2 subsections to follow the SIP → CUBE → PE → graph
  flow: Why KernBench → Device and execution model → Latency model
  → Modeled hardware configuration. Readers now meet the device
  hierarchy before the graph abstraction that re-uses it.
- §2.2 Device and execution model: starts with the SIP/CUBE/PE
  hierarchy paragraphs (each anchoring fig:sip-arch, fig:cube-arch,
  fig:pe-arch); then the runtime-API/sim-engine/components bullet
  list; then the atomic-vs-composite command distinction (corrects
  the prior over-narrow framing that read every PE command as
  composite -- atomic single-engine commands exist too, and PE_CPU
  itself runs control-plane work directly).
- §2.3 Latency model: opens with the four-contribution decomposition
  (per-node overhead, per-edge transmission, drain, queuing delay)
  and the latency_model schematic; retains existing The hardware as
  a graph / From graph to DES / Latency contributions / Congestion
  / Control-plane cost model / Accuracy paragraphs. Accuracy
  paragraph now closes on KernBench's sufficiency for *relative*
  HW/SW design trade-offs given analytic + external-simulator
  agreement.
- New figures and assets:
  - figures/sip_architecture.pdf  (SIP-level graph view)
  - figures/cube_architecture.pdf (CUBE-level zoom-in)
  - figures/latency_model.png      (conceptual latency-model
                                   schematic with per-node /
                                   per-edge / drain / queuing-delay
                                   colour coding)
  - figures/pe_architecture.png    (carried over)
- Source-of-truth generator for the latency schematic:
  scripts/paper/paper_latency_model_diagram.py (a report-only
  harness under scripts/paper/ per the /paper isolation rule).
- main.tex preamble: \usepackage{tikz} added (kept from prior
  sequence-diagram draft -- harmless now that the latency model is
  a PNG; left in to keep paragraph numbering stable).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-06-15 17:17:19 -07:00
mukesh 756680f4e6 gqa: fold decode_long into the Case 4 kernel; drop 1D-chain dead code 
Option B fold: the lrab math previously in
``_gqa_attention_decode_long.py`` (used only as a thin re-export by
the Case 4 wrapper and as the import source for ``_merge_running``
in Cases 1 / 3) now lives inline in the Case 4 kernel file. ``sub_w``
is hardcoded to 4 (the 4×2 cube sub-mesh geometry; root cube 6);
the legacy ``sub_w=0`` 1D-chain backward-compat path is removed
along with its dedicated tests — the dropped ``single_user_*`` /
``multi_user_*`` panels that exercised it are already gone.

Production changes:
  - inline lrab math into _gqa_attention_decode_long_ctx_cube_sp_pe_sp.py
    (drops sub_w param; _ROOT_CUBE=6 baked in)
  - inline _merge_running into Cases 1 (cube_sp_pe_tp) and 3
    (cube_repl_pe_sp) so they no longer depend on the deleted file
  - delete src/kernbench/benches/_gqa_attention_decode_long.py
  - remove dead _run_decode_panel / _DECODE_* constants /
    decode-side _PANEL_DISPATCH / _make_bench_fn decode branch
    from milestone_gqa_headline.py (only the prefill panel remains)
  - update _gqa_attention_decode_opt2.py docstring reference

Test changes:
  - delete 7 legacy test_gqa_*.py files that pre-dated the 4-cases
    architectural split (coverage now subsumed by the 16 4-cases tests)
  - remove test_opt2_matches_opt3_data_mode + _run_decode_data helper
    from test_gqa_decode_opt2.py (the parity check required sub_w=0
    which no longer exists; opt2's other 4 tests preserved)

20/20 tests pass (16 4-cases + 4 opt2 smoke/dispatch).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 16:32:37 -07:00
mukesh ddee28a499 gqa(decode-4cases): rename bench/kernels/panels with long_ctx token 
The 4-cases comparative study is specifically about long-context
decode (LLaMA-3.1-70B, S_kv=128K, T_q=1); the long_ctx token in the
names makes the scope explicit and aligns with the existing
_gqa_attention_decode_long convention.

Renames (mechanical; no semantic change):
  - bench file:   milestone_gqa_decode_4cases.py
                  → milestone_gqa_decode_long_ctx_4cases.py
  - bench name:   milestone-gqa-decode-4cases
                  → milestone-gqa-decode-long-ctx-4cases
  - output dir:   1H_milestone_output/gqa_decode_4cases/
                  → 1H_milestone_output/gqa_decode_long_ctx_4cases/
  - env vars:     GQA_DECODE_4CASES_RUN / _TOPOLOGY
                  → GQA_DECODE_LONG_CTX_4CASES_RUN / _TOPOLOGY
  - 4 kernel files _gqa_attention_decode_<case>.py
                  → _gqa_attention_decode_long_ctx_<case>.py
  - 4 kernel functions gqa_attention_decode_<case>_kernel
                  → gqa_attention_decode_long_ctx_<case>_kernel
  - 4 dispatch kinds  decode_<case> → decode_long_ctx_<case>
  - 4 panel names     single_kv_group_decode_gqa_<case>
                      → single_kv_group_decode_long_ctx_gqa_<case>
  - 4 helper functions _run_decode_panel_<case>
                       → _run_decode_panel_long_ctx_<case>
  - test file renamed in lockstep

Also: Case 4 smoke test now uses its case-specific helper
_run_decode_panel_long_ctx_cube_sp_pe_sp (consistent with Cases 1-3)
instead of the legacy _run_decode_panel from milestone_gqa_headline.
16 tests pass.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 16:09:21 -07:00
mukesh 3c155be8e6 gqa(decode-4cases): give Case 4 its own thin kernel file for naming symmetry 
Cases 1-3 each have a dedicated _gqa_attention_decode_<case>.py
kernel file; Case 4 previously reached into _gqa_attention_decode_long.py
via the headline bench's _run_decode_panel helper, breaking the
one-file-per-case convention. Adds _gqa_attention_decode_cube_sp_pe_sp.py
as a 20-line wrapper that bakes in sub_w=4 (the C=8 lrab geometry)
and gives Case 4 its own kind ("decode_cube_sp_pe_sp") and helper
(_run_decode_panel_cube_sp_pe_sp). decode_long.py is unchanged
(still serves the legacy decode_long tests). 16 tests pass.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 15:58:11 -07:00
mukesh 0ef4fde5d8 gqa(decode-4cases): Case 1 — Cube-SP × PE-TP (inter-CUBE lrab only; PE-TP B=1 waste) (5C.A) 
Adds the fourth and final 4-cases panel: KV split S_kv-wise across
the 8 cubes (cube=row_wise, S_local = S_kv/C), replicated within
each cube (pe=replicate). PE-TP at B=1 means only PE 0 of each cube
has work; PEs 1-7 early-return (slide-11 PE-TP-at-B=1 waste). No
intra-CUBE comm; inter-CUBE 8-way reduce reuses the lrab-adapted
center-root pattern (root cube 6) — same structural cost as Case 4's
inter-CUBE phase (21 ipcq_copy). 16 tests pass (4 per case).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 15:30:00 -07:00
mukesh ae942f6959 gqa(decode-4cases): Case 3 — Cube-Repl × PE-SP (intra-CUBE AR only; 8× memory) (5C.C) 
Adds the third 4-cases panel: KV replicated per cube (8× memory
waste), PEs SP on S_kv within each cube, intra-CUBE 8-way reduce on
(m, ℓ, O), and no inter-CUBE comm (every cube ends with full answer;
designated writer = cube 0). Reuses the row-chain + col-bridge
intra-CUBE pattern that anchors Case 4 (21 ipcq_copy per cube × 8
cubes = 168 total). 12 tests pass (4 Case 4 + 4 Case 2 + 4 Case 3).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 15:18:41 -07:00
mukesh 5672c8f3ef gqa(decode-4cases): Case 2 — Cube-Repl × PE-TP (no comm; 8× memory) (5C.B) 
Second case in the GQA decode 4-cases comparative study per
GQA_full_deck.pptx slide 11. Case 2 replicates K, V across all 8
cubes × 8 PEs (the slide-11 8 KB/tok/PE memory waste) and has zero
inter-rank comm. For B=1 (single-user decode default), only PE 0
of CUBE 0 has work — the inherent PE-TP waste slide 11 calls out.

Changes:
- New kernel: src/kernbench/benches/_gqa_attention_decode_cube_repl_pe_tp.py
  Simplest of the 4 cases. Active rank loads full Q/K/V from HBM,
  computes attention via S_kv tile sweep with online-softmax merge,
  writes O. All non-(0,0) ranks early-return. No tl.send/recv.
- src/kernbench/benches/milestone_gqa_decode_4cases.py:
    - Add panel single_kv_group_decode_gqa_cube_repl_pe_tp (Case 2)
      to _PANELS and _PANEL_DISPATCH.
    - Add _run_decode_panel_cube_repl_pe_tp helper: DPPolicy K/V/Q/O
      = cube=replicate, pe=replicate (models 8× memory waste).
    - Extend _make_bench_fn to dispatch kind="decode_cube_repl_pe_tp"
      to the new runner.
- tests/attention/test_milestone_gqa_decode_4cases.py:
    4 new tests assert Case 2 contract: panel registered, smoke
    completion, zero ipcq_copy (no comm), single dma_write from cube 0.

Verification: 8/8 tests pass (4 Case 4 anchor + 4 new Case 2).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 14:59:58 -07:00
mukesh 65c365f858 bench(milestone-gqa-headline): drop misleading single_user_/multi_user_ panels 
The legacy panel names suggested batched serving semantics they never
had — all four modeled a single user with KV sharded differently
(C=1 single-cube; C=4 multi-cube Cube-SP), at toy dims (T_q=4, S_kv≤128).
The single-KV-group C=8 panel + the new milestone-gqa-decode-4cases
bench cover the meaningful comparisons; pytest regression already
covers C=1/C=4 configurations end-to-end at richer scale.

Changes:
- milestone_gqa_headline.py: drop the 4 legacy panels; _PANELS now
  contains only single_kv_group_prefill_gqa_c8_p8. Update docstring.
- tests/attention/test_milestone_gqa_headline.py: drop the 3 legacy-
  panel architectural tests (Ring-KV traffic, root-only decode write,
  per-CUBE distributed output) and test_decode_panels_use_real_gqa
  (no decode panels in this bench anymore). Equivalent properties
  are asserted in test_milestone_gqa_single_kv_group_prefill_panel.py
  (64 dma_writes, 896 ipcq_copy) and test_milestone_gqa_decode_4cases.py
  (1 dma_write at cube 6, 189 ipcq_copy).
- tests/attention/test_milestone_gqa_single_kv_group_prefill_panel.py:
  drop test_existing_prefill_panel_runner_backward_compat (it exercised
  multi_user_prefill_gqa which no longer exists).
- scripts/paper/paper_plot_gqa.py: replace the 4 legacy _LABELS entries
  with the single single_kv_group_prefill_gqa_c8_p8 label.
- Regenerate 1H_milestone_output/gqa_headline/sweep.json from the new
  panel set.

Verification: 9/9 tests pass across the 3 affected test files.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 14:50:25 -07:00
mukesh c164645aee gqa(decode-4cases): Case 4 anchor in dedicated bench (5C.D) 
First milestone of the decode 4-cases comparative study per
GQA_full_deck.pptx slides 11-17. Case 4 (Cube-SP × PE-SP, the optimal
case per slide 11) is structurally the existing _gqa_attention_decode_long
at sub_w=4 (Increment 2's lrab-adapted center-root reduce). This commit
wires it into a dedicated bench so the remaining cases land alongside.

Changes:
- New bench: src/kernbench/benches/milestone_gqa_decode_4cases.py
  Houses the 4 case panels under one milestone-gqa-decode-4cases entry
  (gated by GQA_DECODE_4CASES_RUN=1; output to
   1H_milestone_output/gqa_decode_4cases/sweep.json). Cases 1-3 are
  TBD in subsequent sub-increments (5C.A/B/C).
- New panel: single_kv_group_decode_gqa_cube_sp_pe_sp
  C=8, P=8, sub_w=4, T_q=1, S_kv=131_072, d_head=128, h_q=8, h_kv=1.
- src/kernbench/benches/milestone_gqa_headline.py: _run_decode_panel
  extended with keyword-only sub_w/T_q/d_head/h_q/h_kv overrides
  (defaults preserve existing-panel behaviour).
- tests/attention/test_milestone_gqa_decode_4cases.py: 4 new tests
  asserting registration, smoke completion, reduce-to-root at the lrab
  center cube (cube 6), and the predicted 189-ipcq Case-4 traffic
  pattern (168 intra-CUBE + 21 inter-CUBE lrab Phase 1+2).
- tests/attention/test_milestone_gqa_headline.py: rename
  test_sweep_json_has_four_panels -> test_sweep_json_has_expected_panels
  and switch hardcoded 4 to len(PANELS) (the panel set grew to 5
  with Increment 5's single_kv_group_prefill_gqa_c8_p8).

Deviation noted: slide 13 prescribes AllReduce on (m,ℓ,O); our kernel
does reduce-to-root (only the lrab center cube has the answer) per
ADR-0060 §4. Treated as the kernbench Case-4 baseline.

Verification: all 4 new tests pass; 90 regression tests pass; the
previously-failing test_sweep_json_has_four_panels now passes under
its renamed form.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 14:31:44 -07:00
mukesh 5b4d9cb597 bench(milestone-gqa-headline): scale single_kv_group prefill panel T_q=S_kv=1K (scratch budget) 
The headline T_q=S_kv=32K target overflows the 1 MB per-PE scratch
pool: at T_q_local=4K and tile_s=1024 the scores matrix alone is 8 MB.
The prefill kernel's bootstrap section also leaves K_t/V_t/scores/
exp_scores persistent (outside tl.scratch_scope), inflating baseline.

Scale-down to T_q=S_kv=1K (T_q_local=128, fits comfortably) preserves
the C=8 + P=8 architecture demonstration; the true LLaMA 32K headline
awaits a future increment to add Q-axis tiling and tighten bootstrap
scratch discipline.

Verified end-to-end: kernbench run --bench milestone-gqa-headline now
produces sweep.json with all 5 panels. The new panel shows
ipcq_copy=896 (matches (C-1)·n_tiles·2·C·P = 7·1·2·8·8) and
dma_write=64 (one per PE, head-parallel + intra-CUBE PE-SP).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 14:00:30 -07:00
mukesh 9e1242039b gqa: single-KV-group LLaMA-3.1-70B prefill milestone (Increments 1-5) 
End-to-end wires C=8 P=8 d_head=128 prefill with snake-ring inter-CUBE
SFR, intra-CUBE PE-SP (all 64 ranks active), and the milestone bench
panel. Decode kernel gains lrab-adapted center-root reduce for the
2×4 sub-mesh per ADR-0060 §4.2.

Increment 1 — SFR multi-row snake
  src/kernbench/ccl/sfr_config.py: configure_sfr_intercube_ring gains
  submesh_shape / submesh_origin kwargs; installs a Hamiltonian snake
  ring through a rectangular sub-mesh (every hop is 1-hop physical
  neighbour). Backward-compat: 1D-row behaviour preserved when
  submesh_shape is None.
  tests/test_intercube_snake_ring.py (12 tests)

Increment 2 — Decode lrab-adapted center-root reduce
  src/kernbench/benches/_gqa_attention_decode_long.py: new sub_w param
  (default 0 = existing 1D-chain). sub_w >= 2 selects the ADR-0060
  §4.2 prescribed lrab-adapted Phase 1+2 reduce (bidirectional row +
  bidirectional col converge to the center cube), with log-sum-exp
  _merge_running replacing the plain + of lrab.
  tests/attention/test_gqa_decode_long_2d_reduce.py (4 tests)

Increment 3 — Prefill kernel at C=8 (no production change)
  Verified by inspection that the existing prefill_long kernel +
  Increment 1's snake SFR already work at C=8 without any kernel
  edit. The kernel speaks logical W/E; the snake routes it.
  tests/attention/test_gqa_prefill_long_c8_snake.py (3 tests)

Increment 4 — Intra-CUBE PE-SP in prefill (all 64 ranks)
  src/kernbench/benches/_gqa_attention_prefill_long.py: new P param
  (default 1 = existing PE-0-only). P > 1 splits T_q query-axis-wise
  across the P PEs of each CUBE; output rows are disjoint per PE so
  no intra-CUBE reduce is needed; each PE drives its own same-lane
  ring (P parallel rings).
  tests/attention/test_gqa_prefill_long_pe_sp.py (5 tests)

Increment 5 — LLaMA-scale milestone bench panel
  src/kernbench/benches/milestone_gqa_headline.py: new panel
  single_kv_group_prefill_gqa_c8_p8 (C=8, P=8, T_q=S_kv=32K,
  d_head=128). _run_prefill_panel extended with P/T_q/d_head
  defaults; routes snake SFR when C > mesh_w.
  tests/attention/test_milestone_gqa_single_kv_group_prefill_panel.py (3 tests)

Total: 4 production files modified, 5 new test files, 27 new tests.
Followed the Phase 1/2 protocol per CLAUDE.md throughout.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-06-15 13:42:31 -07:00
187 changed files with 15072 additions and 3642 deletions
Expand all files Collapse all files
docs/adr-ko/ADR-0064-perf-cpu-issue-cost-model.md
+68
View File
@@ -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 표 유지
docs/adr/ADR-0014-dev-pe-pipeline-execution-model.md
+9 -2
View File
@@ -184,8 +184,15 @@ The empty-`ops` form is the legacy single-op path.
- `DMA_READ`: `simpy.Resource(capacity=1)`.
- `DMA_WRITE`: `simpy.Resource(capacity=1)`.
- Both channels run concurrently (READ ∥ WRITE allowed).
- Within a channel, requests serialize (READ ∥ READ disallowed; same
for WRITE).
- Within a channel, requests serialize **at the issue path** (READ ∥ READ
issued out-of-order is disallowed; same for WRITE). The channel is
held only until the request is enqueued onto the next hop (router),
then released — it does **not** block until the HBM round-trip
completes. Multiple in-flight requests are therefore allowed, and
HBM-level serialization is the responsibility of the HBM controller's
per-PC `available_at` timestamps (ADR-0033 D1). This is what allows
back-to-back small-tile DMAs to amortize the per-request head latency
through the fabric.
- `vc_comm` is an orthogonal channel for IPCQ traffic defined in
ADR-0023 D8 — out of scope for this ADR.
docs/adr/ADR-0064-perf-cpu-issue-cost-model.md
+74
View File
@@ -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
docs/diagrams/allreduce_latency_plots/AllReduce_LRAB_2DMesh_6SiP_2x3.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 38 KiB

After

Width:  |  Height:  |  Size: 38 KiB

docs/diagrams/allreduce_latency_plots/AllReduce_LRAB_2Dtorus_6SiP_2x3.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 36 KiB

After

Width:  |  Height:  |  Size: 36 KiB

docs/diagrams/allreduce_latency_plots/AllReduce_LRAB_2Dtorus_6SiP_2x3_with_TCM_SRAM_HBM.csv
+12 -12
View File
@@ -1,13 +1,13 @@
buffer_kind,sip_topology,n_sips,n_elem,bytes_per_pe,latency_ns
hbm,torus_2d,6,128,256,2120.040000000012
hbm,torus_2d,6,1024,2048,2717.2783333333473
hbm,torus_2d,6,8192,16384,7315.184999999989
hbm,torus_2d,6,32768,65536,23081.26500000037
sram,torus_2d,6,128,256,2060.040000000012
sram,torus_2d,6,1024,2048,2909.2783333333473
sram,torus_2d,6,8192,16384,9523.184999999869
sram,torus_2d,6,32768,65536,32201.265000000385
tcm,torus_2d,6,128,256,1964.040000000012
tcm,torus_2d,6,1024,2048,2477.2783333333473
tcm,torus_2d,6,8192,16384,6403.185000000109
tcm,torus_2d,6,32768,65536,19865.265000000378
hbm,torus_2d,6,128,256,2345.040000000012
hbm,torus_2d,6,1024,2048,2942.2783333333473
hbm,torus_2d,6,8192,16384,7540.184999999989
hbm,torus_2d,6,32768,65536,23306.26500000037
sram,torus_2d,6,128,256,2285.040000000012
sram,torus_2d,6,1024,2048,3134.2783333333527
sram,torus_2d,6,8192,16384,9748.184999999869
sram,torus_2d,6,32768,65536,32426.265000000385
tcm,torus_2d,6,128,256,2189.040000000012
tcm,torus_2d,6,1024,2048,2702.2783333333473
tcm,torus_2d,6,8192,16384,6628.18500000005
tcm,torus_2d,6,32768,65536,20090.265000000378
1 buffer_kind sip_topology n_sips n_elem bytes_per_pe latency_ns
2 hbm torus_2d 6 128 256 2120.040000000012 2345.040000000012
3 hbm torus_2d 6 1024 2048 2717.2783333333473 2942.2783333333473
4 hbm torus_2d 6 8192 16384 7315.184999999989 7540.184999999989
5 hbm torus_2d 6 32768 65536 23081.26500000037 23306.26500000037
6 sram torus_2d 6 128 256 2060.040000000012 2285.040000000012
7 sram torus_2d 6 1024 2048 2909.2783333333473 3134.2783333333527
8 sram torus_2d 6 8192 16384 9523.184999999869 9748.184999999869
9 sram torus_2d 6 32768 65536 32201.265000000385 32426.265000000385
10 tcm torus_2d 6 128 256 1964.040000000012 2189.040000000012
11 tcm torus_2d 6 1024 2048 2477.2783333333473 2702.2783333333473
12 tcm torus_2d 6 8192 16384 6403.185000000109 6628.18500000005
13 tcm torus_2d 6 32768 65536 19865.265000000378 20090.265000000378
docs/diagrams/allreduce_latency_plots/AllReduce_LRAB_2Dtorus_6SiP_2x3_with_TCM_SRAM_HBM.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 75 KiB

After

Width:  |  Height:  |  Size: 77 KiB

docs/diagrams/allreduce_latency_plots/AllReduce_LRAB_Ring1D_6SiP_1x6.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 37 KiB

After

Width:  |  Height:  |  Size: 37 KiB

docs/diagrams/allreduce_latency_plots/comparison_mesh_vs_ring_vs_2DTorus_vs_theoretical_vs_fsim.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 86 KiB

After

Width:  |  Height:  |  Size: 92 KiB

docs/diagrams/allreduce_latency_plots/summary.csv
+36 -36
View File
@@ -1,37 +1,37 @@
algorithm,sip_topology,n_sips,n_elem,bytes_per_pe,bytes_per_sip,latency_ns
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8,16,256,2666.552500000015
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32,64,1024,2747.7400000000152
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,64,128,2048,2855.990000000018
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,128,256,4096,3072.490000000019
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,512,1024,16384,3337.1133333333582
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,1024,2048,32768,3708.0333333333692
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,2048,4096,65536,4449.873333333393
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,4096,8192,131072,5933.020000000124
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8192,16384,262144,8900.379999999863
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,16384,32768,524288,14835.099999999224
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32768,65536,1048576,26704.540000000765
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,49152,98304,1572864,38573.97999999701
lrab_hierarchical_allreduce,ring_1d,6,8,16,256,2365.255833333347
lrab_hierarchical_allreduce,ring_1d,6,32,64,1024,2436.9433333333473
lrab_hierarchical_allreduce,ring_1d,6,64,128,2048,2532.526666666683
lrab_hierarchical_allreduce,ring_1d,6,128,256,4096,2723.693333333349
lrab_hierarchical_allreduce,ring_1d,6,512,1024,16384,3048.635000000021
lrab_hierarchical_allreduce,ring_1d,6,1024,2048,32768,3393.4016666666957
lrab_hierarchical_allreduce,ring_1d,6,2048,4096,65536,4082.401666666714
lrab_hierarchical_allreduce,ring_1d,6,4096,8192,131072,5458.80166666677
lrab_hierarchical_allreduce,ring_1d,6,8192,16384,262144,8216.934999999943
lrab_hierarchical_allreduce,ring_1d,6,16384,32768,524288,13733.201666665835
lrab_hierarchical_allreduce,ring_1d,6,32768,65536,1048576,24765.73500000064
lrab_hierarchical_allreduce,ring_1d,6,49152,98304,1572864,35798.268333331536
lrab_hierarchical_allreduce,torus_2d,6,8,16,256,1700.6025000000095
lrab_hierarchical_allreduce,torus_2d,6,32,64,1024,1753.2900000000102
lrab_hierarchical_allreduce,torus_2d,6,64,128,2048,1823.540000000012
lrab_hierarchical_allreduce,torus_2d,6,128,256,4096,1964.040000000012
lrab_hierarchical_allreduce,torus_2d,6,512,1024,16384,2196.8183333333463
lrab_hierarchical_allreduce,torus_2d,6,1024,2048,32768,2477.2783333333473
lrab_hierarchical_allreduce,torus_2d,6,2048,4096,65536,3038.1983333333583
lrab_hierarchical_allreduce,torus_2d,6,4096,8192,131072,4159.5050000000665
lrab_hierarchical_allreduce,torus_2d,6,8192,16384,262144,6403.185000000109
lrab_hierarchical_allreduce,torus_2d,6,16384,32768,524288,10890.5449999995
lrab_hierarchical_allreduce,torus_2d,6,32768,65536,1048576,19865.265000000378
lrab_hierarchical_allreduce,torus_2d,6,49152,98304,1572864,28839.98500000059
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8,16,256,2918.5525000000157
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32,64,1024,2999.740000000016
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,64,128,2048,3107.990000000019
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,128,256,4096,3324.4900000000207
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,512,1024,16384,3589.1133333333582
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,1024,2048,32768,3960.0333333333692
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,2048,4096,65536,4701.873333333393
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,4096,8192,131072,6185.020000000124
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8192,16384,262144,9152.379999999861
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,16384,32768,524288,15087.099999999224
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32768,65536,1048576,26956.540000000765
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,49152,98304,1572864,38825.97999999701
lrab_hierarchical_allreduce,ring_1d,6,8,16,256,2628.2558333333477
lrab_hierarchical_allreduce,ring_1d,6,32,64,1024,2699.943333333348
lrab_hierarchical_allreduce,ring_1d,6,64,128,2048,2795.526666666683
lrab_hierarchical_allreduce,ring_1d,6,128,256,4096,2986.693333333351
lrab_hierarchical_allreduce,ring_1d,6,512,1024,16384,3311.635000000021
lrab_hierarchical_allreduce,ring_1d,6,1024,2048,32768,3656.4016666666957
lrab_hierarchical_allreduce,ring_1d,6,2048,4096,65536,4345.401666666714
lrab_hierarchical_allreduce,ring_1d,6,4096,8192,131072,5721.801666666768
lrab_hierarchical_allreduce,ring_1d,6,8192,16384,262144,8479.934999999887
lrab_hierarchical_allreduce,ring_1d,6,16384,32768,524288,13996.201666665835
lrab_hierarchical_allreduce,ring_1d,6,32768,65536,1048576,25028.73500000064
lrab_hierarchical_allreduce,ring_1d,6,49152,98304,1572864,36061.26833333154
lrab_hierarchical_allreduce,torus_2d,6,8,16,256,1925.6025000000104
lrab_hierarchical_allreduce,torus_2d,6,32,64,1024,1978.290000000011
lrab_hierarchical_allreduce,torus_2d,6,64,128,2048,2048.540000000012
lrab_hierarchical_allreduce,torus_2d,6,128,256,4096,2189.040000000012
lrab_hierarchical_allreduce,torus_2d,6,512,1024,16384,2421.8183333333463
lrab_hierarchical_allreduce,torus_2d,6,1024,2048,32768,2702.2783333333473
lrab_hierarchical_allreduce,torus_2d,6,2048,4096,65536,3263.1983333333583
lrab_hierarchical_allreduce,torus_2d,6,4096,8192,131072,4384.5050000000665
lrab_hierarchical_allreduce,torus_2d,6,8192,16384,262144,6628.18500000005
lrab_hierarchical_allreduce,torus_2d,6,16384,32768,524288,11115.5449999995
lrab_hierarchical_allreduce,torus_2d,6,32768,65536,1048576,20090.265000000378
lrab_hierarchical_allreduce,torus_2d,6,49152,98304,1572864,29064.985000000597
1 algorithm sip_topology n_sips n_elem bytes_per_pe bytes_per_sip latency_ns
2 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 8 16 256 2666.552500000015 2918.5525000000157
3 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 32 64 1024 2747.7400000000152 2999.740000000016
4 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 64 128 2048 2855.990000000018 3107.990000000019
5 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 128 256 4096 3072.490000000019 3324.4900000000207
6 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 512 1024 16384 3337.1133333333582 3589.1133333333582
7 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 1024 2048 32768 3708.0333333333692 3960.0333333333692
8 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 2048 4096 65536 4449.873333333393 4701.873333333393
9 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 4096 8192 131072 5933.020000000124 6185.020000000124
10 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 8192 16384 262144 8900.379999999863 9152.379999999861
11 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 16384 32768 524288 14835.099999999224 15087.099999999224
12 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 32768 65536 1048576 26704.540000000765 26956.540000000765
13 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 49152 98304 1572864 38573.97999999701 38825.97999999701
14 lrab_hierarchical_allreduce ring_1d 6 8 16 256 2365.255833333347 2628.2558333333477
15 lrab_hierarchical_allreduce ring_1d 6 32 64 1024 2436.9433333333473 2699.943333333348
16 lrab_hierarchical_allreduce ring_1d 6 64 128 2048 2532.526666666683 2795.526666666683
17 lrab_hierarchical_allreduce ring_1d 6 128 256 4096 2723.693333333349 2986.693333333351
18 lrab_hierarchical_allreduce ring_1d 6 512 1024 16384 3048.635000000021 3311.635000000021
19 lrab_hierarchical_allreduce ring_1d 6 1024 2048 32768 3393.4016666666957 3656.4016666666957
20 lrab_hierarchical_allreduce ring_1d 6 2048 4096 65536 4082.401666666714 4345.401666666714
21 lrab_hierarchical_allreduce ring_1d 6 4096 8192 131072 5458.80166666677 5721.801666666768
22 lrab_hierarchical_allreduce ring_1d 6 8192 16384 262144 8216.934999999943 8479.934999999887
23 lrab_hierarchical_allreduce ring_1d 6 16384 32768 524288 13733.201666665835 13996.201666665835
24 lrab_hierarchical_allreduce ring_1d 6 32768 65536 1048576 24765.73500000064 25028.73500000064
25 lrab_hierarchical_allreduce ring_1d 6 49152 98304 1572864 35798.268333331536 36061.26833333154
26 lrab_hierarchical_allreduce torus_2d 6 8 16 256 1700.6025000000095 1925.6025000000104
27 lrab_hierarchical_allreduce torus_2d 6 32 64 1024 1753.2900000000102 1978.290000000011
28 lrab_hierarchical_allreduce torus_2d 6 64 128 2048 1823.540000000012 2048.540000000012
29 lrab_hierarchical_allreduce torus_2d 6 128 256 4096 1964.040000000012 2189.040000000012
30 lrab_hierarchical_allreduce torus_2d 6 512 1024 16384 2196.8183333333463 2421.8183333333463
31 lrab_hierarchical_allreduce torus_2d 6 1024 2048 32768 2477.2783333333473 2702.2783333333473
32 lrab_hierarchical_allreduce torus_2d 6 2048 4096 65536 3038.1983333333583 3263.1983333333583
33 lrab_hierarchical_allreduce torus_2d 6 4096 8192 131072 4159.5050000000665 4384.5050000000665
34 lrab_hierarchical_allreduce torus_2d 6 8192 16384 262144 6403.185000000109 6628.18500000005
35 lrab_hierarchical_allreduce torus_2d 6 16384 32768 524288 10890.5449999995 11115.5449999995
36 lrab_hierarchical_allreduce torus_2d 6 32768 65536 1048576 19865.265000000378 20090.265000000378
37 lrab_hierarchical_allreduce torus_2d 6 49152 98304 1572864 28839.98500000059 29064.985000000597
docs/diagrams/allreduce_latency_plots/topology.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 194 KiB

After

Width:  |  Height:  |  Size: 194 KiB

docs/diagrams/gemm_plots/gemm_mac_utilization_measured.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 40 KiB

After

Width:  |  Height:  |  Size: 29 KiB

docs/diagrams/gemm_plots/gemm_mac_utilization_theoretical_vs_measured.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 46 KiB

After

Width:  |  Height:  |  Size: 32 KiB

docs/diagrams/gemm_plots/gemm_stage_breakdown.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 42 KiB

After

Width:  |  Height:  |  Size: 31 KiB

docs/diagrams/pe2pe_latency_plots/latency_intercube_C0PE0_to_C1PE0_horizontal.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 53 KiB

After

Width:  |  Height:  |  Size: 52 KiB

docs/diagrams/pe2pe_latency_plots/latency_intercube_C0PE0_to_C4PE0_vertical.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 52 KiB

After

Width:  |  Height:  |  Size: 45 KiB

docs/diagrams/pe2pe_latency_plots/latency_intracube_PE0_to_PE1_horizontal.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 51 KiB

After

Width:  |  Height:  |  Size: 53 KiB

docs/diagrams/pe2pe_latency_plots/latency_intracube_PE0_to_PE4_vertical.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 52 KiB

After

Width:  |  Height:  |  Size: 47 KiB

docs/diagrams/pe2pe_latency_plots/overview.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 137 KiB

After

Width:  |  Height:  |  Size: 137 KiB

docs/diagrams/pe2pe_latency_plots/summary.csv
+80 -80
View File
@@ -1,81 +1,81 @@
hop,label,size_bytes,path,total_ns
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),128,ipcq,24.88749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),128,raw,33.57999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),256,ipcq,28.13749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),256,raw,36.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),384,ipcq,29.88749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),384,raw,37.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),512,ipcq,31.63749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),512,raw,38.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),768,ipcq,35.13749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),768,raw,40.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),1024,ipcq,38.63749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),1024,raw,42.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),2048,ipcq,52.63749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),2048,raw,50.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),4096,ipcq,80.63750000000073
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),4096,raw,66.08000000000175
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),8192,ipcq,136.63750000000073
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),8192,raw,98.08000000000175
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),10240,ipcq,164.63750000000073
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),10240,raw,114.08000000000175
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),128,ipcq,38.49749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),128,raw,47.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),256,ipcq,43.24749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),256,raw,51.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),384,ipcq,44.99749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),384,raw,52.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),512,ipcq,46.74749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),512,raw,53.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),768,ipcq,50.24749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),768,raw,55.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),1024,ipcq,53.74749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),1024,raw,57.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),2048,ipcq,67.74749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),2048,raw,65.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),4096,ipcq,95.74750000000131
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),4096,raw,81.19000000000233
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),8192,ipcq,151.7475000000013
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),8192,raw,113.19000000000233
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),10240,ipcq,179.7475000000013
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),10240,raw,129.19000000000233
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),128,ipcq,81.15999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),128,raw,89.28999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),256,ipcq,88.65999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),256,raw,95.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),384,ipcq,90.90999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),384,raw,96.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),512,ipcq,93.15999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),512,raw,97.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),768,ipcq,97.65999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),768,raw,99.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),1024,ipcq,103.15999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),1024,raw,102.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),2048,ipcq,125.15999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),2048,raw,114.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),4096,ipcq,169.15999999999985
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),4096,raw,138.54000000000087
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),8192,ipcq,257.15999999999985
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),8192,raw,186.54000000000087
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),10240,ipcq,301.15999999999985
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),10240,raw,210.54000000000087
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),128,ipcq,103.15999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),128,raw,111.28999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),256,ipcq,112.65999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),256,raw,119.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),384,ipcq,114.90999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),384,raw,120.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),512,ipcq,117.15999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),512,raw,121.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),768,ipcq,121.65999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),768,raw,123.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),1024,ipcq,127.15999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),1024,raw,126.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),2048,ipcq,149.15999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),2048,raw,138.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),4096,ipcq,193.15999999999985
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),4096,raw,162.54000000000087
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),8192,ipcq,281.15999999999985
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),8192,raw,210.54000000000087
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),10240,ipcq,325.15999999999985
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),10240,raw,234.54000000000087
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),128,ipcq,42.88749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),128,raw,51.57999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),256,ipcq,46.13749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),256,raw,54.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),384,ipcq,47.88749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),384,raw,55.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),512,ipcq,49.63749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),512,raw,56.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),768,ipcq,53.13749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),768,raw,58.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),1024,ipcq,56.63749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),1024,raw,60.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),2048,ipcq,70.63749999999891
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),2048,raw,68.07999999999811
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),4096,ipcq,98.63750000000073
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),4096,raw,84.08000000000175
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),8192,ipcq,154.63750000000073
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),8192,raw,116.08000000000175
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),10240,ipcq,182.63750000000073
latency_intracube_PE0_to_PE1_horizontal,Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal),10240,raw,132.08000000000175
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),128,ipcq,56.49749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),128,raw,65.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),256,ipcq,61.24749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),256,raw,69.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),384,ipcq,62.99749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),384,raw,70.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),512,ipcq,64.74749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),512,raw,71.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),768,ipcq,68.24749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),768,raw,73.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),1024,ipcq,71.74749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),1024,raw,75.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),2048,ipcq,85.74749999999585
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),2048,raw,83.18999999999505
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),4096,ipcq,113.74750000000131
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),4096,raw,99.19000000000233
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),8192,ipcq,169.7475000000013
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),8192,raw,131.19000000000233
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),10240,ipcq,197.7475000000013
latency_intracube_PE0_to_PE4_vertical,Intra-cube PE-to-PE latency: PE0 → PE4 (vertical),10240,raw,147.19000000000233
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),128,ipcq,99.15999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),128,raw,107.28999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),256,ipcq,106.65999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),256,raw,113.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),384,ipcq,108.90999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),384,raw,114.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),512,ipcq,111.15999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),512,raw,115.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),768,ipcq,115.65999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),768,raw,117.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),1024,ipcq,121.15999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),1024,raw,120.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),2048,ipcq,143.15999999999804
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),2048,raw,132.53999999999724
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),4096,ipcq,187.15999999999985
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),4096,raw,156.54000000000087
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),8192,ipcq,275.15999999999985
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),8192,raw,204.54000000000087
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),10240,ipcq,319.15999999999985
latency_intercube_C0PE0_to_C1PE0_horizontal,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal),10240,raw,228.54000000000087
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),128,ipcq,121.15999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),128,raw,129.28999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),256,ipcq,130.65999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),256,raw,137.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),384,ipcq,132.90999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),384,raw,138.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),512,ipcq,135.15999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),512,raw,139.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),768,ipcq,139.65999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),768,raw,141.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),1024,ipcq,145.15999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),1024,raw,144.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),2048,ipcq,167.15999999999804
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),2048,raw,156.53999999999724
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),4096,ipcq,211.15999999999985
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),4096,raw,180.54000000000087
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),8192,ipcq,299.15999999999985
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),8192,raw,228.54000000000087
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),10240,ipcq,343.15999999999985
latency_intercube_C0PE0_to_C4PE0_vertical,Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical),10240,raw,252.54000000000087
1 hop label size_bytes path total_ns
2 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 128 ipcq 24.88749999999891 42.88749999999891
3 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 128 raw 33.57999999999811 51.57999999999811
4 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 256 ipcq 28.13749999999891 46.13749999999891
5 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 256 raw 36.07999999999811 54.07999999999811
6 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 384 ipcq 29.88749999999891 47.88749999999891
7 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 384 raw 37.07999999999811 55.07999999999811
8 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 512 ipcq 31.63749999999891 49.63749999999891
9 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 512 raw 38.07999999999811 56.07999999999811
10 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 768 ipcq 35.13749999999891 53.13749999999891
11 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 768 raw 40.07999999999811 58.07999999999811
12 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 1024 ipcq 38.63749999999891 56.63749999999891
13 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 1024 raw 42.07999999999811 60.07999999999811
14 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 2048 ipcq 52.63749999999891 70.63749999999891
15 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 2048 raw 50.07999999999811 68.07999999999811
16 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 4096 ipcq 80.63750000000073 98.63750000000073
17 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 4096 raw 66.08000000000175 84.08000000000175
18 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 8192 ipcq 136.63750000000073 154.63750000000073
19 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 8192 raw 98.08000000000175 116.08000000000175
20 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 10240 ipcq 164.63750000000073 182.63750000000073
21 latency_intracube_PE0_to_PE1_horizontal Intra-cube PE-to-PE latency: PE0 → PE1 (horizontal) 10240 raw 114.08000000000175 132.08000000000175
22 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 128 ipcq 38.49749999999585 56.49749999999585
23 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 128 raw 47.18999999999505 65.18999999999505
24 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 256 ipcq 43.24749999999585 61.24749999999585
25 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 256 raw 51.18999999999505 69.18999999999505
26 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 384 ipcq 44.99749999999585 62.99749999999585
27 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 384 raw 52.18999999999505 70.18999999999505
28 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 512 ipcq 46.74749999999585 64.74749999999585
29 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 512 raw 53.18999999999505 71.18999999999505
30 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 768 ipcq 50.24749999999585 68.24749999999585
31 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 768 raw 55.18999999999505 73.18999999999505
32 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 1024 ipcq 53.74749999999585 71.74749999999585
33 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 1024 raw 57.18999999999505 75.18999999999505
34 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 2048 ipcq 67.74749999999585 85.74749999999585
35 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 2048 raw 65.18999999999505 83.18999999999505
36 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 4096 ipcq 95.74750000000131 113.74750000000131
37 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 4096 raw 81.19000000000233 99.19000000000233
38 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 8192 ipcq 151.7475000000013 169.7475000000013
39 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 8192 raw 113.19000000000233 131.19000000000233
40 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 10240 ipcq 179.7475000000013 197.7475000000013
41 latency_intracube_PE0_to_PE4_vertical Intra-cube PE-to-PE latency: PE0 → PE4 (vertical) 10240 raw 129.19000000000233 147.19000000000233
42 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 128 ipcq 81.15999999999804 99.15999999999804
43 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 128 raw 89.28999999999724 107.28999999999724
44 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 256 ipcq 88.65999999999804 106.65999999999804
45 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 256 raw 95.53999999999724 113.53999999999724
46 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 384 ipcq 90.90999999999804 108.90999999999804
47 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 384 raw 96.53999999999724 114.53999999999724
48 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 512 ipcq 93.15999999999804 111.15999999999804
49 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 512 raw 97.53999999999724 115.53999999999724
50 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 768 ipcq 97.65999999999804 115.65999999999804
51 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 768 raw 99.53999999999724 117.53999999999724
52 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 1024 ipcq 103.15999999999804 121.15999999999804
53 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 1024 raw 102.53999999999724 120.53999999999724
54 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 2048 ipcq 125.15999999999804 143.15999999999804
55 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 2048 raw 114.53999999999724 132.53999999999724
56 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 4096 ipcq 169.15999999999985 187.15999999999985
57 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 4096 raw 138.54000000000087 156.54000000000087
58 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 8192 ipcq 257.15999999999985 275.15999999999985
59 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 8192 raw 186.54000000000087 204.54000000000087
60 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 10240 ipcq 301.15999999999985 319.15999999999985
61 latency_intercube_C0PE0_to_C1PE0_horizontal Inter-cube PE-to-PE latency: Cube0.PE0 → Cube1.PE0 (horizontal) 10240 raw 210.54000000000087 228.54000000000087
62 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 128 ipcq 103.15999999999804 121.15999999999804
63 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 128 raw 111.28999999999724 129.28999999999724
64 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 256 ipcq 112.65999999999804 130.65999999999804
65 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 256 raw 119.53999999999724 137.53999999999724
66 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 384 ipcq 114.90999999999804 132.90999999999804
67 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 384 raw 120.53999999999724 138.53999999999724
68 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 512 ipcq 117.15999999999804 135.15999999999804
69 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 512 raw 121.53999999999724 139.53999999999724
70 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 768 ipcq 121.65999999999804 139.65999999999804
71 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 768 raw 123.53999999999724 141.53999999999724
72 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 1024 ipcq 127.15999999999804 145.15999999999804
73 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 1024 raw 126.53999999999724 144.53999999999724
74 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 2048 ipcq 149.15999999999804 167.15999999999804
75 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 2048 raw 138.53999999999724 156.53999999999724
76 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 4096 ipcq 193.15999999999985 211.15999999999985
77 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 4096 raw 162.54000000000087 180.54000000000087
78 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 8192 ipcq 281.15999999999985 299.15999999999985
79 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 8192 raw 210.54000000000087 228.54000000000087
80 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 10240 ipcq 325.15999999999985 343.15999999999985
81 latency_intercube_C0PE0_to_C4PE0_vertical Inter-cube PE-to-PE latency: Cube0.PE0 → Cube4.PE0 (vertical) 10240 raw 234.54000000000087 252.54000000000087
docs/report/1H-codesign-paper/build/main.pdf
BIN
View File
Binary file not shown.
docs/report/1H-codesign-paper/figures/allreduce_buffer_kind.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 75 KiB

After

Width:  |  Height:  |  Size: 77 KiB

docs/report/1H-codesign-paper/figures/allreduce_comparison.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 86 KiB

After

Width:  |  Height:  |  Size: 92 KiB

docs/report/1H-codesign-paper/figures/cube_architecture.pdf
BIN
View File
Binary file not shown.
docs/report/1H-codesign-paper/figures/cube_architecture.svg
+312
View File
@@ -0,0 +1,312 @@
<svg xmlns="http://www.w3.org/2000/svg" width="860" height="798" viewBox="55 2 860 798">
<title>cube</title>
<rect width="970" height="900" fill="#ffffff"/>
<text x="485" y="22" text-anchor="middle" font-family="monospace" font-size="18" font-weight="bold" fill="#1f2937">CUBE TOPOLOGY — 17.0×14.0mm | 6×6 Router Mesh | n_to_one mode | 64 pseudo-ch</text>
<text x="485" y="40" text-anchor="middle" font-family="monospace" font-size="15" fill="#ffffff">Per-PE: 8 ch × 32.0 GB/s = 256.0 GB/s | Cube total: 64 × 32.0 = 2048.0 GB/s</text>
<rect x="60" y="60" width="850.0" height="700.0" rx="6" fill="none" stroke="#94a3b8" stroke-width="2" stroke-dasharray="8,4"/>
<rect x="260" y="285" width="450" height="250" rx="6" fill="#ecfdf5" stroke="#047857" stroke-width="2" opacity="0.6"/>
<text x="485" y="395" text-anchor="middle" font-family="monospace" font-size="16" font-weight="bold" fill="#047857">HBM_CTRL | 64 pseudo channels</text>
<text x="485" y="412" text-anchor="middle" font-family="monospace" font-size="14" fill="#059669">Total BW: 2048 GB/s</text>
<rect x="270.0" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
<rect x="283.4" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
<rect x="296.9" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
<rect x="310.3" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
<rect x="323.8" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
<rect x="337.2" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
<rect x="350.6" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
<rect x="364.1" y="289" width="12.9" height="8" rx="1" fill="#3b82f6" opacity="0.8"/>
<rect x="377.5" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
<rect x="390.9" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
<rect x="404.4" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
<rect x="417.8" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
<rect x="431.2" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
<rect x="444.7" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
<rect x="458.1" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
<rect x="471.6" y="289" width="12.9" height="8" rx="1" fill="#60a5fa" opacity="0.8"/>
<rect x="485.0" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
<rect x="498.4" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
<rect x="511.9" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
<rect x="525.3" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
<rect x="538.8" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
<rect x="552.2" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
<rect x="565.6" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
<rect x="579.1" y="289" width="12.9" height="8" rx="1" fill="#475569" opacity="0.8"/>
<rect x="592.5" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
<rect x="605.9" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
<rect x="619.4" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
<rect x="632.8" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
<rect x="646.2" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
<rect x="659.7" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
<rect x="673.1" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
<rect x="686.6" y="289" width="12.9" height="8" rx="1" fill="#94a3b8" opacity="0.8"/>
<text x="324" y="286" text-anchor="middle" font-family="monospace" font-size="12" fill="#3b82f6">PE0×8ch</text>
<text x="431" y="286" text-anchor="middle" font-family="monospace" font-size="12" fill="#60a5fa">PE1×8ch</text>
<text x="539" y="286" text-anchor="middle" font-family="monospace" font-size="12" fill="#475569">PE2×8ch</text>
<text x="646" y="286" text-anchor="middle" font-family="monospace" font-size="12" fill="#94a3b8">PE3×8ch</text>
<rect x="270.0" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
<rect x="283.4" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
<rect x="296.9" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
<rect x="310.3" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
<rect x="323.8" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
<rect x="337.2" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
<rect x="350.6" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
<rect x="364.1" y="523" width="12.9" height="8" rx="1" fill="#f59e0b" opacity="0.8"/>
<rect x="377.5" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
<rect x="390.9" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
<rect x="404.4" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
<rect x="417.8" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
<rect x="431.2" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
<rect x="444.7" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
<rect x="458.1" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
<rect x="471.6" y="523" width="12.9" height="8" rx="1" fill="#fbbf24" opacity="0.8"/>
<rect x="485.0" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
<rect x="498.4" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
<rect x="511.9" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
<rect x="525.3" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
<rect x="538.8" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
<rect x="552.2" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
<rect x="565.6" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
<rect x="579.1" y="523" width="12.9" height="8" rx="1" fill="#ef4444" opacity="0.8"/>
<rect x="592.5" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
<rect x="605.9" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
<rect x="619.4" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
<rect x="632.8" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
<rect x="646.2" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
<rect x="659.7" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
<rect x="673.1" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
<rect x="686.6" y="523" width="12.9" height="8" rx="1" fill="#f87171" opacity="0.8"/>
<text x="324" y="539" text-anchor="middle" font-family="monospace" font-size="12" fill="#f59e0b">PE4×8ch</text>
<text x="431" y="539" text-anchor="middle" font-family="monospace" font-size="12" fill="#fbbf24">PE5×8ch</text>
<text x="539" y="539" text-anchor="middle" font-family="monospace" font-size="12" fill="#ef4444">PE6×8ch</text>
<text x="646" y="539" text-anchor="middle" font-family="monospace" font-size="12" fill="#f87171">PE7×8ch</text>
<line x1="135" y1="135" x2="285" y2="135" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="135" y1="135" x2="135" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="285" y1="135" x2="435" y2="135" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="285" y1="135" x2="285" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="435" y1="135" x2="585" y2="135" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="435" y1="135" x2="435" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="585" y1="135" x2="685" y2="135" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="585" y1="135" x2="585" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="685" y1="135" x2="835" y2="135" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="685" y1="135" x2="685" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="835" y1="135" x2="835" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="135" y1="260" x2="285" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="135" y1="260" x2="135" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="285" y1="260" x2="435" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="285" y1="260" x2="285" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="435" y1="260" x2="585" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="435" y1="260" x2="435" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="585" y1="260" x2="685" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="585" y1="260" x2="585" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="685" y1="260" x2="835" y2="260" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="685" y1="260" x2="685" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="835" y1="260" x2="835" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="135" y1="335" x2="285" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="135" y1="335" x2="135" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="285" y1="335" x2="685" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="285" y1="335" x2="285" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="685" y1="335" x2="835" y2="335" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="685" y1="335" x2="685" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="835" y1="335" x2="835" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="135" y1="485" x2="285" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="135" y1="485" x2="135" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="285" y1="485" x2="685" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="285" y1="485" x2="285" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="685" y1="485" x2="835" y2="485" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="685" y1="485" x2="685" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="835" y1="485" x2="835" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="135" y1="560" x2="285" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="135" y1="560" x2="135" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="285" y1="560" x2="435" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="285" y1="560" x2="285" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="435" y1="560" x2="585" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="435" y1="560" x2="435" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="585" y1="560" x2="685" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="585" y1="560" x2="585" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="685" y1="560" x2="835" y2="560" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="685" y1="560" x2="685" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="835" y1="560" x2="835" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="135" y1="685" x2="285" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="285" y1="685" x2="435" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="435" y1="685" x2="585" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="585" y1="685" x2="685" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<line x1="685" y1="685" x2="835" y2="685" stroke="#94a3b8" stroke-width="1" opacity="0.4"/>
<circle cx="135" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="135" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c0</text>
<rect x="119" y="81" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
<text x="135" y="92" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE0</text>
<line x1="135" y1="127" x2="149" y2="97" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
<circle cx="285" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="285" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c1</text>
<rect x="269" y="81" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
<text x="285" y="92" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE1</text>
<line x1="285" y1="127" x2="299" y2="97" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
<circle cx="435" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="435" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c2</text>
<circle cx="585" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="585" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c3</text>
<circle cx="685" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="685" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c4</text>
<circle cx="835" cy="135" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="835" y="138" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r0c5</text>
<circle cx="135" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="135" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c0</text>
<circle cx="285" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="285" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c1</text>
<circle cx="435" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="435" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c2</text>
<rect x="419" y="206" width="32" height="16" rx="3" fill="#ffffff" stroke="#f59e0b" stroke-width="1"/>
<text x="435" y="217" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#f59e0b">M_CPU</text>
<line x1="435" y1="252" x2="449" y2="222" stroke="#f59e0b" stroke-width="1" opacity="0.6"/>
<circle cx="585" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="585" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c3</text>
<circle cx="685" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="685" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c4</text>
<rect x="669" y="206" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
<text x="685" y="217" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE2</text>
<line x1="685" y1="252" x2="699" y2="222" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
<circle cx="835" cy="260" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="835" y="263" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r1c5</text>
<rect x="819" y="206" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
<text x="835" y="217" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE3</text>
<line x1="835" y1="252" x2="849" y2="222" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
<circle cx="135" cy="335" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="135" y="338" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r2c0</text>
<circle cx="285" cy="335" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="285" y="338" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r2c1</text>
<circle cx="685" cy="335" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="685" y="338" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r2c4</text>
<circle cx="835" cy="335" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="835" y="338" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r2c5</text>
<circle cx="135" cy="485" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="135" y="488" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r3c0</text>
<rect x="119" y="523" width="32" height="16" rx="3" fill="#ffffff" stroke="#d97706" stroke-width="1"/>
<text x="135" y="534" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#d97706">SRAM</text>
<line x1="135" y1="493" x2="149" y2="523" stroke="#d97706" stroke-width="1" opacity="0.6"/>
<circle cx="285" cy="485" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="285" y="488" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r3c1</text>
<circle cx="685" cy="485" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="685" y="488" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r3c4</text>
<circle cx="835" cy="485" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="835" y="488" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r3c5</text>
<circle cx="135" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="135" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c0</text>
<rect x="119" y="598" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
<text x="135" y="609" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE4</text>
<line x1="135" y1="568" x2="149" y2="598" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
<circle cx="285" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="285" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c1</text>
<rect x="269" y="598" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
<text x="285" y="609" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE5</text>
<line x1="285" y1="568" x2="299" y2="598" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
<circle cx="435" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="435" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c2</text>
<circle cx="585" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="585" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c3</text>
<circle cx="685" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="685" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c4</text>
<circle cx="835" cy="560" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="835" y="563" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r4c5</text>
<circle cx="135" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="135" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c0</text>
<circle cx="285" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="285" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c1</text>
<circle cx="435" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="435" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c2</text>
<circle cx="585" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="585" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c3</text>
<circle cx="685" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="685" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c4</text>
<rect x="669" y="723" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
<text x="685" y="734" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE6</text>
<line x1="685" y1="693" x2="699" y2="723" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
<circle cx="835" cy="685" r="17" fill="#ffffff" stroke="#94a3b8" stroke-width="1"/>
<text x="835" y="688" text-anchor="middle" font-family="monospace" font-size="12" fill="#1f2937">r5c5</text>
<rect x="819" y="723" width="32" height="16" rx="3" fill="#ffffff" stroke="#a855f7" stroke-width="1"/>
<text x="835" y="734" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#a855f7">PE7</text>
<line x1="835" y1="693" x2="849" y2="723" stroke="#a855f7" stroke-width="1" opacity="0.6"/>
<polyline points="135,143 208,216 251,216 324,289" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
<text x="239" y="216" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
<polyline points="285,143 358,216 358,216 431,289" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
<text x="368" y="216" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
<polyline points="685,268 674,278 549,278 539,289" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
<text x="622" y="278" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
<polyline points="835,268 824,278 657,278 646,289" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
<text x="751" y="278" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
<polyline points="135,552 146,542 313,542 324,531" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
<text x="239" y="542" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
<polyline points="285,552 296,542 421,542 431,531" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
<text x="368" y="542" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
<polyline points="685,677 612,604 612,604 539,531" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
<text x="622" y="604" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
<polyline points="835,677 762,604 719,604 646,531" fill="none" stroke="#10b981" stroke-width="1.5" opacity="0.6" stroke-dasharray="4,3"/>
<text x="751" y="604" font-family="monospace" font-size="12" fill="#047857">256GB/s</text>
<rect x="65" y="360" width="50" height="100" rx="3" fill="#ffffff" stroke="#475569" stroke-width="1.5" opacity="0.9"/>
<text x="90" y="357" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#475569">UCIe-W</text>
<rect x="67" y="362" width="46" height="23" rx="2" fill="#cbd5e1" opacity="0.7"/>
<text x="90" y="376" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c0</text>
<polyline points="127,135 120,142 120,366 113,374" fill="none" stroke="#cbd5e1" stroke-width="1" opacity="0.5"/>
<rect x="67" y="386" width="46" height="23" rx="2" fill="#94a3b8" opacity="0.7"/>
<text x="90" y="400" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c1</text>
<polyline points="127,260 120,267 120,390 113,398" fill="none" stroke="#94a3b8" stroke-width="1" opacity="0.5"/>
<rect x="67" y="410" width="46" height="23" rx="2" fill="#64748b" opacity="0.7"/>
<text x="90" y="424" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c2</text>
<polyline points="127,560 120,553 120,428 113,422" fill="none" stroke="#64748b" stroke-width="1" opacity="0.5"/>
<rect x="67" y="434" width="46" height="23" rx="2" fill="#374151" opacity="0.7"/>
<text x="90" y="448" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c3</text>
<polyline points="127,685 120,678 120,452 113,446" fill="none" stroke="#374151" stroke-width="1" opacity="0.5"/>
<rect x="435" y="65" width="100" height="50" rx="3" fill="#ffffff" stroke="#475569" stroke-width="1.5" opacity="0.9"/>
<text x="485" y="62" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#475569">UCIe-N</text>
<rect x="437" y="67" width="23" height="46" rx="2" fill="#cbd5e1" opacity="0.7"/>
<text x="448" y="93" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c0</text>
<polyline points="135,127 142,120 442,120 448,113" fill="none" stroke="#cbd5e1" stroke-width="1" opacity="0.5"/>
<rect x="461" y="67" width="23" height="46" rx="2" fill="#94a3b8" opacity="0.7"/>
<text x="472" y="93" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c1</text>
<polyline points="285,127 292,120 466,120 472,113" fill="none" stroke="#94a3b8" stroke-width="1" opacity="0.5"/>
<rect x="485" y="67" width="23" height="46" rx="2" fill="#64748b" opacity="0.7"/>
<text x="496" y="93" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c2</text>
<polyline points="685,127 678,120 504,120 496,113" fill="none" stroke="#64748b" stroke-width="1" opacity="0.5"/>
<rect x="509" y="67" width="23" height="46" rx="2" fill="#374151" opacity="0.7"/>
<text x="520" y="93" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c3</text>
<polyline points="835,127 828,120 528,120 520,113" fill="none" stroke="#374151" stroke-width="1" opacity="0.5"/>
<rect x="855" y="360" width="50" height="100" rx="3" fill="#ffffff" stroke="#475569" stroke-width="1.5" opacity="0.9"/>
<text x="880" y="357" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#475569">UCIe-E</text>
<rect x="857" y="362" width="46" height="23" rx="2" fill="#cbd5e1" opacity="0.7"/>
<text x="880" y="376" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c0</text>
<polyline points="843,135 850,142 850,367 857,374" fill="none" stroke="#cbd5e1" stroke-width="1" opacity="0.5"/>
<rect x="857" y="386" width="46" height="23" rx="2" fill="#94a3b8" opacity="0.7"/>
<text x="880" y="400" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c1</text>
<polyline points="843,260 850,267 850,391 857,398" fill="none" stroke="#94a3b8" stroke-width="1" opacity="0.5"/>
<rect x="857" y="410" width="46" height="23" rx="2" fill="#64748b" opacity="0.7"/>
<text x="880" y="424" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c2</text>
<polyline points="843,560 850,553 850,428 857,422" fill="none" stroke="#64748b" stroke-width="1" opacity="0.5"/>
<rect x="857" y="434" width="46" height="23" rx="2" fill="#374151" opacity="0.7"/>
<text x="880" y="448" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c3</text>
<polyline points="843,685 850,678 850,452 857,446" fill="none" stroke="#374151" stroke-width="1" opacity="0.5"/>
<rect x="435" y="705" width="100" height="50" rx="3" fill="#ffffff" stroke="#475569" stroke-width="1.5" opacity="0.9"/>
<text x="485" y="702" text-anchor="middle" font-family="monospace" font-size="14" font-weight="bold" fill="#475569">UCIe-S</text>
<rect x="437" y="707" width="23" height="46" rx="2" fill="#cbd5e1" opacity="0.7"/>
<text x="448" y="733" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c0</text>
<polyline points="135,693 142,700 442,700 448,707" fill="none" stroke="#cbd5e1" stroke-width="1" opacity="0.5"/>
<rect x="461" y="707" width="23" height="46" rx="2" fill="#94a3b8" opacity="0.7"/>
<text x="472" y="733" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c1</text>
<polyline points="285,693 292,700 466,700 472,707" fill="none" stroke="#94a3b8" stroke-width="1" opacity="0.5"/>
<rect x="485" y="707" width="23" height="46" rx="2" fill="#64748b" opacity="0.7"/>
<text x="496" y="733" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c2</text>
<polyline points="685,693 678,700 504,700 496,707" fill="none" stroke="#64748b" stroke-width="1" opacity="0.5"/>
<rect x="509" y="707" width="23" height="46" rx="2" fill="#374151" opacity="0.7"/>
<text x="520" y="733" text-anchor="middle" font-family="monospace" font-size="10" fill="#1f2937">c3</text>
<polyline points="835,693 828,700 528,700 520,707" fill="none" stroke="#374151" stroke-width="1" opacity="0.5"/>
<rect x="60" y="775" width="10" height="10" rx="2" fill="#3b82f6" stroke="#94a3b8" stroke-width="0.5"/>
<text x="74" y="784" font-family="monospace" font-size="13" fill="#1f2937">PE Router</text>
<rect x="147" y="775" width="10" height="10" rx="2" fill="#f59e0b" stroke="#94a3b8" stroke-width="0.5"/>
<text x="161" y="784" font-family="monospace" font-size="13" fill="#1f2937">M_CPU / SRAM</text>
<rect x="255" y="775" width="10" height="10" rx="2" fill="#475569" stroke="#94a3b8" stroke-width="0.5"/>
<text x="269" y="784" font-family="monospace" font-size="13" fill="#1f2937">UCIe</text>
<rect x="307" y="775" width="10" height="10" rx="2" fill="#1f2937" stroke="#94a3b8" stroke-width="0.5"/>
<text x="321" y="784" font-family="monospace" font-size="13" fill="#1f2937">Relay</text>
<rect x="366" y="775" width="10" height="10" rx="2" fill="#10b981" stroke="#94a3b8" stroke-width="0.5"/>
<text x="380" y="784" font-family="monospace" font-size="13" fill="#1f2937">HBM Link</text>
<rect x="446" y="775" width="10" height="10" rx="2" fill="#1f2937" stroke="#94a3b8" stroke-width="0.5"/>
<text x="460" y="784" font-family="monospace" font-size="13" fill="#1f2937">Mesh Link</text>
</svg>
Side by Side

After

Width:  |  Height:  |  Size: 30 KiB

docs/report/1H-codesign-paper/figures/gemm_composite_vs_async_tflops.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 75 KiB

docs/report/1H-codesign-paper/figures/gemm_hbm_bw_util.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 32 KiB

docs/report/1H-codesign-paper/figures/gemm_mac_utilization_measured.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 40 KiB

After

Width:  |  Height:  |  Size: 29 KiB

docs/report/1H-codesign-paper/figures/gemm_mac_utilization_theoretical_vs_measured.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 46 KiB

After

Width:  |  Height:  |  Size: 32 KiB

docs/report/1H-codesign-paper/figures/gemm_per_pe_tflops.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 29 KiB

docs/report/1H-codesign-paper/figures/gemm_stage_breakdown.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 42 KiB

After

Width:  |  Height:  |  Size: 34 KiB

docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_latency.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 58 KiB

docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_memory.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 51 KiB

docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_parallelism.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 74 KiB

docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_6cases_traffic.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 45 KiB

docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_composite.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 136 KiB

docs/report/1H-codesign-paper/figures/gqa_decode_streaming.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 136 KiB

docs/report/1H-codesign-paper/figures/gqa_long_ctx_6cases_hbm_budget.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 47 KiB

docs/report/1H-codesign-paper/figures/gqa_long_ctx_6cases_kv_sharding_diagram.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 135 KiB

docs/report/1H-codesign-paper/figures/gqa_long_ctx_6cases_kv_sharding_table.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 123 KiB

docs/report/1H-codesign-paper/figures/gqa_long_ctx_6cases_memory_comm_analytical.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 132 KiB

docs/report/1H-codesign-paper/figures/gqa_long_ctx_6cases_memory_comm_paired.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 183 KiB

docs/report/1H-codesign-paper/figures/gqa_long_ctx_6cases_summary.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 182 KiB

docs/report/1H-codesign-paper/figures/gqa_prefill_compute_bound.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 145 KiB

docs/report/1H-codesign-paper/figures/ipcq_alternatives_architecture.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 253 KiB

docs/report/1H-codesign-paper/figures/ipcq_alternatives_architecture_flow.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 305 KiB

docs/report/1H-codesign-paper/figures/ipcq_alternatives_architecture_stacked.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 381 KiB

docs/report/1H-codesign-paper/figures/ipcq_alternatives_decision_matrix.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 189 KiB

docs/report/1H-codesign-paper/figures/ipcq_alternatives_latency.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 160 KiB

docs/report/1H-codesign-paper/figures/latency_model.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 61 KiB

docs/report/1H-codesign-paper/figures/pe_architecture.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 1014 KiB

docs/report/1H-codesign-paper/figures/probe_pe_dma_summary.txt
+57
View File
@@ -0,0 +1,57 @@
=== H2D Write Latency (IO->HBM, data=32768B) ===
Case Target Hops Actual Ovhd Drain Wire Ovhd% Drain% Eff.BW BN.BW Util%
-------------------------------------------------------------------------------------------------------------------
h2d-1hop cube0.pe0 1 289.53 50.0 256.0 0.18 17.3% 88.4% 113.17 128.0 88.4%
h2d-2hop cube4.pe0 2 326.04 78.0 256.0 0.20 23.9% 78.5% 100.50 128.0 78.5%
h2d-3hop cube8.pe0 3 362.56 106.0 256.0 0.20 29.2% 70.6% 90.38 128.0 70.6%
h2d-4hop cube12.pe0 4 399.06 134.0 256.0 0.21 33.6% 64.1% 82.11 128.0 64.1%
[v] Monotonic increase: PASS
BW Saturation (Util% by data size):
Case 4KB 16KB 64KB 256KB 1MB
------------------------------------------------------------------
h2d-1hop 38.9% 71.8% 91.1% 97.6% 99.4%
h2d-2hop 29.0% 62.1% 86.8% 96.3% 99.1%
h2d-3hop 23.2% 54.7% 82.8% 95.1% 98.7%
h2d-4hop 19.3% 48.8% 79.2% 93.8% 98.4%
=== D2H Read Latency (HBM->IO, data=32768B) ===
Case Source Hops Actual Ovhd Drain Wire Ovhd% Drain% Eff.BW BN.BW Util%
-------------------------------------------------------------------------------------------------------------------
d2h-1hop cube0.pe0 1 571.20 23.0 256.0 0.04 4.0% 44.8% 57.37 128.0 44.8%
d2h-2hop cube4.pe0 2 635.72 51.0 256.0 0.05 8.0% 40.3% 51.54 128.0 40.3%
d2h-3hop cube8.pe0 3 700.24 79.0 256.0 0.06 11.3% 36.6% 46.80 128.0 36.6%
d2h-4hop cube12.pe0 4 764.76 107.0 256.0 0.07 14.0% 33.5% 42.85 128.0 33.5%
[v] Monotonic increase: PASS
BW Saturation (Util% by data size):
Case 4KB 16KB 64KB 256KB 1MB
------------------------------------------------------------------
d2h-1hop 9.2% 28.9% 61.9% 86.7% 96.3%
d2h-2hop 7.8% 25.2% 57.4% 84.4% 95.6%
d2h-3hop 6.7% 22.4% 53.5% 82.2% 94.9%
d2h-4hop 5.9% 20.1% 50.2% 80.1% 94.2%
[v] D2H >= H2D (reverse data path): PASS
=== PE DMA Latency (pe_dma -> router -> HBM, data=32768B) ===
Case Target Actual Ovhd Drain Wire Ovhd% Drain% Eff.BW BN.BW Util%
----------------------------------------------------------------------------------------------------------------------------
pe-local-hbm c0.pe0->c0.slice0 141.00 2.0 128.0 0.00 1.4% 90.8% 232.40 256.0 90.8%
pe-same-half-hbm c0.pe0->c0.slice1 147.87 4.0 128.0 0.03 2.7% 86.6% 221.60 256.0 86.6%
pe-cross-half-hbm c0.pe0->c0.slice4 161.17 10.0 128.0 0.09 6.2% 79.4% 203.31 256.0 79.4%
pe-cross-cube-hbm-best c0.pe0->c1.slice0 330.52 30.0 256.0 0.01 9.1% 77.5% 99.14 128.0 77.5%
pe-cross-cube-hbm-worst c0.pe0->c15.slice0 677.12 180.0 256.0 0.06 26.6% 37.8% 48.39 128.0 37.8%
* Local BN: 256.0 GB/s, Remote BN: 256.0 GB/s
[v] Cross-cube best < worst: PASS (330.52ns < 677.12ns)
BW Saturation (Util% by data size):
Case 4KB 16KB 64KB 256KB 1MB
------------------------------------------------------------------
pe-local-hbm 88.9% 97.0% 99.2% 99.8% 100.0%
pe-same-half-hbm 79.9% 94.1% 98.5% 99.6% 99.9%
pe-cross-half-hbm 61.3% 86.4% 96.2% 99.0% 99.8%
pe-cross-cube-hbm-best 51.6% 81.0% 94.5% 98.6% 99.6%
pe-cross-cube-hbm-worst 15.1% 41.6% 74.0% 91.9% 97.8%
============================================================
docs/report/1H-codesign-paper/figures/sip_architecture.pdf
BIN
View File
Binary file not shown.
docs/report/1H-codesign-paper/figures/sip_architecture.svg
+94
View File
@@ -0,0 +1,94 @@
<svg xmlns="http://www.w3.org/2000/svg" width="508" height="495" viewBox="70 40 508 495">
<title>sip</title>
<rect width="648" height="648" fill="#ffffff"/>
<line x1="108.0" y1="144.0" x2="252.0" y2="144.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="180.0" y="140.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="108.0" y1="144.0" x2="108.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="108.0" y="200.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="252.0" y1="144.0" x2="396.0" y2="144.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="324.0" y="140.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="252.0" y1="144.0" x2="252.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="252.0" y="200.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="396.0" y1="144.0" x2="540.0" y2="144.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="468.0" y="140.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="396.0" y1="144.0" x2="396.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="396.0" y="200.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="540.0" y1="144.0" x2="540.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="540.0" y="200.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="108.0" y1="264.0" x2="252.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="180.0" y="260.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="108.0" y1="264.0" x2="108.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="108.0" y="320.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="252.0" y1="264.0" x2="396.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="324.0" y="260.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="252.0" y1="264.0" x2="252.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="252.0" y="320.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="396.0" y1="264.0" x2="540.0" y2="264.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="468.0" y="260.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="396.0" y1="264.0" x2="396.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="396.0" y="320.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="540.0" y1="264.0" x2="540.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="540.0" y="320.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="108.0" y1="384.0" x2="252.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="180.0" y="380.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="108.0" y1="384.0" x2="108.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="108.0" y="440.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="252.0" y1="384.0" x2="396.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="324.0" y="380.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="252.0" y1="384.0" x2="252.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="252.0" y="440.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="396.0" y1="384.0" x2="540.0" y2="384.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="468.0" y="380.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="396.0" y1="384.0" x2="396.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="396.0" y="440.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="540.0" y1="384.0" x2="540.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="540.0" y="440.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="108.0" y1="504.0" x2="252.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="180.0" y="500.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="252.0" y1="504.0" x2="396.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="324.0" y="500.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<line x1="396.0" y1="504.0" x2="540.0" y2="504.0" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="468.0" y="500.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">1.0mm 512GB/s</text>
<polyline points="324.0,56.0 108.0,56.0 108.0,144.0" fill="none" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="216.0" y="96.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">2.5mm 512GB/s</text>
<polyline points="324.0,56.0 252.0,56.0 252.0,144.0" fill="none" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="288.0" y="96.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">2.5mm 512GB/s</text>
<polyline points="324.0,56.0 396.0,56.0 396.0,144.0" fill="none" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="360.0" y="96.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">2.5mm 512GB/s</text>
<polyline points="324.0,56.0 540.0,56.0 540.0,144.0" fill="none" stroke="#000000" stroke-width="1" opacity="0.8"/>
<text x="432.0" y="96.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#000000">2.5mm 512GB/s</text>
<rect x="84.0" y="128.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="108.0" y="148.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (0,0)</text>
<rect x="228.0" y="128.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="252.0" y="148.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (1,0)</text>
<rect x="372.0" y="128.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="396.0" y="148.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (2,0)</text>
<rect x="516.0" y="128.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="540.0" y="148.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (3,0)</text>
<rect x="84.0" y="248.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="108.0" y="268.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (0,1)</text>
<rect x="228.0" y="248.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="252.0" y="268.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (1,1)</text>
<rect x="372.0" y="248.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="396.0" y="268.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (2,1)</text>
<rect x="516.0" y="248.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="540.0" y="268.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (3,1)</text>
<rect x="84.0" y="368.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="108.0" y="388.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (0,2)</text>
<rect x="228.0" y="368.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="252.0" y="388.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (1,2)</text>
<rect x="372.0" y="368.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="396.0" y="388.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (2,2)</text>
<rect x="516.0" y="368.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="540.0" y="388.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (3,2)</text>
<rect x="84.0" y="488.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="108.0" y="508.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (0,3)</text>
<rect x="228.0" y="488.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="252.0" y="508.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (1,3)</text>
<rect x="372.0" y="488.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="396.0" y="508.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (2,3)</text>
<rect x="516.0" y="488.0" width="48.0" height="32.0" rx="4" fill="#cbd5e1" stroke="#000000" stroke-width="1"/>
<text x="540.0" y="508.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#1e293b">CUBE (3,3)</text>
<rect x="308.0" y="50.0" width="32.0" height="12.0" rx="4" fill="#0ea5e9" stroke="#000000" stroke-width="1"/>
<text x="324.0" y="60.0" text-anchor="middle" font-family="monospace" font-size="10" fill="#ffffff">IO io0</text>
</svg>
Side by Side

After

Width:  |  Height:  |  Size: 10 KiB

docs/report/1H-codesign-paper/main.tex
+3
View File
@@ -13,8 +13,11 @@
\usepackage{hyperref}
\hypersetup{colorlinks=true,linkcolor=blue!50!black,citecolor=blue!50!black,urlcolor=blue!50!black}
\usepackage{caption}
\usepackage{subcaption}
\captionsetup{font=small,labelfont=bf}
\usepackage{microtype}
\usepackage{tikz}
\usetikzlibrary{arrows.meta,positioning,calc,fit}
\graphicspath{{figures/}}
docs/report/1H-codesign-paper/sections/01-introduction.tex
+6 -3
View File
@@ -25,9 +25,12 @@ serves as the common evaluation platform for all mechanisms and kernels
discussed in this study.
This report focuses on Grouped-Query Attention (GQA), one of the most
performance- and bandwidth-critical components of LLM inference. Modern
decoder-only models such as Llama~3 and Mistral have largely transitioned
from GPT-3-style multi-head attention (MHA) to GQA, in which multiple
performance- and bandwidth-critical components of LLM inference. GQA
dominates inference-time memory traffic and KV-cache capacity in modern
LLM serving, making it the primary bandwidth bottleneck on memory-centric
architectures such as AHBM. Modern decoder-only models such as Llama~3
and Mistral have largely transitioned from GPT-3-style multi-head
attention (MHA) to GQA, in which multiple
query heads share a single KV head to reduce KV cache capacity and
memory-bandwidth requirements. While GQA improves system efficiency at
the model level, mapping it efficiently onto AHBM introduces three
docs/report/1H-codesign-paper/sections/01-introduction_ko.tex
+73
View File
@@ -0,0 +1,73 @@
\section{서론}
\label{sec:intro}
AHBM은 연산 유닛을 HBM 스택 내부에 직접 통합한 메모리 중심 가속기
아키텍처이다. 각 처리 요소(PE)는 전용 HBM 슬라이스와 짝지어지며,
TCM과 SRAM을 거쳐 메모리와 MAC 어레이 사이의 데이터를 단계적으로
이동시킨다. 이러한 메모리 중심 구조에서 커널 성능은 연산 처리량만이
아니라, 메모리 계층과 PE 사이에 데이터를 얼마나 효율적으로 배치하고,
이동시키며, 공유하는가에 의해 결정된다. 따라서 AHBM에서의 AI 커널
최적화는 커널 알고리즘과 아키텍처 메커니즘을 함께 발전시키는
하드웨어--소프트웨어 코디자인을 필요로 한다.
상세한 성능 분석과 신속한 설계 탐색을 위해, 우리는
\textbf{KernBench}---AHBM을 위한 소스 수준 이산 사건 시뮬레이션
플랫폼---을 개발하였다. KernBench는 메모리 시스템 지연, PE 실행
모델, PE 간 통신, 호스트 측 오케스트레이션을 포함한 AHBM 실행 모델을
구현하면서, 커널과 호스트 소프트웨어를 소스 코드로부터 직접 실행한다.
이를 통해 실행 거동을 세밀하게 관찰할 수 있을 뿐만 아니라, 컴파일러나
런타임과 같은 상위 소프트웨어 스택과 독립적으로 하드웨어--소프트웨어
코디자인 선택을 체계적으로 평가할 수 있다. 본 보고서에 제시된 모든
결과는 KernBench를 사용하여 얻은 것이며, KernBench는 본 연구에서
다루는 모든 메커니즘과 커널의 공통 평가 플랫폼 역할을 한다.
본 보고서는 LLM 추론에서 가장 성능 및 대역폭 임계 구성요소 중 하나인
Grouped-Query Attention(GQA)에 초점을 맞춘다. Llama~3, Mistral과
같은 최신 디코더 전용 모델은 GPT-3 형태의 다중-헤드 어텐션(MHA)으로
부터 GQA---여러 질의 헤드가 하나의 KV 헤드를 공유함으로써 KV 캐시
용량과 메모리 대역폭 요구량을 줄이는 방식---로 대거 전환되었다. GQA는
모델 수준에서 시스템 효율을 개선하지만, 이를 AHBM에 효율적으로
매핑하기 위해서는 세 가지 아키텍처 요구사항이 발생한다: PE 간 통신을
최소화하기 위한 KV 캐시와 가중치의 최적 배치, 불가피한 PE 간 트래픽에
대한 저오버헤드 지원, 그리고 각 PE 내에서 메모리 접근과 연산의
효율적인 파이프라이닝이다.
이러한 요구사항을 충족하기 위해, 본 보고서는 세 가지 하드웨어--
소프트웨어 코디자인 메커니즘을 제안한다. 첫째, GQA-aware 데이터
배치는 TCM/SRAM/HBM 계층에 KV 캐시와 가중치를 분산 배치하여 통신
오버헤드를 줄이고 데이터 지역성을 향상시킨다. 둘째, PE\_IPCQ는 리덕션
등 통신 집약적 연산을 위한 효율적인 온디바이스 집합 통신 프리미티브를
제공한다. 셋째, composite-command GEMM 파이프라인은 PE\_SCHEDULER의
제어 하에 각 PE 내부에서 메모리 이동과 연산을 긴밀하게 파이프라이닝
하여, 명령 오버헤드를 줄이면서도 MAC 어레이의 가동률을 효율적으로
유지한다.
연산 측 인에이블러(composite-command GEMM 파이프라인,
\S\ref{sec:gemm})와 통신 측 인에이블러(PE\_IPCQ,
\S\ref{sec:allreduce})는 우선 독립적으로 개발 및 평가된다. 이후 융합
GQA 커널(\S\ref{sec:gqa})은 이들을 GQA-aware 데이터 배치와 결합하여
AHBM 상에서 종단간 어텐션 구현을 시연한다. 대응 관계는 직접적이다:
어텐션의 $QK^{\top}$와 $PV$ 곱은 정확히 composite-command 파이프라인이
이득을 주는 GEMM이며, KV 리덕션은 정확히 PE\_IPCQ가 이득을 주는 집합
연산이다. 이러한 메커니즘들이 결합되어, 융합 GQA 커널은 AHBM의 HBM
대역폭을 효율적으로 활용할 수 있게 된다.
본 연구에서 GQA가 주된 동기 부여 워크로드 역할을 하지만, 도출된
메커니즘들은 훨씬 광범위한 AI 커널 분류군에서 재사용 가능한 구성
요소로 의도되었다. PE\_IPCQ는 분산 및 통신 집약적 워크로드 전반에
걸쳐 집합 통신을 지원할 수 있으며, composite-command 실행은 GEMM
기반 커널, 피드포워드 네트워크(FFN), 정규화, 그 외 융합 연산자
파이프라인에 적용 가능하다. 계층적 데이터 배치 프레임워크와 함께 이
메커니즘들은 AHBM에서의 향후 AI 커널과 통신 라이브러리의 기반을
형성한다. 2026년 하반기에는 이 기반을 FFN 및 MoE 주도 워크로드와
완전한 LLM 실행의 종단간 최적화로 확장할 예정이다.
본 보고서의 나머지 부분은 다음과 같이 구성된다.
\S\ref{sec:platform}에서는 KernBench 플랫폼과 본 연구 전반에 걸쳐
사용된 AHBM 구성을 기술한다. \S\ref{sec:gemm},
\S\ref{sec:allreduce}, \S\ref{sec:gqa}에서는 각각 composite-command
GEMM 파이프라인, PE\_IPCQ 집합 통신, 융합 GQA 커널을 다룬다.
\S\ref{sec:discussion}에서는 이러한 결과의 광범위한 아키텍처적 함의를
논의한다. 마지막으로 \S\ref{sec:conclusion}과 \S\ref{sec:future}에서는
주요 결과를 요약하고 FFN, MoE, 전체 모델 최적화에 관한 향후 과제를
제시한다.
docs/report/1H-codesign-paper/sections/02-platform.tex
+337 -100
View File
@@ -1,15 +1,75 @@
\begin{figure*}[t]
\centering
\begin{subfigure}[b]{0.495\textwidth}
\centering
\includegraphics[width=\linewidth,height=0.7\linewidth,keepaspectratio]{sip_architecture.pdf}
\caption{SIP Architecture}
\label{fig:sip-arch}
\end{subfigure}\hfill
\begin{subfigure}[b]{0.42\textwidth}
\centering
\includegraphics[width=\linewidth,height=0.84\linewidth]{cube_architecture.pdf}
\caption{CUBE Architecture}
\label{fig:cube-arch}
\end{subfigure}
\vspace{0.6em}
\begin{subfigure}[t]{\textwidth}
\centering
\includegraphics[width=\linewidth]{pe_architecture.png}
\caption{PE level (zoom-in of one PE in (b)): \textsf{PE\_CPU}
dispatches commands to \textsf{PE\_SCHED}, which routes tile-token
streams through \textsf{PE\_DMA}, \textsf{PE\_FETCH\_STORE}, and
\textsf{GEMM}/\textsf{MATH} engines; \textsf{PE\_IPCQ} is the
on-device collective control plane.}
\label{fig:pe-arch}
\end{subfigure}
\caption{Modeled hardware graph at the SIP, CUBE, and PE levels.
This figure is an \emph{illustrative topology} chosen for
readability; the \emph{experimental configuration} used throughout
this report (port counts, slice counts, and link parameters) is
specified in \S\ref{sec:hw} / Table~\ref{tab:hw}, and numbers in the
two may differ. KernBench is not tied to this particular
arrangement: each box is a modeled component node, each line a
directed link with bandwidth and propagation attributes, and any
topology that respects those attributes is supported.}
\label{fig:hw-arch}
\end{figure*}
\begin{figure*}[t]
\centering
\includegraphics[width=\textwidth]{latency_model.png}
\caption{Conceptual schematic of the latency model. Two source nodes
(Requester A/B) inject flits through a chain of routers into a
destination node; on the shared edge between routers, flits from the
two transactions are interleaved flit-by-flit by the wire's FIFO
arrival order. End-to-end latency is the sum of four contributions:
\textbf{per-node overhead} (the switch's fixed processing cost,
shown in light yellow---the same colour as the destination's
processing-logic block), \textbf{per-edge transmission}
($\textit{flit\_size}/\textit{BW}$ on each wire),
\textbf{drain} (per-flit service occupancy at the destination's
channel), and \textbf{queuing delay} (waiting in a FIFO when a
shared resource is busy). The places where queuing actually
accumulates are highlighted in green: the router output queue and
the destination's input queue. This is the model, not a
measurement; specific bandwidths and overheads are listed in
\S\ref{sec:hw} (Table~\ref{tab:hw}).}
\label{fig:latency-model}
\end{figure*}
\section{The KernBench Platform}
\label{sec:platform}
All results in this report are produced on \emph{KernBench}, a
system-level discrete-event simulator for LLM kernels running on AHBM.
This section explains why the platform exists, how it executes a
kernel, how its latency model works---specifically how the hardware is
viewed as a graph, how that graph is driven by a discrete-event engine,
and how congestion is captured---and the concrete hardware configuration
used for every experiment that follows.
This section explains why the platform exists, the device and
execution model it presents to a kernel writer, how its latency model
turns that execution into a number, and the concrete hardware
configuration used for every experiment that follows.
\subsection{Why KernBench: source-level kernels without a software stack}
\subsection{Why KernBench}
\label{sec:why}
In a production end-to-end (E2E) stack, kernel performance is entangled
@@ -33,51 +93,100 @@ without the confound of compiler maturity or framework overhead. The
cost is that KernBench numbers are \emph{not} E2E latencies; they are
the achievable-kernel latencies an ideal software stack would expose.
\subsection{Execution model}
\subsection{Device and execution model}
\label{sec:exec}
KernBench models AHBM as a hierarchy of SIPs, CUBEs, and processing
elements (PEs). At the system level, multiple SIPs are connected
through inter-package links, while each SIP contains a collection of
CUBEs joined by an on-package interconnect
(Fig.~\ref{fig:hw-arch}\subref{fig:sip-arch}).
Each CUBE contains eight PEs, shared SRAM, HBM controllers, an
\textsf{M\_CPU} control processor, and an intra-CUBE router mesh
(Fig.~\ref{fig:hw-arch}\subref{fig:cube-arch}). Together these
components form the execution substrate for all kernels evaluated in
this report. This organization reflects the memory-centric nature of
AHBM: each PE is paired with a dedicated slice of HBM bandwidth and
local on-PE storage (TCM), so compute lives next to the data it
consumes rather than fetching it through a far-away memory
controller. The platform's performance question is therefore not
``how many FLOPs can the chip do'' but ``how well can a kernel keep
each compute engine fed from its locally-attached memory while
moving the unavoidable traffic between PEs efficiently.''
Within a PE, commands are dispatched by \textsf{PE\_CPU} to
\textsf{PE\_SCHED}, which routes work to specialized execution
engines---DMA, FETCH/STORE, GEMM, vector-math, and IPCQ
(Fig.~\ref{fig:hw-arch}\subref{fig:pe-arch}). Commands come in two
flavours. \emph{Atomic} commands target a single engine---a plain
DMA read, a GEMM tile, a vector-math op, an IPCQ send/receive---and
are the natural unit for short or special-purpose work; the
\textsf{PE\_CPU} itself also runs control-plane work directly.
\emph{Composite} commands, by contrast, carry an ordered pipeline of
operations across multiple engines (\textsf{DMA\_READ} $\rightarrow$
\textsf{FETCH} $\rightarrow$ \textsf{GEMM}/\textsf{MATH} $\rightarrow$
\textsf{STORE} $\rightarrow$ \textsf{DMA\_WRITE}) that the
\textsf{PE\_SCHED} tiles and streams without per-tile redispatch. The
composite form is the substrate for the GEMM optimization
(\S\ref{sec:gemm}) and, combined with on-PE collectives, for fused
attention (\S\ref{sec:gqa}).
KernBench is layered along the flow of a request:
\begin{itemize}
\item The \textbf{runtime API} is host-facing and
topology-agnostic---it deploys tensors and launches kernels but knows
nothing about routing or interconnect.
\item The \textbf{simulation engine} schedules discrete events, routes
every request through the modeled graph, and tracks completion via
per-request correlation IDs.
\item The \textbf{simulation engine} schedules discrete events and
routes every request through the modeled graph.
\item The \textbf{components} are device-side nodes that model
hardware behavior: the per-PE blocks (scheduler, DMA, GEMM and
hardware behaviour: the per-PE blocks (scheduler, DMA, GEMM and
vector-math engines, TCM, IPCQ), the NoC routers, the HBM
controllers, and the inter-chiplet links.
\end{itemize}
Within a PE, work is expressed as \emph{composite commands}: a single
command carries an ordered pipeline of operations
(\textsf{DMA\_READ} $\rightarrow$ \textsf{FETCH} $\rightarrow$
\textsf{GEMM}/\textsf{MATH} $\rightarrow$ \textsf{STORE} $\rightarrow$
\textsf{DMA\_WRITE}) that the PE scheduler tiles and streams. This
composite mechanism is the substrate for the GEMM optimization
(\S\ref{sec:gemm}) and, combined with on-PE collectives, for fused
attention (\S\ref{sec:gqa}). Data and timing are handled in two passes,
so that a kernel's numeric results and its latency are computed
consistently but independently.
Data and timing are handled in two passes, so that a kernel's
numeric results and its latency are computed consistently but
independently. \textbf{Pass~1 (timing)} runs the kernel under
the discrete-event engine: memory ops (\textsf{tl.load},
\textsf{tl.store}) execute against a host-side \textsf{MemoryStore}
and return real tensor data, while compute ops (GEMM, vector-math)
only emit records into an op-log carrying their operands, shapes,
dtypes, and scheduled time. \textbf{Pass~2 (data)} replays that
op-log offline in numpy, producing the actual numeric outputs and
comparing them against a reference computation under per-dtype
tolerances (e.g.\ \texttt{rtol}/\texttt{atol} $=10^{-3}$ for f16).
Pass~2 is optional---runs that need only latency skip it---but when
enabled it guarantees that every reported timing number corresponds
to a kernel whose numeric output has been verified end-to-end.
\subsection{Latency model: a graph traversed by events}
\subsection{Latency model: graph traversal and contention}
\label{sec:latency}
\paragraph{The hardware as a graph.} KernBench views the modeled
hardware as a directed graph. \emph{Nodes} are the components listed
above; \emph{edges} are the interconnect links between them, each
carrying bandwidth (\si{\giga\byte\per\second}) and propagation
(\si{\nano\second}) attributes. The topology is compiled once at
The modeled hardware hierarchy described above is represented
internally as a directed graph. Nodes correspond to hardware
components---PEs, routers, memory controllers, SRAM blocks, IO
chiplets---while edges represent communication links with associated
bandwidth (\si{\giga\byte\per\second}) and propagation
(\si{\nano\second}) attributes. Every operation in KernBench---DMA
transfers, remote memory accesses, collective communication, command
dispatch---is modelled as a traversal through this graph. End-to-end
latency is decomposed into four contributions accumulated along the
traversal path (Fig.~\ref{fig:latency-model}): \textbf{per-node
overhead} at each component, \textbf{per-edge transmission} on each
wire, \textbf{drain} (per-flit service occupancy) at the destination,
and \textbf{queuing delay} at the shared FIFOs that the wire and the
destination share between concurrent transactions.
\paragraph{The hardware as a graph.} The topology is compiled once at
configuration time into this graph and is never mutated during a run.
Every routed request---a DMA, a remote read, an IPCQ message, a
kernel-launch command---is a \emph{traversal} of this graph from a
source node to a destination service, hopping through routers and
links along the way. There are no hidden shortcuts, implicit bypasses,
or magic paths: if a request reaches its destination, the path it took
is explicit in the graph, and the latency it incurred is the sum of
the per-node and per-edge costs paid along that path.
There are no hidden shortcuts, implicit bypasses, or magic paths: if a
request reaches its destination, the path it took is explicit in the
graph, and the latency it incurred is the sum of the per-node and
per-edge costs paid along that path. The same graph representation
applies recursively at every hierarchy level---system, SIP, CUBE, and
PE (Fig.~\ref{fig:hw-arch}).
\paragraph{From graph to discrete-event simulation.} The graph is
driven by a discrete-event engine. Two kinds of events advance
@@ -85,14 +194,14 @@ simulation time: \emph{node events} (component switching overhead,
service completions such as an HBM channel commit or a GEMM tile
finish) and \emph{edge events} (the flit-by-flit serialization of a
payload across a bandwidth-limited link). The engine maintains a
priority queue of pending events ordered by their scheduled time,
fires them one at a time, and treats ties under a deterministic
ordering policy so that the same kernel on the same topology always
yields the same trace. Per-request correlation IDs are stamped at
injection and carried through every hop, so the trace is recoverable
from injection to completion. Every nanosecond in a reported latency
traces back to exactly one of these events on exactly one node or
edge.
priority queue of pending events ordered by their scheduled time and
fires them one at a time, with ties broken under a deterministic
policy so that the same kernel on the same topology always yields the
same trace. Per-request correlation IDs are stamped at injection and
carried through every hop, so the path from injection to completion
is fully traceable. Every modeled latency contribution
corresponds to exactly one of these events on exactly one node or
edge---there is no slack in the budget.
\paragraph{Latency contributions.} Three kinds of latency accumulate
along a traversal: (i) \emph{per-node fixed overhead}---each component
@@ -107,17 +216,83 @@ collective engines hold the request for their service time before
releasing it downstream. Each of these is attached to a specific node
or edge in the graph; together they make up the entire latency budget.
\paragraph{Congestion: where bottlenecks emerge.} The simulator's
Beyond data movement and execution latency, KernBench
also models the control-plane cost required to issue work
to accelerator engines. Since every DMA, GEMM, vector-
math, and IPCQ operation is initiated through the
\textsf{PE\_CPU} $\rightarrow$ \textsf{PE\_SCHED} path,
command dispatch latency contributes to the end-to-end
execution time of all kernels.
\paragraph{Command dispatch overhead model.} The cost incurred by
the \textsf{PE\_CPU} when it dispatches a command to one of the
accelerator engines is modelled structurally rather than with a
per-operation calibration table. The \textsf{PE\_CPU} charges, per
command,
\[
d_{\text{cmd}} = \textsf{FIXED} + b_{\text{logical}} \cdot R,
\]
This linear-in-size form reflects the serialization cost of writing a
command descriptor into the scheduler queue: a fixed per-command
bookkeeping cost plus a byte-wise descriptor-transfer cost.
Concretely, $b_{\text{logical}}$ is the command's hardware-logical byte
size, \textsf{FIXED} captures the fixed per-command cost (queue-tail
update, completion registration) and $R$ captures the per-byte cost of
serializing the command descriptor into the scheduler queue. This
\textsf{FIXED} is distinct from Table~\ref{tab:hw}'s
\SI{2}{\nano\second} / \SI{1}{\nano\second} \textsf{PE\_CPU} /
\textsf{PE\_SCHED} fixed costs: the latter are per-component
traversal overheads each command pays when it transits those nodes,
whereas the \textsf{FIXED} term here is the command-descriptor
issue cost. \textsf{FIXED} depends on the command class: a single-op
command --- one engine operation, i.e.\ a single DMA descriptor, GEMM,
or elementwise issue --- carries a lighter 8-cycle \textsf{FIXED},
while the composite command, which the scheduler expands into a
multi-stage tile-feeder plan, carries 40 cycles
(\S\ref{sec:gemm-vs-async}). With the composite \textsf{FIXED} $= 40$
cycles and $R = 0.0625$ cycles/byte (i.e.\ \SI{16}{\byte\per\cycle},
at \SI{1}{\giga\hertz}) a typical composite lands at roughly
\SI{43}{\nano\second}, and a hard cap on a composite's descriptor size
prevents the model from rewarding arbitrarily large fused commands
beyond what real descriptor queues accept. In the configurations measured here, command issue is
not the bottleneck---data movement is---so this term stays small
relative to DMA and collective time.
\begin{table*}[t]
\centering
\caption{PE\,$\to$\,HBM DMA latency probe at varying hop distances
(\SI{32}{\kibi\byte} transfer; output captured from \texttt{kernbench
probe}). \emph{Util\%} is the achieved bandwidth as a fraction of the
path's bottleneck-edge bandwidth.}
\label{tab:probe-pe-dma}
\small
\begin{tabular}{@{}lrrr@{}}
\toprule
Case & Latency~(\si{\nano\second}) & Util\% (\,32\,KiB) & Util\% (\,1\,MiB) \\
\midrule
PE\,$\to$\,local HBM & 141.0 & 90.8 & 100.0 \\
PE\,$\to$\,same-half HBM & 147.9 & 86.6 & ~99.9 \\
PE\,$\to$\,cross-half HBM & 161.2 & 79.4 & ~99.8 \\
PE\,$\to$\,cross-CUBE (best) & 330.5 & 77.5 & ~99.6 \\
PE\,$\to$\,cross-CUBE (worst) & 677.1 & 37.8 & ~97.8 \\
\bottomrule
\end{tabular}
\end{table*}
\subsubsection{Congestion and contention modeling}
\label{sec:congestion}
The simulator's
sharpness comes from how it models contention for those nodes and
edges. \emph{Every directed edge has a FIFO}: an arriving flit takes
its bandwidth-limited transfer time on top of whatever earlier flits
are still being served, so a busy link queues later traffic behind
earlier traffic rather than transferring everything at peak BW.
\emph{HBM is modeled with per-pseudo-channel parallelism}: a stateless
array of channel-availability timestamps with address-based channel
selection captures the bank-level concurrency that real HBM exposes,
so 64 channels per CUBE deliver real parallelism on uniform addresses
but a hot channel surfaces as the bottleneck on skewed ones.
\emph{HBM is modelled with per-pseudo-channel parallelism}: a
stateless array of channel-availability timestamps with address-based
channel selection captures the bank-level concurrency that real HBM
exposes, so 64 channels per CUBE deliver real parallelism on uniform
addresses but a hot channel surfaces as the bottleneck on skewed ones.
\emph{Every component has a serial worker}: a router carrying two
heavy streams interleaves them at flit granularity in arrival order
rather than fanning out for free, so two concurrent collectives sharing
@@ -127,67 +302,127 @@ it reveals where the real bottlenecks form and which hardware levers
actually relieve them---which is exactly the question the codesign
work in this report turns on.
\paragraph{Control-plane (issue) cost model.} The cost of \emph{issuing}
a command is modeled structurally rather than with a per-operation
calibration table. The PE control processor charges, per command,
\[
d_{\text{cmd}} = \textsf{FIXED} + b_{\text{logical}} \cdot R,
\]
where $b_{\text{logical}}$ is the command's hardware-logical byte size,
\textsf{FIXED} captures the fixed per-command cost (queue-tail update,
completion registration) and $R$ captures the per-byte cost of
serializing the command descriptor into the scheduler queue. The
default anchoring (\textsf{FIXED} $= 40$ cycles, $R = 0.0625$
cycles/byte, i.e.\ \SI{16}{\byte\per\cycle}, at \SI{1}{\giga\hertz})
places a typical composite at roughly \SI{43}{\nano\second}, and a
hard cap on a composite's descriptor size prevents the model from
rewarding arbitrarily large fused commands beyond what real descriptor
queues accept. In the configurations measured here, command issue is
not the bottleneck---data movement is---so this term stays small
relative to DMA and collective time.
\paragraph{Accuracy.} The model is precise about the effects that
\subsection{Accuracy}
\label{sec:accuracy}
The model is precise about the effects that
dominate kernel latency on this class of hardware: per-edge bandwidth
occupancy and flit-level serialization, HBM pseudo-channel parallelism,
and per-component switching overhead. Two independent cross-checks
drawn from the experiments in this report confirm that this precision
translates into physically reasonable kernel latencies. First, in the
GEMM study (\S\ref{sec:gemm}), simulator-measured MAC efficiency
tracks an analytic ideal-pipeline model within roughly
\SIrange{10}{20}{\percent} across a wide range of tile counts; the
residual gap is attributable to pipeline-fill and DMA effects the
analytic model omits. Second, in the all-reduce study
(\S\ref{sec:allreduce}, Fig.~\ref{fig:allreduce-cmp}), simulator
latency for a 2D-torus over six devices follows the expected
startup-plus-per-packet shape across the entire payload sweep---tight
at small payloads where startup dominates, and within a single-digit
multiplicative factor at the largest payloads, where the residual gap
is explained by per-router switching the analytic shape elides. A
single-device point from an external full-system simulator (FSIM) at
the largest payload sits an order of magnitude above the KernBench
multi-device torus, illustrating the well-known gap between an
achievable-kernel number and a full end-to-end-stack number rather
than a model error. The known simplifications---FIFO router
arbitration (instead of round-robin), HBM scheduler without
write-buffer reordering, no bank conflict, no refresh or thermal
effects, and no upstream backpressure---are the price of a
deterministic, inspectable model. They bound the absolute accuracy but
do not distort the \emph{relative} comparisons (tiling A vs.\ B,
topology X vs.\ Y, with vs.\ without composite command, mesh vs.\
torus) that this report is built on.
translates into physically reasonable kernel latencies.
First, in the GEMM study (\S\ref{sec:gemm}), simulator-measured MAC
efficiency tracks an analytic ideal-pipeline model within
\SI{1.4}{ppt} across every swept shape---from a single-tile
$M{=}K{=}N{=}32$ at $\sim\SI{7.7}{\percent}$ up to the deep-$K$
$K{=}3072$ case at $\sim\SI{90}{\percent}$. The residual gap is
fill/tail overhead the closed-form pipeline model omits.
Second, in the all-reduce study (\S\ref{sec:allreduce},
Fig.~\ref{fig:allreduce-cmp}), simulator latency for a 2D-torus over
six devices follows the expected startup-plus-per-packet shape across
the entire payload sweep---tight at small payloads where startup
dominates, and within a single-digit multiplicative factor at the
largest payloads, where the residual gap is explained by per-router
switching the analytic shape elides.
Third, the simulator's per-traversal behaviour can be probed
directly. Running \texttt{kernbench probe} on the modelled topology
issues a sequence of PE-to-HBM DMA reads at progressively greater
hop distances and reports the per-component overhead, per-edge
serialization, and per-PC drain that the model charges
(Table~\ref{tab:probe-pe-dma}). Three properties stand out. First,
the reported latency increases \emph{monotonically} with hop
count---from \SI{141}{\nano\second} at the local HBM slice to
\SI{677}{\nano\second} at the worst-case remote-CUBE slice---with the
increment per added hop matching the per-router overhead and the
UCIe-link cost in Table~\ref{tab:hw}. Second, the bandwidth-saturation
curves track the wire's serialization model: at \SI{1}{\mebi\byte}
transfers the local PE DMA reaches \SI{100}{\percent} of the per-edge
bandwidth limit, while the longest cross-CUBE path saturates at
\SI{97.8}{\percent}---exactly the gap that the per-edge propagation
and per-router overhead model would predict. Third, the simulator
passes the internal-consistency invariants the probe builds in
(monotonic hop progression; D2H reads at least as long as the
equivalent H2D writes because of the reverse-path acknowledgement;
cross-CUBE best less than cross-CUBE worst). These are properties
that hold \emph{only} when the underlying graph traversal, FIFO
contention, and propagation models are internally
self-consistent---a model error in any of them would surface as a
non-monotonic or under-utilising curve here long before it polluted a
kernel-level measurement.
The known simplifications---FIFO router arbitration (instead of
round-robin), HBM scheduler without write-buffer reordering, no bank
conflict, no refresh or thermal effects, and no upstream
backpressure---are the price of a deterministic, inspectable model.
These simplifications bound the absolute accuracy, but the agreement
with both analytic models and the probe's internal-consistency
checks above indicates that KernBench is sufficiently accurate for
evaluating the \emph{relative} hardware--software design trade-offs
(tiling A vs.\ B, topology X vs.\ Y, with vs.\ without composite
command, mesh vs.\ torus) that are the primary objective of this
work.
\subsection{Modeled hardware configuration}
\label{sec:hw}
Table~\ref{tab:hw} summarizes the hardware configuration used for
every experiment in this report. It is read directly from the
simulator's topology description; per-experiment workload parameters
(matrix shapes, collective sizes, sequence lengths) are stated in
their respective sections rather than here.
Table~\ref{tab:hw} summarizes the hardware configuration
used throughout this report. The intent of this
configuration is not to model a specific product, but to
represent a realistic memory-centric accelerator and to
provide a consistent baseline for evaluating hardware--
software co-design mechanisms.
The compute capability within each CUBE is provisioned
such that inference workloads can effectively saturate
the available HBM bandwidth. The aggregate compute
throughput is therefore balanced against memory-system
bandwidth rather than being intentionally over- or
under-provisioned. Concretely, 8 PEs $\times$
\SI{8}{\tera\flop\per\second} give roughly
\SI{64}{\tera\flop\per\second} of f16 compute per CUBE
against \SI{2048}{\giga\byte\per\second} of HBM bandwidth,
which places the balanced point at about
$\sim$31~FLOP/byte; inference decode kernels typically
operate well below that, so HBM bandwidth---not compute---is
the natural ceiling on this configuration. For reference, the
GQA decode kernels evaluated in \S\ref{sec:gqa} operate at
arithmetic intensity well below this balance point, so their
ceiling is set by HBM and inter-PE traffic rather than by GEMM
throughput. HBM bandwidth
is distributed evenly across the PEs within a CUBE, with
each PE responsible for servicing approximately one-eighth
of the CUBE's memory bandwidth through its dedicated HBM
channels.
The on-chip interconnect is configured using bandwidth
and latency parameters representative of commercially
available mesh-network IPs. The goal is not to study a
particular NoC implementation, but rather to evaluate
kernel behavior under a realistic baseline communication
fabric.
For die-to-die communication, UCIe-A bandwidth
characteristics are used as the reference point for
inter-CUBE links. Inter-SIP communication is modeled
using PCIe-class links. Together, these assumptions
provide a representative communication hierarchy for
evaluating collective communication and distributed
kernel execution.
Unless otherwise stated, all experiments use this
configuration. Workload-specific parameters such as
matrix dimensions, sequence lengths, and collective
payload sizes are introduced in their respective sections.
\begin{table}[t]
\centering
\caption{Modeled hardware configuration (shared by all experiments).}
\caption{Modeled hardware configuration}
\label{tab:hw}
\small
\begin{tabular}{@{}ll@{}}
@@ -205,11 +440,12 @@ GEMM engine peak & \SI{8}{\tera\flop\per\second} (f16) \\
TCM (on-PE) & \SI{16}{\mega\byte}, \SI{512}{\giga\byte\per\second} R/W \\
\quad kernel scratch & \SI{1}{\mega\byte} \\
DMA engines & 1 read + 1 write \\
CPU / scheduler overhead & \SI{2}{\nano\second} / \SI{1}{\nano\second} \\
\textsf{PE\_CPU} fixed cost & \SI{2}{\nano\second} \\
\textsf{PE\_SCHED} fixed cost & \SI{1}{\nano\second} \\
\midrule
\multicolumn{2}{@{}l}{\emph{Memory (per CUBE)}} \\
HBM capacity & \SI{48}{\giga\byte} (8 slices) \\
HBM aggregate BW & \SI{1024}{\giga\byte\per\second} \\
HBM aggregate BW & \SI{2048}{\giga\byte\per\second} \\
HBM pseudo-channels & 64 (8 per PE), \SI{32}{\giga\byte\per\second} each \\
SRAM (shared) & \SI{32}{\mega\byte}, \SI{128}{\giga\byte\per\second} link \\
HBM burst & \SI{256}{\byte} \\
@@ -220,7 +456,8 @@ Inter-CUBE (UCIe PHY) & \SI{512}{\giga\byte\per\second}, \SI{8}{\nano\second}, X
Inter-SIP (PCIe) & \SI{768}{\giga\byte\per\second} per endpoint \\
\midrule
\multicolumn{2}{@{}l}{\emph{Command-issue cost model (defaults)}} \\
FIXED per command & 40 cycles \\
FIXED per single-op command & 8 cycles \\
FIXED per composite command & 40 cycles \\
per-byte rate $R$ & 0.0625 cycles/byte (\SI{16}{\byte\per\cycle}) \\
composite size cap & \SI{1024}{\byte} \\
\bottomrule
docs/report/1H-codesign-paper/sections/03-gemm.tex
+354 -63
View File
@@ -1,8 +1,6 @@
\section{GEMM Acceleration via the Composite Command}
\label{sec:gemm}
\subsection{Why it is needed}
GEMM is the compute core of every transformer block---the QKV
projections, the attention score and context products, and the
feed-forward matrices are all matrix multiplications. On a tiled
@@ -21,20 +19,17 @@ in command overhead?
\subsection{Design}
The answer is the \emph{composite command}. A single command carries the
ordered tile pipeline
\[
\textsf{DMA\_READ}\rightarrow\textsf{FETCH}\rightarrow\textsf{GEMM}
\rightarrow\textsf{STORE}\rightarrow\textsf{DMA\_WRITE},
\]
The answer is the \emph{composite command}. A single command carries
the ordered five-stage tile pipeline (\textsf{DMA\_READ},
\textsf{FETCH}, \textsf{GEMM}, \textsf{STORE}, \textsf{DMA\_WRITE}),
and the PE scheduler splits the payload into hardware tiles
(here $32\times64\times32$), emitting one tile token per tile. Subsequent
stages are reached by \emph{token self-routing} between the on-PE engines,
so a tile flows DMA\,$\rightarrow$\,fetch\,$\rightarrow$\,GEMM\,$%
\rightarrow$\,store without returning to the scheduler between stages.
Because the whole pipeline is described by one command, the issue cost is
paid once per GEMM rather than once per tile-stage, and the scheduler is
free to keep every stage busy on different tiles simultaneously---tile
(here $32\times64\times32$), emitting one tile token per tile.
Subsequent stages are reached by \emph{token self-routing} between
the on-PE engines, so a tile flows through the chain without
returning to the scheduler between stages. Because the whole
pipeline is described by one command, the issue cost is paid once
per GEMM rather than once per tile-stage, and the scheduler is free
to keep every stage busy on different tiles simultaneously---tile
$i$'s GEMM overlaps tile $i{+}1$'s DMA read. A multi-operation composite
additionally lets an epilogue (for example a vector-math step) ride the
same tile loop, firing per $K$-tile, per output tile, or once per kernel
@@ -45,66 +40,362 @@ kernel to fuse its softmax work into the GEMM pipeline
\subsection{Results}
We sweep eight GEMM shapes spanning square, tall, wide, and deep-$K$
geometries, under three operand-staging variants
(\textsf{ref\_ref}, both operands streamed from HBM; \textsf{load\_ref},
one operand resident in TCM; \textsf{load\_load}, both resident).
Figure~\ref{fig:gemm-util} reports MAC utilization and efficiency, and
Figure~\ref{fig:gemm-stages} breaks the kernel into per-stage engine
busy time.
geometries under two operand-staging variants that bracket the
realistic LLM cases. In \textsf{load\_ref}, the activation $A$ is
pre-staged on chip and only the weight $W$ streams from HBM during
the composite (the ``activation in TCM, weights from HBM'' case
typical of decoding with a small batch). In \textsf{ref\_ref}, both
operands stream from HBM during the composite (the ``cold both
sides'' case, where the activation does not fit on-chip or is itself
the output of a producer composite that wrote back to HBM). Both
variants run the same composite command; only the up-front staging
of $A$ differs.
We evaluate the composite GEMM along two axes that together describe
how well it lands on the hardware: \emph{HBM bandwidth utilization}
(does the workload saturate the per-PE HBM bandwidth, i.e.\ is it
memory-bound on this configuration?) and \emph{achieved GEMM
throughput} (what fraction of the \SI{8}{\tera\flop\per\second}
per-PE peak does the kernel actually deliver). Both metrics use the
composite window as the denominator, so the up-front \textsf{tl.load}
of $A$ in \textsf{load\_ref} is excluded — only HBM traffic and
compute that happen \emph{inside} the composite count.
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gemm_hbm_bw_util.png}
\caption{Per-PE HBM bandwidth utilization during the composite window,
for the two staging variants. The ceiling is the per-PE HBM bandwidth
(\SI{256}{\giga\byte\per\second}; dashed line at \SI{100}{\percent}).
\textsf{ref\_ref} streams both $A$ and $W$ and so always sits at a
higher BW-utilization than \textsf{load\_ref} on the same shape;
both variants saturate ($>\SI{85}{\percent}$) once the kernel
has enough total bytes to keep the link busy (deep-$K$ $K{=}3072$,
or low-reuse output-dominated $M{=}128,K{=}8,N{=}128$). Small or
single-tile shapes do not have enough total traffic to saturate the
link and are bottlenecked by pipeline fill rather than HBM.}
\label{fig:gemm-bw}
\end{figure}
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gemm_per_pe_tflops.png}
\caption{Per-PE GEMM throughput delivered during the composite window
(reference line at the \SI{8}{\tera\flop\per\second} per-PE engine
peak). At the deep-$K$ corner $M{=}32,K{=}3072,N{=}32$ the
activation-pre-staged kernel reaches
\SI{7.18}{\tera\flop\per\second} ($\sim\SI{90}{\percent}$ of peak);
the both-from-HBM variant drops to
\SI{3.83}{\tera\flop\per\second} ($\sim\SI{48}{\percent}$) at the
same shape because the second operand doubles HBM pressure and the
kernel hits the BW ceiling shown in Fig.~\ref{fig:gemm-bw}.
Under-tile shapes ($M{=}K{=}N{=}32$, $M{=}8$, $K{=}8$) are
hard-capped by tile-fill at \SI{50}{\percent}, \SI{25}{\percent},
\SI{12.5}{\percent} of peak regardless of staging.}
\label{fig:gemm-tflops}
\end{figure}
The composite reaches the hardware roofline in two distinct regimes.
\emph{Memory-bound}: $K{=}3072$ deep-$K$ saturates HBM (\SI{88}{\percent}
load\_ref, \SI{94}{\percent} ref\_ref) and the low-reuse
$M{=}128,K{=}8,N{=}128$ corner saturates even harder
($>\SI{96}{\percent}$ in both), because back-to-back output
writes dominate. \emph{Compute-rich}: at the same deep-$K$ shape,
\textsf{load\_ref} delivers \SI{7.18}{\tera\flop\per\second}
($\sim\SI{90}{\percent}$ of the per-PE peak) — the configuration's
clean win. The \emph{distance between the two variants} at the same
shape is the cost of \emph{not} pre-staging $A$: at $K{=}3072$ the
throughput halves (\SI{7.18}{\tera\flop\per\second} $\to$
\SI{3.83}{\tera\flop\per\second}) because the second operand doubles
HBM pressure
and the kernel hits the BW ceiling. This is the operational
take-away — when the activation can live on-chip the composite
delivers near-peak GEMM; when it cannot the BW ceiling dominates and
throughput halves.
Figure~\ref{fig:gemm-util} validates that simulator-measured
efficiency tracks an analytic ideal-pipeline model (within
\SI{1.4}{ppt} across every swept shape), and
Figure~\ref{fig:gemm-stages} resolves the per-stage engine wall-clock
that backs the above interpretation.
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gemm_mac_utilization_theoretical_vs_measured.png}
\caption{GEMM MAC utilization and efficiency, theoretical vs.\ measured.
Tile-fill sets the ceiling: under-tile shapes (marked $\ast$) such as
$M{=}K{=}N{=}32$, $M{=}8$, and $K{=}8$ cannot fill the MAC tile and cap at
\SI{50}{\percent}, \SI{25}{\percent}, \SI{12.5}{\percent}. For
tile-filling shapes, efficiency climbs with tile count---from
\textasciitilde\SI{23}{\percent} at one tile to \textasciitilde%
\SI{78}{\percent} measured at 48 tiles ($K{=}3072$)---and the measured
bars track the analytic prediction within
\SIrange{10}{20}{\percent}.}
\caption{GEMM MAC utilization and efficiency, theoretical vs.\ measured
(\textsf{load\_ref} staging). Tile-fill sets the ceiling: under-tile
shapes (marked $\ast$) such as $M{=}K{=}N{=}32$, $M{=}8$, and $K{=}8$
cannot fill the MAC tile and cap at \SI{50}{\percent},
\SI{25}{\percent}, \SI{12.5}{\percent}. For tile-filling shapes,
efficiency climbs with tile count---from
\textasciitilde\SI{15}{\percent} at one tile to \textasciitilde%
\SI{90}{\percent} measured at 48 tiles ($K{=}3072$)---and the
measured bars track the analytic ideal-pipeline prediction within
\SI{1.4}{ppt} across every shape.}
\label{fig:gemm-util}
\end{figure}
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gemm_stage_breakdown.png}
\caption{Per-stage engine wall-clock (DMA in, Fetch, GEMM, DMA out) under
\textsf{load\_ref} staging. For the deep-$K$ shape ($K{=}3072$, 48 tiles)
the Fetch and GEMM stages are large and comparable
(\textasciitilde\SI{770}{} and \SI{785}{\nano\second}) while DMA-out is
negligible---a compute-rich, well-pipelined regime. For the low-reuse
shape ($M{=}128,K{=}8,N{=}128$) DMA-out grows to
\textasciitilde\SI{350}{\nano\second} and compute is small---a
data-movement-bound regime.}
\caption{Per-stage engine wall-clock (DMA in, Fetch, GEMM, DMA out)
under \textsf{load\_ref} staging. For the deep-$K$ shape ($K{=}3072$,
48 tiles) DMA in, Fetch, and GEMM are all close to
\textasciitilde\SI{785}{\nano\second}, running concurrently in a
tightly pipelined regime while DMA-out is negligible. For the
low-reuse shape ($M{=}128,K{=}8,N{=}128$, 16 output tiles) DMA-out
grows to \textasciitilde\SI{336}{\nano\second} while GEMM compute is
\textasciitilde\SI{262}{\nano\second}---a data-movement-bound regime
where the output write becomes the largest stage.}
\label{fig:gemm-stages}
\end{figure}
Two regularities stand out. (i) Utilization is governed by how completely
the problem fills the MAC tile: the three under-tile shapes are hard-capped
well below \SI{100}{\percent}, independent of how the kernel is issued.
(ii) Among tile-filling shapes, efficiency is governed by tile count---more
tiles amortize the one-time pipeline fill and issue cost, so the deep-$K$
shape reaches \textasciitilde\SI{78}{\percent} of peak while a single-tile
shape reaches only \textasciitilde\SI{23}{\percent}. The stage breakdown
explains why: with 48 tiles the GEMM and Fetch stages overlap and stay
busy, whereas the low-reuse shape spends most of its time moving data in
and out.
\subsection{Analysis and meaning}
The composite command does not manufacture bandwidth or MAC throughput; it
removes the two software-shaped obstacles between a GEMM and its hardware
roofline. By paying issue cost once per GEMM and chaining tile-stages
through token self-routing, it lets the scheduler keep the pipeline full,
so compute-rich shapes actually reach the efficiency their arithmetic
intensity allows ($\sim$\SI{78}{\percent} measured at 48 tiles), and
data-bound shapes actually reach their DMA bound instead of stalling on
command overhead. The close agreement between measured and theoretical
efficiency (Figure~\ref{fig:gemm-util}) is also the report's primary
validation that KernBench's latency model is faithful in the regime that
matters. The hardware implication is concrete: a single-command,
self-routing tile pipeline is the issue mechanism that makes the MAC array
usable, and it is a prerequisite---not a luxury---for the fused attention
kernel of \S\ref{sec:gqa}.
The composite command does not manufacture bandwidth or MAC
throughput; it removes the two software-shaped obstacles between a
GEMM and its hardware roofline. By paying the issue cost once per
GEMM rather than once per tile-stage, and by chaining tile-stages
through token self-routing, it lets the scheduler keep every stage
busy on a different tile, so compute-rich shapes actually reach the
efficiency their arithmetic intensity allows
($\sim$\SI{90}{\percent} of GEMM peak at the deep-$K$
\textsf{load\_ref} corner) and data-bound shapes actually reach
their HBM-BW ceiling instead of stalling on command overhead.
The split between \textsf{load\_ref} and \textsf{ref\_ref} also makes
the hardware-software boundary explicit: when the activation fits in
on-chip storage the composite is BW-headroom-rich and saturates the
GEMM engine; when it does not, both operands compete for the same
HBM port and the kernel is BW-bound. This is the lever the GQA
kernel of \S\ref{sec:gqa} reaches for next — keeping the right
working set on-chip so the composite pipeline lands in the
compute-rich regime rather than the BW-bound one.
\subsection{Why composite, and not kernel-orchestrated async loading?}
\label{sec:gemm-vs-async}
A reader familiar with double-buffered GEMM kernels on conventional
hardware may ask: why a hardware-side composite command at all? Why
isn't the obvious kernel-level pattern --- async-load each operand,
overlap with compute, accumulate --- sufficient?
To answer this concretely we contrast composite against two
kernel-orchestrated baselines that have access to the same single-op
primitives the platform exposes (\textsf{tl.load} for async DMA into
TCM, \textsf{tl.dot} for a single-op GEMM command on TCM-resident
operands, \textsf{tl.store} for a DMA write-back). Both baselines
pre-stage activation $A$ identically to \textsf{load\_ref}, so the
window measured is the engine-pipeline window with $A$'s up-front DMA
excluded for all three kernels.
\paragraph{Async-full (naive).} A single \textsf{tl.load(A)} followed by
a single \textsf{tl.load(B)} (async, queued behind $A$),
\textsf{tl.dot(A, B)}, and \textsf{tl.store(out)}. The kernel issues
four commands total. The decisive constraint is that
\textsf{tl.dot}'s \textsf{\_await\_pending(b)} blocks the GEMM
command until the \emph{entire} $B$ has landed in TCM --- there is no
way at the runtime API surface to express ``start computing on
tile 0 of $B$ while tile 1 is still in flight.'' Load-of-$B$ and
GEMM therefore serialize.
\paragraph{Async-tiled (chunked prefetch).} The kernel-level workaround is to
split $B$ along $K$ into \textsf{TILE\_K}-sized chunks, issue async
\textsf{tl.load}s for those chunks, issue one \textsf{tl.dot} per
chunk (each blocking only on its own $b_i$), and accumulate via
\textsf{out = out + tl.dot(...)}. This is the standard
double-buffered-GEMM pattern transcribed to the single-op primitives.
We measure two versions of this kernel that differ only in
\emph{prefetch depth}:
\textbf{depth-2 (TCM-bounded)} keeps at most two B-chunks in flight
at any time by issuing the next \textsf{tl.load} just before each
\textsf{tl.dot}; \textbf{depth-$\infty$ (queue-all)} issues all $N_K$
B-chunk loads up front so the DMA engine has the deepest possible
request queue. For a $K{=}3072$ shape both versions emit
48 $A$-chunk loads + 48 $B$-chunk loads + 48 \textsf{tl.dot}s + 47
elementwise adds + 1 store = 192 host-side commands; the only
difference is the temporal interleaving of B-load and dot dispatches.
\paragraph{Why two depths.} The depth distinction matters because the
async-full kernel and the depth-$\infty$ async-tiled kernel both pin the
\emph{entire} $B$ in TCM simultaneously --- $K \cdot N \cdot 2$
bytes. The on-PE scratch is capped at \SI{1}{\mebi\byte}
(\texttt{topology.yaml: pe\_tcm.kernel\_scratch\_mb=1}), so an LLM-scale
attention $B = K_{\text{KV}} \times d_{\text{head}}$ at
$K_{\text{KV}}=4096, d_{\text{head}}=128$ already needs
\SI{1}{\mebi\byte}, and at $K_{\text{KV}}=8192$ it needs \SI{2}{\mebi\byte}
--- past the cap. async-full and queue-all async-tiled are therefore
not just slower than composite but \emph{architecturally infeasible}
at LLM context length. The depth-2 async-tiled kernel is the only
kernel-level
variant whose peak TCM footprint stays
$O(2 \cdot \textsf{TILE\_K} \cdot N)$ regardless of $K$, the same
order as composite's per-tile streaming buffer. It is the apples-to-apples
comparison.
\paragraph{Per-PE throughput.} Figure~\ref{fig:gemm-async} reports
per-PE TFLOP/s for all four kernels side-by-side. The single-op
fast-path in the dispatch cost model (\S\ref{sec:congestion}) is
enabled, so every single-op command the async kernels emit --- DMA
descriptors and \textsf{tl.dot}/\textsf{tl.add} alike --- is charged the
light 8-cycle \textsf{FIXED}, not the 40-cycle composite control-path
cost; neither async-tiled variant carries an inflated per-command cost.
\paragraph{The $K{=}3072$ corner, concretely.} We take this shape
($M{=}32, K{=}3072, N{=}32$) as the running example throughout the
mechanism discussion because it is the regime where the four
kernels spread the widest --- the deepest $K$ in the sweep maps
onto $K / \textsf{TILE\_K} = 48$ hardware tiles, so per-tile costs
amplify into the largest measurable gap. The work content is
identical for all four kernels: $\sim$6.3 M f16 MACs and
$\sim$386 KiB of $B$ traffic from HBM, which together require
$\sim$786 ns of GEMM-engine compute and $\sim$781 ns of DMA on a
saturated per-PE link. What differs is the number of host commands
the same work is decomposed into --- 2 for composite (one
\textsf{tl.load(A)} plus one composite), 4 for async-full, and 192
for either async-tiled variant (48 $A$-loads + 48 $B$-loads + 48
\textsf{tl.dot}s + 47 elementwise adds + 1 store). The
engine-pipeline-window throughput tracks that decomposition closely:
composite reaches \SI{7.18}{\tera\flop\per\second} (post-overlap
limit, only \SI{10}{\percent} below the \SI{8}{\tera\flop\per\second}
per-PE GEMM peak), async-full \SI{3.91}{\tera\flop\per\second}
(DMA and compute serialize on a single big dot), and both async-tiled
variants $\sim$\SI{2.53}{\tera\flop\per\second} (192 commands' worth of
structural dispatch cost --- even at the light per-command rate ---
accumulates on the wall). The next paragraph attributes those gaps to
specific simulator mechanisms.
\begin{figure*}[t]
\centering
\includegraphics[width=\linewidth]{gemm_composite_vs_async_tflops.png}
\caption{Per-PE achieved TFLOP/s for the same shape sweep run under
four issuance patterns: composite (one command, scheduler streams
per-tile internally), async-full (one \textsf{tl.dot} on
fully-loaded $B$), async-tiled with depth-2 double-buffer
(TCM-bounded; the only kernel-level variant that scales to LLM
context length), and async-tiled with depth-$\infty$
(all B-tiles queued up front; included as a sanity check that the
prefetch depth is \emph{not} what separates the async-tiled kernel from
composite). All curves exclude $A$'s up-front DMA from the
measurement window. Composite wins at every full-tile shape; the
depth-2 and depth-$\infty$ async-tiled kernels deliver
\emph{indistinguishable} throughput (e.g. \SI{2.522}{} vs.\
\SI{2.544}{\tera\flop\per\second} at $K{=}3072$), confirming that
prefetch depth is not the lever --- the structural per-command
dispatch cost is. The gap between composite and the async-tiled
kernels grows with $K_{\text{useful}}$
(\textbf{$\sim$$2.8\times$} at $K{=}3072$, where the async-tiled
kernels emit 192 host commands while composite emits one).
The under-tile corner $M{=}128,K{=}8,N{=}128$ inverts: composite's
per-tile orchestration overhead exceeds the per-tile useful work,
and all three async kernels beat it.}
\label{fig:gemm-async}
\end{figure*}
\paragraph{Decomposing the gap.} Three structural mechanisms separate
composite from the kernel-level baselines, and they layer.
\emph{1. Inter-engine token routing happens below the host-side
dispatch path.} The composite encodes the full
\textsf{DMA\_READ}$\to$\textsf{FETCH}$\to$\textsf{GEMM}$\to$\textsf{STORE}$\to$\textsf{DMA\_WRITE}
pipeline once. The scheduler's tile-feeder loop then emits one
\emph{tile token} per HW tile inside that one composite, and each
token self-routes between engines after each stage finishes. The
per-tile token routing is a scheduler-internal event, not a fresh
host command, so it does not pay the structural CPU dispatch cost.
At $K{=}3072$ the composite emits 48 tile tokens that flow
fully-pipelined through five stages each --- 240 inter-engine
hand-offs total --- behind a single command from the host's point of
view.
\emph{2. \textsf{tl.dot} cannot replicate that per-tile pipeline at the
kernel level.} A single-op GEMM command is handled on the GEMM engine as
a single monolithic compute timeout for the supplied $M{\times}K{\times}N$;
there is no internal token loop that would let a streaming DMA of
$B[i{+}1]$ overlap with the GEMM of $B[i]$ inside one
\textsf{tl.dot}. The user can only recover inter-tile overlap by
emitting one \textsf{tl.dot} per chunk --- which is exactly what the
async-tiled baseline does, at the price of $N$ host-side commands.
\emph{3. The host-side dispatch cost the async-tiled baseline pays is
structural, not modelling slack.} KernBench charges every host-emitted
command a structural CPU dispatch cost $d_{\text{cmd}} = \textsf{FIXED} +
b_{\text{logical}} \cdot R$ (\S\ref{sec:congestion}). The cost model is
deliberately charitable to the async kernels here: only a
\textsf{CompositeCmd} pays the 40-cycle control-path \textsf{FIXED} (it
alone drives a scheduler-built tile-feeder plan), while \emph{every}
single-op command --- the 96 DMA descriptors, the 48 \textsf{tl.dot}s,
and the 47 elementwise adds alike --- pays only the light 8-cycle
\textsf{FIXED}, calibrated to descriptor-ring-push / single-instruction
issue patterns in modern accelerators (NVIDIA Hopper TMA $\sim$1 ISA
cycle, a single \textsf{mma.sync} one instruction, AMD AQL packet writes
$\sim$5--15 cycles). So the async-tiled baseline is \emph{not} penalized
by an inflated per-command cost on \emph{any} of its operations. Yet it
still emits 192 host commands against composite's one, so $\sim$192
light dispatches accumulate to $\sim$\SI{1.5}{\micro\second} of
structural PE\_CPU time that composite never pays --- composite hides
its 48 tiles' 240 inter-engine hand-offs as scheduler-internal events
(mechanism 1). Layered on top, the dependency chain on the running
accumulator serializes the 47 adds on the math engine. The net result
is that the async-tiled kernel runs $\sim$2.8$\times$ slower than
composite at $K{=}3072$ despite getting the inter-chunk overlap right
--- down from $\sim$6.3$\times$ under the earlier uniform-40 cost model,
because D8 removed the per-command overcharge, but \emph{not} closed:
the residual gap is the command-count structure (one composite vs.\ 192
single-ops), not a modelling artifact.
\emph{4. Prefetch depth is not the lever, command count is.} The
depth-2 and depth-$\infty$ async-tiled kernels land within \SI{1}{\percent}
of each other at every measured shape (Figure~\ref{fig:gemm-async}).
This is the diagnostic against a natural objection: ``surely the
async-tiled kernel was just under-prefetching; deepen the queue and the
DMA$\to$GEMM overlap recovers.'' Deepening the prefetch queue
changes \emph{when} DMA descriptors hit the engine but not their
total number or the structural dispatch cost they each pay. The
$\sim$2.8$\times$ gap to composite is not a prefetch-depth gap; it
is the gap between
``one composite command with internal per-tile token routing'' and
``$N_K$ host-side dot/add commands, each charged separately.''
The depth-2 kernel additionally constrains peak TCM occupancy to
$O(2 \cdot \textsf{TILE\_K} \cdot N)$, matching composite's per-tile
streaming buffer; the depth-$\infty$ kernel needs the full $B$ in
TCM, which makes it infeasible at LLM context length even if the
throughput were competitive.
\paragraph{Where the composite advantage doesn't apply.} The shape
$M{=}128,K{=}8,N{=}128$ inverts the picture: composite delivers
\SI{0.66}{\tera\flop\per\second} and both async kernels reach
\SI{1.21}{\tera\flop\per\second}. The reason is consistent with the
analysis above and explicit in the simulator state. Composite emits
16 tile tokens (one per output tile) for this shape, each carrying
$K_{\text{useful}}{=}8$ MACs across a TILE\_K$=64$ pipeline ---
$12.5\%$ of the hardware tile's MAC slots are useful, the rest is
K-padding. The per-tile inter-engine hand-off cost stays the same
regardless. When the per-tile useful work is small enough that
hand-off overhead exceeds the GEMM work itself, a single-op
\textsf{tl.dot} --- which submits one monolithic GEMM command with no
per-tile orchestration --- wins. The take-away is the bound on
composite's value: it amortizes useful per-tile compute, not
padding-dominated under-tile shapes. Real kernels at this corner are
better served by reshape-into-batched-GEMM transforms that move
under-tile $K$ into a tile-filling dimension before reaching the GEMM
engine, which is exactly what the GQA decode kernel of
\S\ref{sec:gqa} does for its $K_{\text{useful}}=\text{head\_dim}=128$
inner reduction.
\paragraph{Summary of the comparison.} Composite is the only one of
the four kernels to combine (a) macro-command dispatch at the host
boundary (amortizing the structural CPU cost across all the work a
single GEMM does), (b) scheduler-internal per-HW-tile streaming of
DMA$\rightleftarrows$compute, and (c) TCM-bounded streaming buffer.
Kernel-orchestrated async kernels can have any two of those, not all
three: async-full pays one host dispatch (a) but forfeits per-tile
overlap (b) and pins all of $B$ in TCM (c); depth-$\infty$ async-tiled
achieves inter-chunk overlap but at $N_K$ host dispatches and
full-$B$ TCM occupancy; depth-2 async-tiled fixes the TCM
footprint (c) but still pays $N_K$ host dispatches. The two corners
where async catches up
(small-$K$ where composite has nothing useful to amortize;
under-tile shapes where per-tile useful work is sub-token) are
diagnostic of \emph{where composite is the wrong tool}, not of a
slack the user kernel could close.
docs/report/1H-codesign-paper/sections/04-allreduce.tex
+222 -47
View File
@@ -1,46 +1,215 @@
\section{PE\_IPCQ and Collective Communication}
\label{sec:allreduce}
\subsection{Why it is needed}
Distributing a transformer across devices turns every tensor-parallel
layer into a collective: partial results computed on different PEs, CUBEs,
and SIPs must be summed and redistributed with an all-reduce. If that
collective is handled by the host or by a generic DMA path, three problems
appear. The reduction traffic competes with the kernel's own compute DMA
on the same links, causing head-of-line blocking; there is no efficient
peer-to-peer ring primitive, so data takes extra hops; and the ordering is
hard to make deterministic. The hardware question is how to perform
collectives \emph{on the device}, overlapped with compute and reproducible
run-to-run.
and SIPs must be summed and redistributed with an all-reduce. Underneath
the algorithm this is fundamentally a PE-to-PE problem---many short
messages flowing between neighbors as the reduction proceeds. The natural
software realization is a per-direction ring buffer whose head and tail
pointers the producer and consumer update atomically and poll, but our
H2 2025 report measured this scheme end-to-end and found that the
atomic-pointer traffic together with the consumer's polling loop dominate
the per-message cost: the queue itself becomes the bottleneck well before
the link runs out of bandwidth, and a pure-SW collective spends most of
its time on metadata rather than on actually moving partials. A second,
orthogonal problem is link sharing---if the collective rides the kernel's
generic DMA path it competes with the GEMM's compute traffic on the same
wires, so a large tile transfer head-of-line-blocks a pending reduction.
This section asks how a dedicated hardware primitive can lift queue
management off the software path entirely and, in the same design, stop
collective traffic from stalling behind compute traffic, so the all-reduce
runs overlapped with compute and at the interconnect's physical limit.
\subsection{Design}
KernBench models a dedicated per-PE collective engine, \textbf{PE\_IPCQ}
(inter-PE communication queue). It is a control-plane block: it owns the
ring-buffer address arithmetic, head/tail pointers, peer-pointer caches,
backpressure, and the four-direction (N/S/E/W) neighbor map, with eight
ring buffers per PE (four directions $\times$ \{tx, rx\}). Crucially,
PE\_IPCQ does \emph{not} move data itself---it delegates the actual
transfer to PE\_DMA, keeping a clean control/data split. To stop
collective traffic from blocking compute, PE\_DMA is extended into a
two-channel virtual-channel model: \texttt{vc\_compute} carries tile
load/store for GEMM and vector math, \texttt{vc\_comm} carries IPCQ sends,
each with an independent state machine. The same physical link is shared
but progresses in chunks (\SI{256}{\byte}), so a large GEMM DMA does not
lock the link end-to-end against a pending reduction. On top of this
substrate the collective runs a hierarchical local-reduce / global
all-reduce-broadcast schedule across whatever inter-device topology the
configuration specifies.
The proposed block is \textbf{PE\_IPCQ} (inter-PE communication queue),
a small controller dropped into every PE next to PE\_DMA, PE\_GEMM,
PE\_MATH, and PE\_TCM. It is a control-plane block---it holds no
payload data---and consists of three pieces: a per-direction
\emph{QPair register file} (\textasciitilde\SI{576}{\byte} of flip-flops
covering up to eight directions), a combinational slot-address
generator and backpressure comparator, and a credit injector/receiver
wired to the NoC. Each QPair holds the local pointers
(\texttt{my\_head}, \texttt{my\_tail}), shadowed views of the peer's
(\texttt{peer\_head\_cache}, \texttt{peer\_tail\_cache}), the local and
peer rx-buffer physical bases, ring depth and slot size (both
power-of-two), and the peer's credit-target address. The ring data
itself lives in a reserved slot region of TCM (or PE-local HBM, or
cube-shared SRAM), addressed by the QPair registers; PE\_IPCQ never
touches the bytes, only the pointers. The PE\_CPU sees the controller
as an MMIO peripheral and the N/S/E/W direction labels as logical
ports---one kernel image runs across a 1D ring, 2D mesh, or 2D torus
because the topology only changes which peers the QPair registers
point at.
\emph{Initialization.} The host CCL backend brings the whole machine
to a usable state before any kernel runs.
\texttt{init\_process\_group(backend="ahbm")} loads \texttt{ccl.yaml},
resolves the algorithm + topology + buffer\_kind + slot configuration,
allocates an rx ring-buffer region on every participating PE so every
rank now knows every other rank's \texttt{rx\_base\_pa}, and fans out
an \texttt{IpcqInitMsg} that writes the QPair register file on each
PE\_IPCQ over MMIO. The same fan-out wires the per-direction
credit-return channel: each PE\_IPCQ records its peer's credit-target
address and a back-pointer that the credit receiver will use to update
the right QPair. After init, every register the runtime needs is
preloaded; nothing in the kernel path allocates memory, walks a table,
or talks to the host.
\emph{Send.} When the kernel executes \texttt{tl.send(dir="E",
src\_addr, nbytes)}, the PE\_CPU performs a single MMIO write into
PE\_IPCQ describing the request (direction, source address/space,
length, sender handle). The controller evaluates backpressure in one
combinational compare---\texttt{(my\_head $-$ peer\_tail\_cache) $<$
n\_slots}---and, on a hit, the slot-address generator returns
\texttt{dst = peer\_rx\_base\_pa + (my\_head \% n\_slots) $\times$
slot\_size} in one to two cycles. PE\_IPCQ then emits one
\texttt{IpcqDmaToken} on its dedicated port to PE\_DMA's
\texttt{vc\_comm} channel, carrying the data descriptor (\texttt{src,
dst, nbytes}) plus a small piggyback header (\texttt{sender\_seq =
my\_head, src\_coord, direction}). \texttt{my\_head} is incremented in
the same cycle---a local flip-flop bump, not a cross-PE atomic---and
\texttt{tl.send} returns fire-and-forget. If the backpressure compare
fails, the controller stalls the CPU in either of two modes selected
at init: \texttt{poll} (CPU re-reads a status CSR) or \texttt{sleep}
(controller asserts a wake event when a credit arrives). Both modes
are benchmarked in the results.
\emph{Transport.} PE\_DMA is the only block that touches the fabric,
and it has been extended for IPCQ in two ways. First, it now exposes
two virtual channels---\texttt{vc\_compute} for GEMM/Math
\texttt{TileToken}s and \texttt{vc\_comm} for \texttt{IpcqDmaToken}s
---with independent state machines and a chunk-level
(\SI{256}{\byte}) weighted round-robin arbiter on the shared physical
link. A large GEMM tile DMA can no longer monopolize the link against
a pending IPCQ send, enabling compute and communication to make
progress independently as an architectural property of the design. Second, on the sender side PE\_DMA packs the piggyback
header into the same flit train as the data, and on the receiver side
it runs the I6 \emph{atomic terminal handler}: pay the per-flit
bottleneck-BW drain, then, in one indivisible block, (i) write the
payload into \texttt{MemoryStore} at \texttt{dst\_addr} (the
receiver's ring slot) and (ii) forward an \texttt{IpcqMetaArrival}
$\{\texttt{sender\_seq, dst\_addr}\}$ to the local PE\_IPCQ over a
PE-internal wire. No yield, no PE\_CPU interaction, no cache-coherence
round trip---the bytes land in TCM and the metadata reaches the
controller in the same cycle.
\emph{Receive and credit return.} PE\_IPCQ's Meta Extractor
range-matches the incoming \texttt{dst\_addr} against each direction's
$[\texttt{rx\_base}, \texttt{rx\_base} + \texttt{n\_slots} \times
\texttt{slot\_size})$ window---unambiguous even when two directions
share a peer, as in a 2-rank bidirectional ring---and updates
\texttt{peer\_head\_cache[d] := max(prev, sender\_seq + 1)}, releasing
any \texttt{tl.recv} blocked on direction $d$. When the kernel
eventually consumes the slot, the controller increments
\texttt{my\_tail} and emits a \SI{16}{\byte} credit packet on the
dedicated fast-path channel preloaded at init; the latency is the
real fabric path (per-node overhead + edge propagation + 16 B /
bottleneck BW), so an in-cube credit returns faster than a cross-SIP
credit and the model carries no magic constants. The sender's PE\_IPCQ
absorbs the credit, advances \texttt{peer\_tail\_cache}, and
de-asserts backpressure if it was stalled. The pointer-synchronization
problem the software baseline solved with explicit atomic RMWs and
polling loops has been split in two and dissolved into existing
traffic: \emph{head} updates ride on the DMA payload itself, with the
receiver's PE\_DMA performing the data + metadata write in a single
atomic step, so the sender never blocks waiting for the receiver to
see it; \emph{tail} updates ride on a dedicated 16 B credit on a
side-channel, with no software in the loop. Send becomes a single
MMIO write, receive becomes a flip-flop read, backpressure becomes
one 64-bit subtract, and the per-message cost in IPCQ is dominated by
the link traversal rather than by metadata bookkeeping---which is
what the H2 2025 measurement showed the pure-software queue could
not achieve. On top of this substrate the collective runs a
hierarchical local-reduce / global all-reduce-broadcast schedule
across whatever inter-device topology the configuration specifies,
with the IPCQ ring buffer placed in on-PE TCM, PE-local HBM, or
cube-shared SRAM---the third knob the results section sweeps.
\subsection{Design alternatives}
\label{sec:ipcq-alternatives}
PE\_IPCQ is one point in a small space of hardware mechanisms for
moving a short message from one PE to a neighbor and signalling its
arrival. Three established alternatives anchor the space, each the
HW realization of a familiar host-networking idea: a \emph{doorbell +
polling} scheme (the classic MMIO doorbell---write the payload by DMA,
write a doorbell, let the peer poll or take an interrupt); a
\emph{hardware message queue} (HMQ, the NVLink-style descriptor engine
that pushes a queue entry to the peer, with large payloads still
riding a second DMA); and a \emph{completion-queue} design (RDMA-CQ,
the InfiniBand/RoCE pattern where a DMA write auto-posts a completion
entry the peer's CQ polls). PE\_IPCQ is the fourth: a hardware ring
with credit return, splitting the control plane into PE\_IPCQ and the
data plane into PE\_DMA, with head updates riding the payload and tail
updates riding a 16\,B side-channel credit (\S\ref{sec:allreduce}).
\begin{figure*}[t]
\centering
\includegraphics[width=0.78\linewidth]{ipcq_alternatives_architecture_flow.png}
\caption{Per-send data and control flow for the four PE-to-PE
signalling mechanisms (sender\,$\rightarrow$\,NoC\,$\rightarrow$\,receiver).
Doorbell and RDMA-CQ each issue two fabric transactions (payload then
doorbell / completion) and leave the peer polling or taking an
interrupt; HMQ adds a dedicated descriptor engine but still moves large
payloads on a second DMA; PE\_IPCQ folds head-pointer signalling into
the payload flit train and returns the tail credit on a side channel,
so a send is one MMIO write and a receive is a flip-flop read. This is
a \emph{design schematic}, not a measured comparison.}
\label{fig:ipcq-arch}
\end{figure*}
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{ipcq_alternatives_decision_matrix.png}
\caption{Why the ring+credit design was chosen, across five criteria:
single-send latency, whether the host CPU sits on the critical path,
whether the receiver must poll or take a wake-up interrupt, whether the
control and data datapaths are duplicated, and whether the mechanism is
right-sized for single-owner PE-to-PE traffic (rather than a
multi-tenant fabric). PE\_IPCQ is the only design that clears every
criterion. The accompanying per-send step-count tally
($\sim$28 control events for IPCQ versus $\sim$38 for HMQ, $\sim$53 for
RDMA-CQ, and $\sim$56 for doorbell+polling) is an \emph{illustrative}
order-of-magnitude comparator over hand-counted pipeline steps---not a
simulator measurement. The measured, simulator-grounded results follow
in the next subsection.}
\label{fig:ipcq-decision}
\end{figure}
The qualitative comparison motivates the design but is not a
quantitative claim: the cycle-step tallies above are hand-counted
control events, deliberately separated from the measured latencies that
follow. Everything in the results subsection runs on the PE\_IPCQ
substrate and is simulator-grounded.
\subsection{Results}
Following the milestone-evaluation convention, the collective sweep builds
its own six-device (six-SIP, $2\times3$) configurations---distinct from
the two-SIP default of Table~\ref{tab:hw}---and measures all-reduce
latency as a function of payload size for three inter-device topologies:
a 1D ring, a 2D mesh (no wrap), and a 2D torus. Table~\ref{tab:allreduce}
and Figure~\ref{fig:allreduce-cmp} report the result.
All measurements in this section run on the PE\_IPCQ substrate
described above; the topology sweep is intended to characterize how
effectively the proposed mechanism exposes the underlying interconnect's
properties, not to compare PE\_IPCQ against an alternative
communication primitive. Following the milestone-evaluation convention,
the collective sweep builds its own six-device (six-SIP, $2\times3$)
configurations---distinct from the two-SIP default of
Table~\ref{tab:hw}---and measures all-reduce latency as a function of
payload size for three inter-device topologies: a 1D ring, a 2D mesh
(no wrap), and a 2D torus (Figure~\ref{fig:allreduce-topo}).
Table~\ref{tab:allreduce} and Figure~\ref{fig:allreduce-cmp} report the
result.
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{allreduce_topology.png}
\caption{The three six-device ($2\times3$) inter-device topologies the
collective sweep runs over, and the hierarchical local-reduce /
global all-reduce-broadcast schedule mapped onto each: a 1D ring, a 2D
mesh (no wrap-around), and a 2D torus (wrap-around links on both axes).
The torus's wrap links shorten the worst-case reduction path, which is
what the latency sweep below rewards.}
\label{fig:allreduce-topo}
\end{figure}
\begin{table}[t]
\centering
@@ -52,11 +221,11 @@ per-PE payload. Lower is better; the torus wins at every size.}
\toprule
\textbf{Bytes/PE} & \textbf{2D mesh} & \textbf{Ring 1D} & \textbf{2D torus} \\
\midrule
256 & 2667 & 2365 & 1701 \\
4{,}096 & 4450 & 4082 & 3038 \\
16{,}384 & 8900 & 8217 & 6403 \\
65{,}536 & 26705 & 24766 & 19865 \\
98{,}304 & 38574 & 35798 & 28840 \\
256 & 4189 & 3883 & 2957 \\
4{,}096 & 5566 & 5240 & 4031 \\
16{,}384 & 10016 & 9376 & 7396 \\
65{,}536 & 27821 & 25925 & 20858 \\
98{,}304 & 39690 & 36957 & 29833 \\
\bottomrule
\end{tabular}
\end{table}
@@ -64,13 +233,19 @@ per-PE payload. Lower is better; the torus wins at every size.}
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{allreduce_comparison.png}
\caption{All-reduce latency vs.\ per-PE payload for the three topologies,
against the analytic torus model and an external full-system simulator
(FSIM) reference. The measured torus tracks the analytic curve within a
small constant factor; the FSIM single-device point
(\SI{366}{\micro\second}) sits an order of magnitude above the
KernBench algorithmic latency, illustrating the difference between an
achievable-kernel number and a full end-to-end-stack number.}
\caption{All-reduce latency vs.\ per-PE payload for the three PE\_IPCQ
topologies, against the analytic torus model. The isolated point in
the top panel (\SI{366}{\micro\second}) is the only data point
available from the H2 2025 software-queue measurement campaign---an
intra-device all-reduce across 16 CUBEs on a single device. A true
6-device inter-SIP measurement under the same software queue was not
collected, so this point in fact \emph{understates} the SW-queue cost
for the multidevice configuration KernBench measures here; even so the
proposed PE\_IPCQ inter-device curves sit roughly an order of
magnitude below it, which is the headline SW-vs-HW comparison this
section makes. The measured torus tracks the analytic curve within a
small constant factor, the gap reflecting real link serialization that
the analytic model idealizes away.}
\label{fig:allreduce-cmp}
\end{figure}
@@ -78,9 +253,9 @@ achievable-kernel number and a full end-to-end-stack number.}
\centering
\includegraphics[width=\linewidth]{allreduce_buffer_kind.png}
\caption{Effect of IPCQ staging-buffer placement (2D torus). At
\SI{64}{\kibi\byte}/PE, TCM staging (\SI{19865}{\nano\second}) beats HBM
(\SI{23081}{\nano\second}) by \textasciitilde\SI{14}{\percent} and SRAM
(\SI{32201}{\nano\second}) by \textasciitilde\SI{38}{\percent}; at small
\SI{64}{\kibi\byte}/PE, TCM staging (\SI{20858}{\nano\second}) beats HBM
(\SI{24074}{\nano\second}) by \textasciitilde\SI{13}{\percent} and SRAM
(\SI{33194}{\nano\second}) by \textasciitilde\SI{37}{\percent}; at small
payloads the three are indistinguishable.}
\label{fig:allreduce-buf}
\end{figure}
docs/report/1H-codesign-paper/sections/05-gqa.tex
+349 -59
View File
@@ -1,8 +1,6 @@
\section{Fused Grouped-Query Attention}
\label{sec:gqa}
\subsection{Why it is needed}
Attention is the 1H focus, and it is where the two preceding optimizations
have to come together. Grouped-Query Attention (GQA) shrinks the KV cache
by sharing each KV head across a group of query heads (here $h_q=8$ query
@@ -16,12 +14,99 @@ realizing it as a fast \emph{fused} kernel needs both building blocks from
this report: efficient GEMM issue (\S\ref{sec:gemm}) for the
$Q\!\cdot\!K^{\top}$ and $P\!\cdot\!V$ products, and an efficient on-device
reduction (\S\ref{sec:allreduce}) for the multi-user and
sequence-parallel KV reductions. This section is the capstone: the fused
sequence-parallel KV reductions. \emph{Fused} here is meant in the
FlashAttention sense---$Q\!\cdot\!K^{\top}$, the online softmax, and
$P\!\cdot\!V$ collapse into a single kernel that never materializes the
score matrix---and, beyond that, the cross-device KV reduction is absorbed
into the same kernel (on PE\_IPCQ) rather than issued as a separate
all-reduce. This section is the capstone: the fused
kernel that uses the composite command and PE\_IPCQ at the same time.
Multi-head attention (MHA) was studied in prior work and serves here as
the established baseline rather than being re-derived.
\subsection{Design}
\subsection{Data Placement Policy}
\label{sec:gqa-placement}
Long-context decode is bound by the KV cache, so the first-order design
question is how to place that cache---and the running softmax state it
feeds---across the two hardware axes the machine exposes: the CUBEs and,
within each CUBE, the PEs (here $C{=}8$ CUBEs $\times$ $P{=}8$ PEs, for
$C\!\cdot\!P{=}64$ attention engines over one KV-head group on the
LLaMA-3.1-70B target). Each axis can \emph{replicate} the KV cache or
\emph{shard} it, and a shard can run along the sequence dimension
$S_{kv}$ or the head dimension $d_{\text{head}}$. The cross product is a
small, enumerable taxonomy; six placements span its meaningful corners
(Figure~\ref{fig:gqa-kv-sharding}).
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gqa_long_ctx_6cases_kv_sharding_diagram.png}
\caption{The six KV-placement strategies, drawn on the
$S_{kv}\!\times\!d_{\text{head}}$ KV tensor (rows = sequence, columns =
head dimension). Cube colour bands and dashed PE dividers show which
axis each level shards. Cases~1--3 either replicate the cache or shard
it on a single axis (8-way at most); Cases~4--6 reach a full 64-way
split, three different ways: Case~4 splits $S_{kv}$ across CUBEs and
$d_{\text{head}}$ across PEs, Case~5 the mirror, and Case~6~$\star$
splits $S_{kv}$ on \emph{both} axes.}
\label{fig:gqa-kv-sharding}
\end{figure}
Two quantities decide which placement is viable, and they pull against
each other. The first is \textbf{per-PE KV memory}. With a per-PE HBM
budget of \SI{6.0}{\giga\byte} and \SI{1.76}{\giga\byte} of attention
weights resident, the KV headroom is \SI{4.24}{\giga\byte} per PE. At a
production context of $S_{kv}{=}1\,\text{M}$ tokens the unsharded cache
is \SI{40}{\giga\byte}/PE (Case~1), an 8-way shard is \SI{5}{\giga\byte}
(Cases~2--3)---both \emph{over} the headroom---while only the 64-way
placements bring it to \SI{640}{\mega\byte}/PE (Cases~4--6), comfortably
inside budget. Memory alone therefore eliminates Cases~1--3 at long
context. The second quantity is \textbf{communication per token}, and it
is what separates the three survivors. Sharding $d_{\text{head}}$
(Cases~4--5) makes each PE hold only a slice of every head, so the
$Q\!\cdot\!K^{\top}$ score is \emph{partial} and must be all-reduced
across the slice owners on every token---a reduction whose volume scales
with $S_{kv}$ ($\sim$\SI{166}{\mega\byte}/token analytically, intra-CUBE
on the NoC for Case~4, inter-CUBE on UCIe for Case~5). Case~6~$\star$
instead shards $S_{kv}$ on both axes, so every PE computes a
\emph{complete} score over its own token range and only the small running
softmax state $(m,\ell,O)$ is merged across PEs
($\sim$\SI{6.2}{\mega\byte}/token)---a $\sim$27$\times$ lighter collective
than the $d_{\text{head}}$-split designs.
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gqa_long_ctx_6cases_summary.png}
\caption{Long-context placement analysis at $S_{kv}{=}1\,\text{M}$ tokens.
\emph{Left:} the per-PE HBM budget---\SI{1.76}{\giga\byte} of attention
weights leave \SI{4.24}{\giga\byte} of KV headroom (red line).
\emph{Middle:} per-PE KV memory per case (log scale); only the 64-way
placements (Cases~4--6, \SI{640}{\mega\byte}) clear the headroom, while
the unsharded (\SI{40}{\giga\byte}) and 8-way (\SI{5}{\giga\byte}) cases
overflow. \emph{Right:} analytical communication per token (log scale);
the $d_{\text{head}}$-split Cases~4--5 pay a partial-score all-reduce
($\sim$\SI{166}{\mega\byte}/token) that the both-axes-$S_{kv}$ split of
Case~6~$\star$ avoids ($\sim$\SI{6.2}{\mega\byte}/token, merging only the
softmax state). On these two axes Case~6 (marked $\star$ in the figure)
is the single placement that lands both inside the memory budget and at
low per-token communication; whether that combination is the right one to
pick is a regime-specific question taken up next.}
\label{fig:gqa-budget}
\end{figure}
These two costs---per-PE KV memory and per-token communication---are
intrinsic properties of each placement, fixed by how it shards the cache
and independent of the workload regime. They do not by themselves name a
winner: a short prompt where the whole cache fits on one PE values low
communication and tolerates replication, whereas a million-token decode
is bound by the memory wall and will pay communication to escape it.
Which placement is appropriate is therefore a per-regime question, which
the short- and long-context subsections that follow answer by running the
options on the simulator and reading off latency, traffic, and the
redundant compute each one induces.
\subsection{Inference with Short-Context Length}
\label{sec:gqa-short}
The fused GQA kernel issues its matrix products as scheduler-managed
composite commands and keeps the online-softmax merge and the cross-device
@@ -37,74 +122,273 @@ overlaps score computation; and per-tile \emph{scratch recycling} that
keeps the running softmax accumulators ($m,\ell,O$) in a persistent arena
while freeing per-tile temporaries, so the kernel fits the
\SI{1}{\mebi\byte} scratch budget across many tiles. A further refinement
that restructures the decode step into two stateful composites (a named
\textsf{softmax\_merge} recipe) is designed but not yet wired into the
measured path; results below reflect the implemented kernel only.
restructures the decode step so the per-tile matrix products and the
online-softmax merge are issued as \emph{composite} commands rather than
hand-tiled primitives; \S\ref{sec:gqa-composite} measures it against the
primitive baseline at long context.
\subsection{Results}
% TODO: CUBE <-> KV-head mapping diagram for the short-context regime
% (h_kv=8 KV heads -> 8 CUBEs, 1:1; intra-CUBE PE usage).
% Bench code: src/kernbench/benches/gqa_helpers/short_ctx/
We measure four headline panels that vary the user count $C$ and the phase:
single- and multi-user prefill ($T_q=4$, $S_{kv}=16$), and single- and
multi-user decode ($P=8$ PEs, $S_{kv}=64$ and $128$), all at $d_{\text{head}}=64$
and $G=8$. For each panel we harvest end-to-end latency
(max event end minus min event start, the same window convention as the
GEMM study) together with the per-engine busy time and the operation mix.
Figure~\ref{fig:gqa-lat} and Figure~\ref{fig:gqa-break} report the result;
the underlying numbers are in Table~\ref{tab:gqa}.
% TODO: prefill performance figure (latency, stage breakdown).
% TODO: decode performance figure (latency, stage breakdown).
% Bench output for short_ctx to be generated.
\begin{table}[t]
\centering
\caption{Fused GQA per-panel latency and operation mix. Compute (GEMM,
MATH) is a tiny fraction of DMA occupancy; IPCQ copies grow with users and
PEs.}
\label{tab:gqa}
\small
\begin{tabular}{@{}lrrrr@{}}
\toprule
\textbf{Panel} & \textbf{Lat.\ (ns)} & \textbf{GEMM} & \textbf{IPCQ} & \textbf{DMA rd} \\
\midrule
prefill C=1 & 445 & 2 & 0 & 3 \\
prefill C=4 (Ring) & 4630 & 32 & 24 & 12 \\
decode C=1, P=8 & 3632 & 16 & 21 & 24 \\
decode C=4, P=8 & 6693 & 64 & 93 & 96 \\
\bottomrule
\end{tabular}
\end{table}
\subsection{Inference with Long-Context Length}
\label{sec:gqa-long}
% TODO: prefill long-context kernel implementation description
% (Sequence-Parallel partition of S_kv, per-case mechanics).
% Bench code: src/kernbench/benches/gqa_helpers/long_ctx/
% TODO: prefill long-context performance figure.
Long-context decode is the regime where the KV cache, not attention
compute, sets serving cost, so the placement question of
\S\ref{sec:gqa-placement} becomes decisive here. To pick the right
placement for this regime we run each of the six options as a fused
decode kernel on the simulator (one decode step on the LLaMA-3.1-70B
single-KV-head-group target, $C{=}8$ CUBEs $\times$ $P{=}8$ PEs) at a
tractable $S_{kv}{=}8192$, and read off end-to-end latency, on-device op
traffic, and the redundant compute each one induces. The swept context is
small enough that all six fit in memory at $S_{kv}{=}8192$; the
placements the long-context memory budget rules out (Cases~1--3) are
drawn in red, run here only to expose their issue and communication
structure.
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gqa_latency_by_panel.png}
\caption{Fused GQA end-to-end latency. Latency grows from
\SI{445}{\nano\second} (single-user prefill) to \SI{6693}{\nano\second}
(four-user decode) as the KV history and the number of participating
devices grow.}
\label{fig:gqa-lat}
\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_latency.png}
\caption{Measured end-to-end decode latency per placement
($S_{kv}{=}8192$; red = ruled out by the long-context memory budget,
blue = the predicted Pareto choice). The fastest raw latency belongs to
Case~3 (\SI{17.8}{\micro\second})---but Case~3 replicates the full KV
cache into every CUBE, so it overflows the per-PE budget at production
context and wastes 8$\times$ the compute
(Figure~\ref{fig:gqa-6cases-par}). Among the placements that actually fit
1\,M-token memory (Cases~4--6), Case~6~$\star$ is the fastest
(\SI{30.6}{\micro\second}, versus \SI{31.4}{} and \SI{34.5}{\micro\second}
for the $d_{\text{head}}$-split Cases~4 and~5)---making it the placement
of choice for long-context decode.}
\label{fig:gqa-6cases-lat}
\end{figure}
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gqa_op_engine_breakdown.png}
\caption{Where the work goes. Left: operation counts---GEMM and IPCQ-copy
volume both scale with users and PEs. Right: summed engine occupancy on a
log scale---the DMA engine dominates by two to three orders of magnitude
over the GEMM and MATH engines in every panel.}
\label{fig:gqa-break}
\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_parallelism.png}
\caption{Redundant compute per placement, measured as active-PE
$\times$ $S_{\text{local}}$ (PE-tokens; lower means less wasted work).
The minimum is \num{8192} PE-tokens---one pass over the sequence.
Case~3 inflates this 8$\times$ to \num{65536} by replicating the KV
cache across all eight CUBEs so every CUBE redundantly re-attends the
whole sequence; the $d_{\text{head}}$-split Cases~4--5 likewise carry
\num{65536} because each token is processed across eight head slices.
Case~6~$\star$ achieves the full 64-way split at the minimal
\num{8192} PE-tokens---fully parallel, no replication.}
\label{fig:gqa-6cases-par}
\end{figure}
The dominant observation is in Figure~\ref{fig:gqa-break}: the compute
engines are almost idle. The GEMM engine accumulates only
\SIrange{2}{33}{\nano\second} of busy time across the panels and the
vector-math engine \SIrange{5}{688}{\nano\second}, while the DMA engine
accumulates \SIrange{72}{15920}{\nano\second}. Fused GQA, as modeled here,
is overwhelmingly data-movement bound. The operation mix shows why the
collective machinery matters: IPCQ-copy count rises from zero (single-user
prefill) to 93 (four-user decode) as the kernel reduces partial outputs
across more PEs and CUBEs, and DMA-read count rises in step as more KV
shards are streamed. The PE control-processor dispatch cost registered as
zero in this configuration---command issue is simply not on the critical
path when data movement is this dominant.
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gqa_decode_long_ctx_6cases_traffic.png}
\caption{Measured on-device op traffic per placement. The unsharded
Case~1 issues no IPCQ copies (each PE has the full cache, nothing to
reduce); the single-axis Case~3 charges 168. Among the 64-way splits,
the $d_{\text{head}}$-split Cases~4--5 charge the most---280 IPCQ copies
and 8 DMA writes each, the partial-score all-reduce that head-slicing
forces---while Case~6~$\star$ needs only 189 IPCQ copies and a single DMA
write, because it merges just the running softmax state $(m,\ell,O)$
rather than partial scores. This is the on-device collective traffic that
PE\_IPCQ and the torus links of \S\ref{sec:allreduce} are provisioned to
absorb at link speed.}
\label{fig:gqa-6cases-traffic}
\end{figure}
\subsection{Analysis and meaning}
The measurements select the right placement for this regime. The fastest
raw latency (Case~3, \SI{17.8}{\micro\second}) comes from replicating the
full KV cache into every CUBE---which is exactly the placement the
long-context memory budget forbids, and which the parallelism panel shows
wastes 8$\times$ the compute. Restricting attention to the placements that
fit production-context memory (the 64-way splits, Cases~4--6), the choice
is Case~6~$\star$: it is the fastest of the three
(\SI{30.6}{\micro\second}), and the op-count panel shows why---its
softmax-state-only reduction charges 189 IPCQ copies and one DMA write
against the 280 copies and 8 DMA writes the $d_{\text{head}}$-split
Cases~4--5 pay for their partial-score all-reduce. For long-context
decode, then, the appropriate data placement is the both-axes sequence
shard (Case~6): it is the cheapest-communicating member of the only
memory-feasible family, and the cross-PE softmax reduction it does pay is
precisely the traffic the communication-side codesign of this report is
built to move quickly.
\subsection{Use of Composite Commands}
\label{sec:gqa-composite}
The decode kernel of \S\ref{sec:gqa-long} issues its local attention as
primitive operations: it walks each PE's $S_{\text{local}}$ token slice in
\SI{1024}{}-token tiles, and for every tile issues a $Q\!\cdot\!K^{\top}$
\textsf{dot}, the online-softmax primitives, and a $P\!\cdot\!V$
\textsf{dot}, merging the running $(m,\ell,O)$ state by hand. The number
of PE\_CPU commands this costs grows with the context. At a production
context of $S_{kv}{=}1\,\text{M}$ tokens the Case-6 64-way split gives
each PE $S_{\text{local}}{=}16384$ tokens, so its local attention is a
$Q\!\cdot\!K^{\top}$ of $(8,128)\!\cdot\!(128,16384)$ and a
$P\!\cdot\!V$ of $(8,16384)\!\cdot\!(16384,128)$---sixteen hand-issued
tiles, each a fresh batch of CPU commands.
The composite command lets the kernel hand that tiling to PE\_SCHEDULER.
We compare three command forms of the \emph{same} Case-6 kernel---identical
placement and identical $(m,\ell,O)$ reduce, differing only in how the
local attention is issued:
\begin{itemize}
\item \textbf{primitive}---the hand-tiled \textsf{dot}/softmax kernel
above (the baseline of \S\ref{sec:gqa-long}).
\item \textbf{composite}---each matrix product is one coarse
\textsf{composite} GEMM over the \emph{whole} $S_{\text{local}}$, with
$K$ and $V$ passed as HBM references so PE\_SCHEDULER streams and
tiles them on the fixed $32\!\times\!64\!\times\!32$ MAC tile; the
softmax stays primitive.
\item \textbf{composite\,+\,softmax\_merge}---additionally folds the
online-softmax merge and $P\!\cdot\!V$ into a single stateful
\textsf{softmax\_merge} recipe composite.
\end{itemize}
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gqa_decode_long_ctx_composite.png}
\caption{Three command forms of the Case-6 decode kernel, swept over
context length ($S_{\text{local}}{=}S_{kv}/64$ per PE). \emph{Right:}
PE\_CPU commands issued. The hand-tiled primitive kernel rises
$O(n_{\text{tiles}})$---from 96 commands at one tile to 426 at the
1\,M-token, sixteen-tile production point---while both composite forms
issue a context-\emph{independent} $O(1)$ count (94 and 98) that
\emph{saturates}: one coarse descriptor offloads the entire per-tile
fan-out. \emph{Left:} the consequence for wall-clock latency is none---all
three land on the same curve (\SI{30.6}{}, \SI{231}{},
\SI{461}{\micro\second} at 8\,K\,/\,64\,K\,/\,128\,K), because decode is
bound by streaming the KV cache, not by issue. Command-count is measured
at emit time (exact, to 1\,M); latency on the data-mode engine over the
tractable range.}
\label{fig:gqa-composite}
\end{figure}
Figure~\ref{fig:gqa-composite} reads off the two quantities that matter,
and they point in opposite directions. The PE\_CPU command count (right)
collapses from a context-growing $O(n_{\text{tiles}})$ to a flat $O(1)$:
at 1\,M tokens the composite form issues \num{94} commands against the
primitive kernel's \num{426}, a $4.5\times$ reduction
($4.2\times$ in modeled dispatch cost), and---crucially---that number no
longer grows with context. The wall-clock latency (left), by contrast,
is unchanged across all three forms: decode is bound by streaming the KV
cache out of HBM, so the command form does not move the critical path.
That juxtaposition is the point. The composite command is not a latency
optimization for this memory-bound decode \emph{at the 64-way production scale} (a
single-rank caveat follows, Figure~\ref{fig:gqa-decode-stream}); it is a
\emph{CPU-issue}
optimization. Its value is removing the per-tile dispatch work that would
otherwise grow without bound as context grows, freeing PE\_CPU to run
ahead and keep the engines fed---which is exactly what lets the
data-movement cost analyzed next show through as the true bottleneck
rather than being masked by issue overhead. The \textsf{softmax\_merge}
recipe folds the online merge into the same descriptor; on this
memory-bound path its marginal cost over the plain GEMM composite is small
(98 vs.\ 94 commands), and like the plain composite it keeps the issued
count flat as context scales.
\paragraph{Isolating the rank: the masked streaming win.} That neutrality
is a property of the \emph{full 64-way} critical path, not of the local
attention: at production scale the inter-CUBE $(m,\ell,O)$ reduce tail and
shared-HBM contention set the wall clock, so a faster local attention does
not surface. Stripping those away---a single rank, no cross-CUBE reduce,
swept over the per-rank context $S_{kv}$ (so $S_{kv}{=}16$\,K here is the
per-PE load of a 1M-token, 64-way-sharded decode)---exposes the local
attention directly (Figure~\ref{fig:gqa-decode-stream}), and the composite
\emph{does} win, by \SI{25}{}--\SI{28}{\percent}. The reason is the
memory-bound mirror of prefill: its scheduler-streamed concurrent per-tile
DMAs keep the HBM pipeline full and reach \SI{233}{\giga\byte\per\second}
---\SI{91}{\percent} of the per-rank \SI{256}{\giga\byte\per\second}
roofline---whereas the primitive kernel's blocking \textsf{tl.dot}
serializes one tile DMA at a time and plateaus at
\SI{166}{\giga\byte\per\second}. So even for memory-bound decode the
composite is not \emph{only} a CPU-issue optimization---it also extracts
bandwidth---but that latency benefit materializes only when the local
attention is on the critical path, which at 64-way production scale it is
not.
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gqa_decode_streaming.png}
\caption{Single-rank memory-bound decode ($T_q{=}1$, $M{=}8$), three
command forms, swept over per-rank context. \emph{Left:} end-to-end
latency---the composite forms run \SI{25}{}--\SI{28}{\percent} below the
primitive, a gap that widens with context. \emph{Right:} achieved HBM
bandwidth against the per-rank \SI{256}{\giga\byte\per\second} roofline.
The primitive's blocking load$\rightarrow$dot serializes the KV stream and
plateaus at \SI{166}{\giga\byte\per\second}; the composite forms pipeline
concurrent per-tile DMAs through the scheduler and reach
\SI{233}{\giga\byte\per\second}. This is the memory-bound mirror of the
prefill result (Figure~\ref{fig:gqa-prefill-cb}): there the composite
approaches the MAC roofline, here the bandwidth roofline. Capped at
$16$\,K---the plain composite materializes the full $(M,S_{kv})$ scores in
TCM, so beyond that only the \textsf{softmax\_merge} recipe, which tiles
the softmax, stays within scratch.}
\label{fig:gqa-decode-stream}
\end{figure}
\paragraph{The compute-bound mirror: prefill.} Decode's \emph{production}
verdict---command form is latency-neutral at 64-way scale---is a property
of its regime, not of the
composite command. A decode step has $T_q{=}1$, so its score and context
products are skinny ($M{=}G\,T_q{=}8$): the MAC array is barely fed and
the kernel is bound by streaming the KV cache. Prefill is the opposite
corner. It processes a block of query positions at once, so $M{=}G\,T_q$
is large and tile-filling, the GEMMs carry real arithmetic intensity
($\sim$$M$ flops/byte, well above the roofline ridge), and the kernel is
\emph{compute-bound}. This is the regime the composite command was built
for (\S\ref{sec:gemm}): it streams the per-HW-tile
DMA$\rightleftarrows$compute pipeline so the MAC array stays fed, whereas
the primitive kernel's blocking \textsf{tl.dot} serializes each tile's
load and compute and starves the array between tiles. We run the same
three command forms on a single-rank compute-bound prefill (FlashAttention
$Q$-block $\times$ $S_{kv}$-tile, online softmax) and sweep the context
length (Figure~\ref{fig:gqa-prefill-cb}).
\begin{figure}[t]
\centering
\includegraphics[width=\linewidth]{gqa_prefill_compute_bound.png}
\caption{Compute-bound prefill, three command forms, swept over context
length ($M{=}8\,T_q$ tile-filling). \emph{Left:} end-to-end latency.
\emph{Right:} MAC utilization (achieved $\div$ the
\SI{8}{\tera\flop\per\second} per-PE peak). The hand-tiled primitive sits
flat at $\sim$\SI{68}{\percent}---its serial load$\to$dot path leaves the
MAC array idle between tiles regardless of context. The composite forms
climb with context (\SI{67}{}$\to$\SI{80}{\percent} plain,
\SI{67}{}$\to$\SI{83}{\percent} with the recipe) because a deeper $P\!\cdot
\!V$ reduction gives more HW tiles to pipeline, and they convert that into
wall-clock: at \num{1024} the recipe form is \SI{646.9}{} vs.\
\SI{794.1}{\micro\second} (\SI{19}{\percent} faster). The margin
\emph{grows} with context---the compute-bound mirror of the GEMM result of
\S\ref{sec:gemm}.}
\label{fig:gqa-prefill-cb}
\end{figure}
The two studies together state the composite command's value precisely. It
has two distinct benefits, and which one matters is set by the workload's
roofline position. The first is \emph{host-issue offload}: one macro
command in place of $O(n_{\text{tiles}})$ fine ones, which removes
PE\_CPU dispatch work and is regime-independent (it shows in the decode
command count). The second is \emph{MAC-array feeding}: the
scheduler-internal per-tile DMA$\rightleftarrows$compute pipeline, which
only converts to latency when the workload is compute-bound enough to have
a MAC array worth keeping busy (it shows in the prefill utilization). A
memory-bound decode exercises only the first; a compute-bound prefill
exercises both. The composite command is the single mechanism that
delivers each where it applies.
\subsection{Comprehensive Analysis}
\label{sec:gqa-analysis}
These panels are the clearest statement of the codesign thesis in the
report. Because the composite command keeps GEMM issue cheap and the MAC
@@ -121,3 +405,9 @@ optimization is what \emph{attacks} it. For an attention-dominated decoder
the meaningful hardware investments are therefore the ones that move data
faster and reduce it on-device---not additional MAC throughput, which this
workload cannot use.
% TODO: cross-regime DP (data parallelism) applicability:
% - Does Case-4 long-context placement compose with batch-level DP
% without further changes?
% - Does the short-context placement compose the same way?
% - Implications for multi-user serving (single vs. mixed regimes).
docs/report/1H-codesign-paper/sections/06-discussion.tex
+4 -1
View File
@@ -40,7 +40,10 @@ hardware levers for this workload:
GQA panels leave the GEMM and vector-math engines two to three orders of
magnitude below the DMA engine in busy time. Adding MAC area would not move
decode latency; the workload cannot use it. This is the single most
actionable finding for an attention-dominated roadmap.
actionable finding for an attention-dominated roadmap. The implication is
that future hardware investment should prioritize communication and
memory-system efficiency over additional compute throughput for
attention-dominated inference workloads.
\paragraph{Caveats.} These conclusions are achievable-kernel results from a
deterministic model, not E2E measurements; absolute numbers carry the
scratch_dot_tiling_experiment.py
+204
View File
@@ -0,0 +1,204 @@
#!/usr/bin/env python3
"""SCRATCH EXPERIMENT (not production; do not commit).
Question: does charging the *primitive* decode kernel for per-HW-tile
(16x16x16) CPU dispatch flip the "composite gives no decode-latency
benefit" conclusion?
We monkeypatch TLContext.dot so that, in the primitive kernel, every
tl.dot whose (M,K,N) exceeds the HW GEMM tile (mac_m/mac_k/mac_n) is
split by the CPU into ceil(M/mac_m)*ceil(K/mac_k)*ceil(N/mac_n)
HW-tile-sized GemmCmds. Each tile GemmCmd is emitted through the normal
_emit() path, so it (a) charges PE_CPU dispatch overhead via
_charge_dispatch (PeCpuOverheadCmd), and (b) blocks like a normal
single-op GemmCmd on PE_GEMM at the cycle-accurate ceil-product latency.
We also inject mac_m/mac_k/mac_n into every pe_gemm topology node so BOTH
the primitive-tiled and the composite variants run on the *same*
cycle-accurate engine (fair comparison). Composite is left untouched:
the CPU emits ONE CompositeCmd, and PE_SCHEDULER tiles internally (no
per-HW-tile CPU dispatch).
Data correctness:
Inputs are ctx.zeros (q/k/v), so every matmul result is zeros and the
DataExecutor replay is trivial. To keep replay numerically correct
regardless, exactly ONE emitted tile per dot carries the *real* full
operands+output handles (so the DataExecutor computes the true (M,N)
result via the recorded handle shapes), while its timing fields
(m,k,n) are the HW-tile size so the engine charges exactly one tile of
cycle time. The remaining n_tiles-1 emitted tiles are timing-only
GemmCmds (16x16x16) writing to throwaway scratch. Net: n_tiles tiles
of engine time + n_tiles dispatch charges, and a correct final output.
Engine mode: enable_data=True (same as the production sweep's
_engine_latency_ns), op_log end-to-end latency.
"""
from __future__ import annotations
import json
from math import ceil
from pathlib import Path
# --- HW GEMM tile under test -------------------------------------------------
MAC_M = 16
MAC_K = 16
MAC_N = 16
# Legacy alt to also try: (8, 16, 32)
S_KV_LATENCY = (8192, 32_768, 65_536, 131_072)
ROOT = Path(__file__).resolve().parent
SWEEP_JSON = (
ROOT / "src" / "kernbench" / "benches" / "1H_milestone_output"
/ "gqa" / "long_ctx" / "sweep_decode_composite.json"
)
# --- mac-dim topology override -----------------------------------------------
def _topo_with_mac(mac_m: int, mac_k: int, mac_n: int):
"""Compiled topology with mac dims injected into every pe_gemm node."""
from kernbench.topology.builder import resolve_topology
handle = resolve_topology("topology.yaml")
g = handle.topology_obj
n = 0
for node in g.nodes.values():
if node.kind == "pe_gemm":
node.attrs["mac_m"] = mac_m
node.attrs["mac_k"] = mac_k
node.attrs["mac_n"] = mac_n
n += 1
print(f" injected mac=({mac_m},{mac_k},{mac_n}) into {n} pe_gemm nodes")
return handle
# --- tiling monkeypatch for TLContext.dot ------------------------------------
def _make_tiled_dot(orig_dot, mac_m: int, mac_k: int, mac_n: int):
from kernbench.common.pe_commands import GemmCmd
def tiled_dot(self, a, b):
if len(a.shape) < 2 or len(b.shape) < 2:
return orig_dot(self, a, b)
m, k = a.shape[-2], a.shape[-1]
k2, n = b.shape[-2], b.shape[-1]
if k != k2:
raise ValueError(f"dot shape mismatch: a.K={k} != b.K={k2}")
n_tiles = ceil(m / mac_m) * ceil(k / mac_k) * ceil(n / mac_n)
out_shape = (*a.shape[:-2], m, n)
out = self._make_compute_out(shape=out_shape, dtype=a.dtype)
self._await_pending(a, b)
if n_tiles <= 1:
self._emit(GemmCmd(a=a, b=b, out=out, m=m, k=k, n=n))
return out
# One real-data tile: full handles (so DataExecutor computes the
# true result), but timing fields = HW tile (one tile of cycles).
self._emit(GemmCmd(a=a, b=b, out=out, m=mac_m, k=mac_k, n=mac_n))
# Remaining timing-only tiles: throwaway scratch, 16x16x16.
scratch = self._make_compute_out(shape=(mac_m, mac_n), dtype=a.dtype)
for _ in range(n_tiles - 1):
self._emit(GemmCmd(a=a, b=b, out=scratch,
m=mac_m, k=mac_k, n=mac_n))
return out
return tiled_dot
# --- latency runner (replicates sweep's _engine_latency_ns) ------------------
def _engine_latency_ns(variant: str, S_kv: int, topo) -> float:
from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_composite import ( # noqa: E501
_end_to_end_ns, _run_panel_fn,
)
from kernbench.runtime_api.bench_runner import run_bench
from kernbench.runtime_api.types import resolve_device
from kernbench.sim_engine.engine import GraphEngine
result = run_bench(
topology=topo, bench_fn=_run_panel_fn(variant, S_kv),
device=resolve_device(None),
engine_factory=lambda t, d: GraphEngine(
getattr(t, "topology_obj", t), enable_data=True,
),
)
if not result.completion.ok:
raise RuntimeError(
f"{variant}@{S_kv} failed: {result.completion}"
)
return _end_to_end_ns(result.engine.op_log)
def _emit_dispatch(variant: str, S_kv: int) -> int:
"""PE_CPU command count at the center rank (cube 6, pe 0)."""
from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_composite import ( # noqa: E501
_emit_dispatch as prod_emit,
)
return prod_emit(variant, S_kv)[0]
def main() -> None:
import kernbench.triton_emu.tl_context as tlc
# Baseline (A): primitive UNTILED latencies from the production sweep
# (mac=0 / TFLOPS model). Read straight off the committed sweep JSON.
sweep = json.loads(SWEEP_JSON.read_text())
base_A = {}
for r in sweep["rows"]:
if r["variant"] == "primitive" and r["latency_ns"] is not None:
base_A[r["S_kv"]] = r["latency_ns"]
print(f"== mac tile = ({MAC_M},{MAC_K},{MAC_N}) ==")
# --- command-count sanity (emit-time, mac-independent) ---------------
orig_dot = tlc.TLContext.dot
print("\n[dispatch counts @ S_kv=131072]")
n_prim_untiled = _emit_dispatch("primitive", 131072)
tlc.TLContext.dot = _make_tiled_dot(orig_dot, MAC_M, MAC_K, MAC_N)
try:
n_prim_tiled = _emit_dispatch("primitive", 131072)
n_comp = None
finally:
tlc.TLContext.dot = orig_dot
n_comp = _emit_dispatch("composite", 131072)
print(f" primitive UNTILED PE_CPU cmds : {n_prim_untiled}")
print(f" primitive TILED PE_CPU cmds : {n_prim_tiled} "
f"(x{n_prim_tiled / max(n_prim_untiled,1):.0f})")
print(f" composite PE_CPU cmds : {n_comp}")
# --- latency sweep ----------------------------------------------------
rows = []
topo = _topo_with_mac(MAC_M, MAC_K, MAC_N)
for S_kv in S_KV_LATENCY:
A = base_A.get(S_kv)
# (C) composite on the mac engine (untouched dot path)
C = _engine_latency_ns("composite", S_kv, topo)
# (B) primitive TILED on the mac engine
tlc.TLContext.dot = _make_tiled_dot(orig_dot, MAC_M, MAC_K, MAC_N)
try:
B = _engine_latency_ns("primitive", S_kv, topo)
finally:
tlc.TLContext.dot = orig_dot
gap_pct = (B - C) / C * 100.0 if C else float("nan")
rows.append((S_kv, A, B, C, gap_pct))
print(f" S_kv={S_kv:>7}: A(untiled)={A!s:>12} "
f"B(tiled)={B:12.2f} C(comp)={C:12.2f} (B-C)/C={gap_pct:+6.1f}%")
# --- final table ------------------------------------------------------
print("\n==================== RESULT TABLE ====================")
print(f"{'S_kv':>8} | {'A untiled(ns)':>14} | {'B tiled(ns)':>14} | "
f"{'C comp(ns)':>14} | {'(B-C)/C':>9}")
print("-" * 72)
for S_kv, A, B, C, gap in rows:
a_s = f"{A:.2f}" if A is not None else "n/a"
print(f"{S_kv:>8} | {a_s:>14} | {B:>14.2f} | {C:>14.2f} | "
f"{gap:>+8.1f}%")
if __name__ == "__main__":
main()
scripts/paper/measure_gqa_decode_placement_comm.py
+362
View File
@@ -0,0 +1,362 @@
"""Measured comm overlay for all 6 GQA decode KV placements.
Runs the simulator for Cases 1-6 (the same 6 placements the chart in
paper_plot_gqa_4cases_summary.py covers analytically) at S_kv = 64 K,
sums actual IPCQ-copy bytes from the engine op_log, projects to
per-token (x80 layers), scales the partial-score-AR component of
Cases 4/5 from S_kv = 64 K -> S_kv = 1 M (linear in S_kv; other
cases are S_kv-independent), adds the constant Wo + FFN AR
(1.25 MB / token), and writes the result to JSON for
paper_plot_gqa_4cases_summary.py to overlay on the analytical bars.
Single layer of decode attention only — the projection × 80 takes
the per-layer measurement to a per-token total.
Usage:
python scripts/paper/measure_gqa_decode_placement_comm.py
"""
from __future__ import annotations
import json
import os
import sys
from pathlib import Path
from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_repl_pe_sp import (
gqa_attention_decode_long_ctx_cube_repl_pe_sp_kernel as _case3_kernel,
)
from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_repl_pe_tp import (
gqa_attention_decode_long_ctx_cube_repl_pe_tp_kernel as _case1_kernel,
)
from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_sp_pe_sp import (
gqa_attention_decode_long_ctx_cube_sp_pe_sp_kernel as _case6_kernel,
)
from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_sp_pe_tp import (
gqa_attention_decode_long_ctx_cube_sp_pe_tp_kernel as _case2_kernel,
)
from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_sp_pe_tp_dhead import (
gqa_attention_decode_long_ctx_cube_sp_pe_tp_dhead_kernel as _case4_kernel,
)
from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_tp_dhead_pe_sp import (
gqa_attention_decode_long_ctx_cube_tp_dhead_pe_sp_kernel as _case5_kernel,
)
from kernbench.benches.gqa_helpers.shared._gqa_panel_helpers import _ccl_cfg
from kernbench.ccl.sfr_config import configure_sfr_intercube_multisip
from kernbench.policy.placement.dp import DPPolicy
from kernbench.runtime_api.bench_runner import run_bench
from kernbench.runtime_api.types import resolve_device
from kernbench.sim_engine.engine import GraphEngine
from kernbench.topology.builder import resolve_topology
_C = 8
_P = 8
_N_LAYERS = 80
_S_KV_MEAS = 64 * 1024 # per-run simulator S_kv (1/16th of headline)
_S_KV_HEADLINE = 1 << 20 # 1 Mi tokens, the chart's headline S_kv
# Per-cube S_kv share for the d_head-TP partial-score AR cost.
# Case 4 (Cube-SP × PE-TP_dhead): per-cube = S_kv / C
# Case 5 (Cube-TP_dhead × PE-SP): per-cube = S_kv (KV replicated across
# cubes for the cube-axis d_head-TP), so partial-score AR scales
# with the full S_kv.
# Headline / measured per-cube ratios give the partial-score-AR scale-up
# from S_kv = 64 K to S_kv = 1 M. (m,ℓ,O) AR is S_kv-independent.
_PARTIAL_SCORE_SCALE = _S_KV_HEADLINE / _S_KV_MEAS # = 16
# Per-token Wo + FFN AR (constant across all cases, comes from the
# attn-output and FFN-down all-reduces NOT measured by the attention-
# only kernel run here).
_WO_PER_LAYER_BYTES = 8 * 1024
_FFN_PER_LAYER_BYTES = 8 * 1024
_WO_FFN_PER_TOKEN_BYTES = (
(_WO_PER_LAYER_BYTES + _FFN_PER_LAYER_BYTES) * _N_LAYERS
) # 1.25 MB
# Total PE count in one KV-head group — average per-PE comm = total / N.
_NUM_PES = _C * _P
# Total partial-score-AR slice produced by the attention compute when
# d_head is sharded. Used to split measured IPCQ traffic into the
# S_kv-scaling component (partial scores) vs the S_kv-independent
# component ((m,ℓ,O) merge). Same as the analytical formula in
# paper_plot_gqa_4cases_summary.py: h_q · S_q · per_cube_S_kv · 2 bytes.
_H_Q = 8
_S_Q = 1
_BYTES_PER_ELEM = 2
_PARAMS = dict(C=_C, P=_P, T_q=_S_Q, S_kv=_S_KV_MEAS,
d_head=128, h_q=_H_Q, h_kv=1)
def _bench_fn_case1(ctx):
"""Case 1: Cube-Repl x PE-repl (PE-TP doesn't shard KV)."""
p = _PARAMS
configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
dp = DPPolicy(cube="replicate", pe="replicate",
num_cubes=p["C"], num_pes=p["P"])
q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp, name="q_c1")
k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp, name="k_c1")
v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp, name="v_c1")
o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp, name="o_c1")
ctx.launch("case1_repl_repl", _case1_kernel,
q, k, v, o,
p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
p["d_head"], p["C"], p["P"],
_auto_dim_remap=False)
def _bench_fn_case2(ctx):
"""Case 2: Cube-SP x PE-repl (PE-TP doesn't shard KV)."""
p = _PARAMS
configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
dp_full = DPPolicy(cube="replicate", pe="replicate",
num_cubes=p["C"], num_pes=p["P"])
dp_kv = DPPolicy(cube="row_wise", pe="replicate",
num_cubes=p["C"], num_pes=p["P"])
q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="q_c2")
k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="k_c2")
v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="v_c2")
o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="o_c2")
ctx.launch("case2_sp_repl", _case2_kernel,
q, k, v, o,
p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
p["d_head"], p["C"], p["P"],
_auto_dim_remap=False)
def _bench_fn_case3(ctx):
"""Case 3: Cube-Repl x PE-SP."""
p = _PARAMS
configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
dp_full = DPPolicy(cube="replicate", pe="replicate",
num_cubes=p["C"], num_pes=p["P"])
dp_kv = DPPolicy(cube="replicate", pe="row_wise",
num_cubes=p["C"], num_pes=p["P"])
q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="q_c3")
k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="k_c3")
v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="v_c3")
o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="o_c3")
ctx.launch("case3_repl_sp", _case3_kernel,
q, k, v, o,
p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
p["d_head"], p["C"], p["P"],
_auto_dim_remap=False)
def _bench_fn_case4(ctx):
p = _PARAMS
configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
dp_full = DPPolicy(cube="replicate", pe="column_wise",
num_cubes=p["C"], num_pes=p["P"])
dp_kv = DPPolicy(cube="row_wise", pe="column_wise",
num_cubes=p["C"], num_pes=p["P"])
q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="q_c4")
k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="k_c4")
v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="v_c4")
o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="o_c4")
ctx.launch("case4_dhead_tp", _case4_kernel,
q, k, v, o,
p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
p["d_head"], p["C"], p["P"],
_auto_dim_remap=False)
def _bench_fn_case5(ctx):
p = _PARAMS
configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
dp_q = DPPolicy(cube="column_wise", pe="replicate",
num_cubes=p["C"], num_pes=p["P"])
dp_kv = DPPolicy(cube="column_wise", pe="row_wise",
num_cubes=p["C"], num_pes=p["P"])
q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_q, name="q_c5")
k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="k_c5")
v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="v_c5")
o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_q, name="o_c5")
ctx.launch("case5_dhead_tp_inter", _case5_kernel,
q, k, v, o,
p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
p["d_head"], p["C"], p["P"],
_auto_dim_remap=False)
def _bench_fn_case6(ctx):
p = _PARAMS
configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
dp_full = DPPolicy(cube="replicate", pe="replicate",
num_cubes=p["C"], num_pes=p["P"])
dp_kv = DPPolicy(cube="row_wise", pe="row_wise",
num_cubes=p["C"], num_pes=p["P"])
q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="q_c6")
k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="k_c6")
v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="v_c6")
o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="o_c6")
ctx.launch("case6_sp_sp", _case6_kernel,
q, k, v, o,
p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
p["d_head"], p["C"], p["P"],
_auto_dim_remap=False)
def _sum_ipcq_bytes(op_log) -> int:
"""Sum nbytes across all ipcq_copy records."""
return sum(
r.params.get("nbytes", 0)
for r in op_log
if r.op_kind == "memory" and r.op_name == "ipcq_copy"
)
def _partial_score_slices(case: int) -> int:
"""Divisor that splits S_kv into partial-score tiles, per the
analytical model in paper_plot_gqa_4cases_summary.py:
partial_score_per_PE = h_q * S_q * (s_kv / slices) * bytes
Case 4 (Cube-SP x PE-TP-dhead): intra-cube AR over d_head shards
on PE axis -> partial tile per PE has per-cube S_kv = s_kv/C.
Case 5 (Cube-TP-dhead x PE-SP): inter-cube AR over d_head shards
on cube axis -> partial tile per PE has per-PE S_kv = s_kv/P.
Cases 1, 2, 3, 6: no partial-score AR (only (m,l,O) merge).
"""
if case == 4:
return _C
if case == 5:
return _P
return 0
def _split_attn_layer_bytes(case: int, total_ipcq_bytes: int,
s_kv: int) -> tuple[int, int]:
"""Split per-layer attention-time IPCQ bytes into:
(partial_score_component, mlo_component).
The partial-score component scales with s_kv (so it must be scaled
when projecting from the measure-time s_kv to the headline s_kv);
the (m,l,O) component is constant in s_kv.
Partial-score size is analytically known per-case (formula in
_partial_score_slices); the remainder is treated as (m,l,O) + any
other S_kv-independent overhead. Per-PE = total / NUM_PES.
"""
per_pe_total = total_ipcq_bytes // _NUM_PES
slices = _partial_score_slices(case)
if slices == 0:
# No partial-score AR for this case.
return 0, per_pe_total
partial_score_per_pe = (
_H_Q * _S_Q * (s_kv // slices) * _BYTES_PER_ELEM
)
partial_score_per_pe = min(partial_score_per_pe, per_pe_total)
mlo_per_pe = per_pe_total - partial_score_per_pe
return partial_score_per_pe, mlo_per_pe
_KERNELS = (
(1, "Case 1 (Cube-Repl x PE-repl)", _bench_fn_case1),
(2, "Case 2 (Cube-SP x PE-repl)", _bench_fn_case2),
(3, "Case 3 (Cube-Repl x PE-SP)", _bench_fn_case3),
(4, "Case 4 (Cube-SP x PE-TP d_head)", _bench_fn_case4),
(5, "Case 5 (Cube-TP d_head x PE-SP)", _bench_fn_case5),
(6, "Case 6 (Cube-SP x PE-SP) [*]", _bench_fn_case6),
)
def main() -> int:
topology = os.environ.get("GQA_1H_TOPOLOGY", "topology.yaml")
topo = resolve_topology(topology)
out: dict = {
"S_kv_measured": _S_KV_MEAS,
"S_kv_headline": _S_KV_HEADLINE,
"n_layers": _N_LAYERS,
"num_pes": _NUM_PES,
"wo_ffn_per_token_bytes": _WO_FFN_PER_TOKEN_BYTES,
"cases": {},
}
print(f"Measuring at S_kv={_S_KV_MEAS:,} ; scaling partial-score AR "
f"to S_kv={_S_KV_HEADLINE:,} (×{int(_PARTIAL_SCORE_SCALE)})")
print()
for case_id, label, bench_fn in _KERNELS:
try:
res = run_bench(
topology=topo, bench_fn=bench_fn,
device=resolve_device(None),
engine_factory=lambda t, d: GraphEngine(
getattr(t, "topology_obj", t), enable_data=True,
),
)
except Exception as e:
print(f" {label:<42} FAIL: {type(e).__name__}: {e}")
return 1
if not res.completion.ok:
print(f" {label:<42} ENGINE FAIL: {res.completion}")
return 1
total_ipcq = _sum_ipcq_bytes(res.engine.op_log)
partial_pe, mlo_pe = _split_attn_layer_bytes(
case_id, total_ipcq, _S_KV_MEAS,
)
# Per-token attention-time comm at S_kv = 1 M:
# (partial_score_per_layer × scale + mlo_per_layer) × 80 layers
scaled_partial_per_token = (
partial_pe * int(_PARTIAL_SCORE_SCALE) * _N_LAYERS
)
mlo_per_token = mlo_pe * _N_LAYERS
attn_per_token = scaled_partial_per_token + mlo_per_token
total_per_token = attn_per_token + _WO_FFN_PER_TOKEN_BYTES
out["cases"][str(case_id)] = {
"label": label,
"total_ipcq_bytes_one_layer": total_ipcq,
"per_pe_partial_score_bytes_one_layer": partial_pe,
"per_pe_mlo_bytes_one_layer": mlo_pe,
"per_pe_attn_bytes_per_token_at_1M": attn_per_token,
"per_pe_total_bytes_per_token_at_1M": total_per_token,
}
print(f" {label:<42} "
f"ipcq_total={total_ipcq:>10,} "
f"per_pe_attn(1L)={(partial_pe + mlo_pe):>9,} "
f"per_pe_total/tok@1M={total_per_token / (1<<20):>7.2f} MB")
out_path = (
Path(__file__).resolve().parents[2]
/ "src" / "kernbench" / "benches"
/ "1H_milestone_output" / "gqa" / "long_ctx"
/ "gqa_long_ctx_6cases_measured_comm.json"
)
out_path.parent.mkdir(parents=True, exist_ok=True)
out_path.write_text(json.dumps(out, indent=2))
print()
print(f"wrote {out_path}")
return 0
if __name__ == "__main__":
sys.exit(main())
scripts/paper/paper_cube_architecture.py
+155
View File
@@ -0,0 +1,155 @@
"""Re-emit cube_view.svg in an academic (white-background, large-font)
palette and convert it to PDF for the 1H-codesign-paper Figure 2.
Source of truth: docs/diagrams/cube_view.svg (generated by
src/kernbench/topology/visualizer.py:_render_cube_view_svg, dark theme).
This script does a targeted color/font/size remap on the dark palette so
the resulting figure prints well on a white paper page. If the upstream
palette or geometry in visualizer.py changes, the maps below must be
reviewed.
"""
from __future__ import annotations
import re
import shutil
import subprocess
from pathlib import Path
REPO = Path(__file__).resolve().parents[2]
SRC_SVG = REPO / "docs" / "diagrams" / "cube_view.svg"
OUT_DIR = REPO / "docs" / "report" / "1H-codesign-paper" / "figures"
OUT_SVG = OUT_DIR / "cube_architecture.svg"
OUT_PDF = OUT_DIR / "cube_architecture.pdf"
# ── 1. Legend-icon safety: legend rects whose general color rule below
# would turn them white-on-white. Apply BEFORE the wholesale fill rules.
LEGEND_FIXUP: list[tuple[str, str]] = [
# "Relay" legend icon (slate-700 router)
('fill="#334155" stroke="#475569" stroke-width="0.5"',
'fill="#94a3b8" stroke="#475569" stroke-width="0.5"'),
# "Mesh Link" legend icon (slate-600)
('fill="#475569" stroke="#475569" stroke-width="0.5"',
'fill="#94a3b8" stroke="#475569" stroke-width="0.5"'),
]
# ── 2. Color remap: dark theme -> academic (white) theme
COLOR_MAP: list[tuple[str, str]] = [
# page background slate-900 -> white
('fill="#0f172a"', 'fill="#ffffff"'),
# title text slate-400 -> slate-800
('fill="#94a3b8"', 'fill="#1f2937"'),
# subtitle text slate-500 -> slate-600
('fill="#64748b"', 'fill="#475569"'),
# router has-attach stroke slate-500 -> slate-600
('stroke="#64748b"', 'stroke="#475569"'),
# router has-attach fill slate-600 -> white
('fill="#475569"', 'fill="#ffffff"'),
# mesh lines + cube boundary stroke slate-600 -> slate-400
('stroke="#475569"', 'stroke="#94a3b8"'),
# router no-attach fill slate-700 -> white
('fill="#334155"', 'fill="#ffffff"'),
# component block dark fills -> white (PE / M_CPU / SRAM / UCIe)
('fill="#2d1f3d"', 'fill="#ffffff"'),
('fill="#451a03"', 'fill="#ffffff"'),
('fill="#1c1917"', 'fill="#ffffff"'),
('fill="#1e1b4b"', 'fill="#ffffff"'),
# HBM zone background emerald-950 -> emerald-50
('fill="#052e16"', 'fill="#ecfdf5"'),
# router label text white -> slate-800
('fill="white"', 'fill="#1f2937"'),
# boost low-alpha emerald annotations so PE/HBM BW labels read on light bg
('fill="#10b98188"', 'fill="#047857"'),
('fill="#05966988"', 'fill="#059669"'),
]
# ── 3. UCIe palette desaturation: bright violet/indigo/fuchsia is too
# attention-grabbing on a white page; remap to a slate gradient.
# NOTE: side-effect on PE2/PE3 HBM port-bar colors (which share
# #8b5cf6/#a78bfa with UCIe) — PE labels still convey identity.
# Applied AFTER the academic color map so the new slate values are
# not picked up by the rules above.
UCIE_DESATURATE: list[tuple[str, str]] = [
# UCIe block stroke/text (violet-500) -> slate-600
('"#8b5cf6"', '"#475569"'),
# UCIe cell 1 / PE3 port bar (violet-400) -> slate-400
('"#a78bfa"', '"#94a3b8"'),
# UCIe cell 0 (indigo-400) -> slate-300
('"#818cf8"', '"#cbd5e1"'),
# UCIe cell 2 (purple-400) -> slate-500
('"#c084fc"', '"#64748b"'),
# UCIe cell 3 (fuchsia-400) -> gray-700
('"#e879f9"', '"#374151"'),
]
# ── 4. Font-size bumps. The CUBE figure is rendered at half-text-width
# (~250pt) inside the side-by-side subfigure in 02-platform.tex, so
# native fonts get crushed ~3x by \linewidth scaling. We push the
# bumps to the legibility limit of the layout (router-label text
# stays inside a slightly enlarged circle; legend items may touch).
FONT_MAP: dict[str, str] = {
"5": "10",
"6": "12",
"7": "14",
"8": "13", # legend rect text — capped by upstream layout spacing
# (advance = 7*len(label)+24 was sized for ~font 8);
# font 13 keeps each item's text inside its slot.
"9": "14",
"10": "15",
"11": "16",
"14": "18", # title — kept moderate so it does not overflow canvas
}
# ── 5. Router circle radius bump (only circles use r="8"). Enlarged so
# the bumped router labels stay inside the circle.
RADIUS_MAP: list[tuple[str, str]] = [
(' r="8"', ' r="17"'),
]
# ── 6. Tighten whitespace: move legend just below dashed box and crop
# the unused canvas margins so LaTeX's \linewidth scaling does not
# shrink the figure text any more than necessary.
LAYOUT_FIXUP: list[tuple[str, str]] = [
# Move legend rects up (dashed box bottom is at y=760)
(' y="865"', ' y="775"'),
# Move legend text baselines up to match (offset = 9 px)
(' y="874"', ' y="784"'),
# Tight crop: 5 px margin around dashed box + cut top/bottom whitespace
('<svg xmlns="http://www.w3.org/2000/svg" width="970" height="900" '
'viewBox="0 0 970 900">',
'<svg xmlns="http://www.w3.org/2000/svg" width="860" height="798" '
'viewBox="55 2 860 798">'),
]
def _bump_font(m: re.Match) -> str:
return f'font-size="{FONT_MAP.get(m.group(1), m.group(1))}"'
def main() -> None:
if not SRC_SVG.exists():
raise SystemExit(f"source SVG missing: {SRC_SVG}")
rsvg = shutil.which("rsvg-convert")
if rsvg is None:
raise SystemExit("rsvg-convert not found (brew install librsvg)")
svg = SRC_SVG.read_text(encoding="utf-8")
for old, new in (LEGEND_FIXUP + COLOR_MAP + UCIE_DESATURATE
+ RADIUS_MAP + LAYOUT_FIXUP):
if old not in svg:
print(f"warn: pattern not present in source SVG: {old}")
svg = svg.replace(old, new)
svg = re.sub(r'font-size="(\d+)"', _bump_font, svg)
OUT_DIR.mkdir(parents=True, exist_ok=True)
OUT_SVG.write_text(svg, encoding="utf-8")
subprocess.run(
[rsvg, "-f", "pdf", "-o", str(OUT_PDF), str(OUT_SVG)],
check=True,
)
print(f"wrote {OUT_SVG.relative_to(REPO)}")
print(f"wrote {OUT_PDF.relative_to(REPO)}")
if __name__ == "__main__":
main()
scripts/paper/paper_gqa_latency.py
-101
View File
@@ -1,101 +0,0 @@
"""Report harness: GQA end-to-end latency + op/engine breakdown.
Isolated under ``scripts/paper/`` for the 1H codesign report only — it is
NOT a registered bench and does not touch other people's benches. It
reuses the existing GQA headline panels (the real GQA kernels wired in
``milestone_gqa_headline``) but, unlike that milestone (which records only
op-counts), it also harvests per-panel end-to-end latency and per-engine
occupancy from ``result.engine.op_log``.
Latency definition (same window convention as ``milestone_1h_gemm``):
end_to_end_ns = max(r.t_end) - min(r.t_start) over all op_log records.
Output: docs/report/1H-codesign-paper/figures/gqa_latency.json
Run:
python scripts/paper/paper_gqa_latency.py
"""
from __future__ import annotations
import json
from pathlib import Path
from kernbench.benches.milestone_gqa_headline import (
_PANEL_DISPATCH,
_PANELS,
_make_bench_fn,
_summarize_op_log,
)
_REPORT_DIR = Path(__file__).resolve().parents[2] / "docs" / "report" / "1H-codesign-paper"
_FIG_DIR = _REPORT_DIR / "figures"
_OUT_JSON = _FIG_DIR / "gqa_latency.json"
# PE engine component suffixes whose occupancy we break out.
_ENGINES = ("pe_gemm", "pe_math", "pe_dma", "pe_fetch_store", "pe_ipcq", "pe_cpu")
def _occupancy_ns(op_log, suffix: str) -> float:
return sum(
r.t_end - r.t_start
for r in op_log
if r.component_id.endswith("." + suffix)
)
def _end_to_end_ns(op_log) -> float:
if not op_log:
return 0.0
return max(r.t_end for r in op_log) - min(r.t_start for r in op_log)
def _run_panel(panel: str, topology: str) -> dict:
from kernbench.runtime_api.bench_runner import run_bench
from kernbench.runtime_api.types import resolve_device
from kernbench.sim_engine.engine import GraphEngine
from kernbench.topology.builder import resolve_topology
topo = resolve_topology(topology)
result = run_bench(
topology=topo,
bench_fn=_make_bench_fn(panel),
device=resolve_device(None),
engine_factory=lambda t, d: GraphEngine(
getattr(t, "topology_obj", t), enable_data=True,
),
)
if not result.completion.ok:
raise RuntimeError(f"panel {panel!r} failed: {result.completion}")
op_log = result.engine.op_log
kind, params = _PANEL_DISPATCH[panel]
return {
"panel": panel,
"kind": kind,
**params,
"latency_ns": _end_to_end_ns(op_log),
"op_log_summary": _summarize_op_log(op_log),
"engine_occupancy_ns": {
eng: _occupancy_ns(op_log, eng) for eng in _ENGINES
},
}
def main() -> None:
topology = "topology.yaml"
rows = [_run_panel(panel, topology) for panel in _PANELS]
_FIG_DIR.mkdir(parents=True, exist_ok=True)
out = {"version": 1, "panels": list(_PANELS), "rows": rows}
_OUT_JSON.write_text(json.dumps(out, indent=2))
print(f"wrote {_OUT_JSON}")
for r in rows:
s = r["op_log_summary"]
print(
f" {r['panel']:24s} latency={r['latency_ns']:10.1f} ns "
f"gemm={s['gemm_count']:3d} ipcq={s['ipcq_copy_count']:3d} "
f"dma_rd={s['dma_read_count']:3d} dma_wr={s['dma_write_count']:2d}"
)
if __name__ == "__main__":
main()
scripts/paper/paper_latency_model_diagram.py
+535
View File
@@ -0,0 +1,535 @@
#!/usr/bin/env python3
"""Conceptual latency-model diagram for the KernBench paper (v5).
Generic naming (Requester Node A/B, Router, Destination Node), with
Router internals visible (two input ports -> switch -> output queue ->
output port) so the queuing delay can be located *at the out port*
rather than at the box edge. The Destination Node shows queue -> drain
slot -> processing logic.
Highlights:
* Same-size Routers, each with explicit switching logic + output queue.
* Edge labels removed -- the wires speak for themselves.
* Annotation lines are strictly vertical; arrow heads use a larger
mutation_scale so no thin line protrudes past the tip; `shrinkA`
pulls the tail away from the text so they no longer overlap.
* "flit-level interleaving on wires" (not "shared FIFO").
* Destination: queue -> drain -> processing logic (no "served").
Output:
docs/report/1H-codesign-paper/figures/latency_model.png
Per /paper isolation: this is a report-only harness under scripts/paper/.
"""
import math
from pathlib import Path
import matplotlib.patches as patches
import matplotlib.pyplot as plt
OUT = Path(
"/Users/ywkang/kernbench/docs/report/1H-codesign-paper/figures/latency_model.png"
)
OUT.parent.mkdir(parents=True, exist_ok=True)
# --- Colours --------------------------------------------------------------
C_A = "#E07B5E"
C_B = "#5E9BD1"
C_NODE = "#FFFFFF"
C_BORD = "#222222"
C_WIRE = "#222222"
C_ANN = "#333333"
C_DIM = "#777777"
C_SWITCH = "#FAFAFA"
C_OQUEUE = "#EFEFEF"
C_PROC = "#F4ECDC"
C_QUEUE = "#3CB371" # medium-sea-green: distinctive vs A/B flit colours
# --- Canvas ---------------------------------------------------------------
fig = plt.figure(figsize=(17.0, 5.0))
ax = fig.add_subplot(111)
ax.set_xlim(0, 33)
ax.set_ylim(2.2, 11.2)
ax.axis("off")
# --- Top header: end-to-end latency formula -----------------------------
ax.text(
16.5, 10.6,
r"End-to-end latency = $\Sigma$ per-node overhead + "
r"$\Sigma$ per-edge transmission + drain + "
r"queuing delay",
ha="center", fontsize=12, color=C_ANN, weight="bold",
)
ax.plot([1.5, 31.5], [10.05, 10.05], color=C_DIM, lw=0.6)
# --- Helpers --------------------------------------------------------------
def box(cx, cy, w, h, label, fs=11, fweight="normal"):
ax.add_patch(patches.FancyBboxPatch(
(cx - w / 2, cy - h / 2), w, h,
boxstyle="round,pad=0.05",
facecolor=C_NODE, edgecolor=C_BORD, linewidth=1.6,
))
if label:
ax.text(cx, cy, label, ha="center", va="center",
fontsize=fs, weight=fweight)
def draw_flit_aligned(cx, cy, angle_rad, w, h, color, label):
"""Draw a flit polygon rotated to align with the wire angle."""
cos_a, sin_a = math.cos(angle_rad), math.sin(angle_rad)
corners = []
for dx, dy in [(-w / 2, -h / 2), (w / 2, -h / 2),
(w / 2, h / 2), (-w / 2, h / 2)]:
rx = dx * cos_a - dy * sin_a
ry = dx * sin_a + dy * cos_a
corners.append((cx + rx, cy + ry))
ax.add_patch(patches.Polygon(
corners, facecolor=color, edgecolor="black", linewidth=0.5,
))
ax.text(cx, cy, label, ha="center", va="center",
fontsize=9, color="white", weight="bold")
def draw_wire_with_flits(x0, y0, x1, y1, n_flits, label_char,
color=None, flit_w=0.66, flit_h=0.95, gap=0.10):
"""Wire (line) + endpoint arrowhead + flits centred on the line.
Wires are drawn in the neutral C_WIRE colour: the wire itself is
the place where the *transmission* delay accumulates (flit_size /
BW), not where flits queue. Queuing happens before the wire, at
the FIFO at the egress side -- coloured separately.
"""
head_back = 0.30
dx, dy = x1 - x0, y1 - y0
wire_len = math.hypot(dx, dy)
ux, uy = dx / wire_len, dy / wire_len
x_line_end = x1 - ux * head_back
y_line_end = y1 - uy * head_back
ax.plot([x0, x_line_end], [y0, y_line_end],
color=C_WIRE, lw=1.6, zorder=1)
head_len, head_half = 0.30, 0.16
bx = x1 - ux * head_len
by = y1 - uy * head_len
px, py = -uy, ux
tri = [
(x1, y1),
(bx + px * head_half, by + py * head_half),
(bx - px * head_half, by - py * head_half),
]
ax.add_patch(patches.Polygon(tri, facecolor=C_WIRE,
edgecolor=C_WIRE, linewidth=0.0))
# Flit train
angle_rad = math.atan2(dy, dx)
train_len = n_flits * flit_w + (n_flits - 1) * gap
centre_p = 0.5
start_p = centre_p - (train_len / 2) / wire_len
step_p = (flit_w + gap) / wire_len
if isinstance(label_char, str):
labels = [label_char] * n_flits
cols = [color] * n_flits
else:
labels = [c[0] for c in label_char]
cols = [c[1] for c in label_char]
for i in range(n_flits):
p = start_p + (i + 0.5) * step_p
cx = x0 + p * dx
cy = y0 + p * dy
draw_flit_aligned(cx, cy, angle_rad,
flit_w, flit_h, cols[i], labels[i])
def draw_router(cx, cy, w, h, two_inputs=True):
"""Same-size Router with explicit internal switching logic and an
output queue. Returns the wire-attachment (x,y) for each input
port and the single output port:
((in1_x, in1_y), (in2_x, in2_y) or None, (out_x, out_y))
"""
# Outer box
box(cx, cy, w, h, "")
ax.text(cx, cy + h / 2 - 0.32, "Router",
ha="center", fontsize=10, weight="bold")
# Input ports (small circles on the left edge)
in_x = cx - w / 2
in_x_internal = in_x + 0.25
if two_inputs:
in_y_top = cy + 0.55
in_y_bot = cy - 0.55
for iy in (in_y_top, in_y_bot):
ax.add_patch(patches.Circle(
(in_x_internal, iy), 0.12,
facecolor="white", edgecolor=C_BORD, linewidth=0.9,
))
else:
in_y_top = None
in_y_bot = cy
ax.add_patch(patches.Circle(
(in_x_internal, in_y_bot), 0.12,
facecolor="white", edgecolor=C_BORD, linewidth=0.9,
))
# Switch (small box, centre-left-ish). sw_cy = cy so that the
# output queue and (single-input) input port are at the same y as
# the router centre -- this keeps all router-to-router edges
# strictly horizontal.
sw_w, sw_h = 0.85, 1.10
sw_cx = cx - 0.55
sw_cy = cy
# The switch is the router's processing logic -- colour it the
# same as the Destination Node's processing-logic block so the
# two read as the same "processing" concept.
ax.add_patch(patches.Rectangle(
(sw_cx - sw_w / 2, sw_cy - sw_h / 2), sw_w, sw_h,
facecolor=C_PROC, edgecolor=C_BORD, linewidth=0.8,
))
ax.text(sw_cx, sw_cy, "switch", ha="center", va="center",
fontsize=7.5, style="italic")
# Short feeder lines (no arrowhead) from input ports to the
# switch. We deliberately drop arrowheads here: the head + port
# circle were too small at this scale and read as an overlap.
sw_left_x = sw_cx - sw_w / 2
if two_inputs:
for iy in (in_y_top, in_y_bot):
ax.plot(
[in_x_internal + 0.12, sw_left_x],
[iy, sw_cy + 0.25 * ((iy - cy) / 0.55)],
color=C_DIM, lw=0.7, zorder=1,
)
else:
ax.plot(
[in_x_internal + 0.12, sw_left_x],
[in_y_bot, sw_cy],
color=C_DIM, lw=0.7, zorder=1,
)
# Output queue (small queue holding a couple of flits). This *is*
# a real queueing location, so its fill takes the C_QUEUE family.
oq_w, oq_h = 0.95, 0.55
oq_cx = cx + 0.55
oq_cy = sw_cy
ax.add_patch(patches.Rectangle(
(oq_cx - oq_w / 2, oq_cy - oq_h / 2), oq_w, oq_h,
facecolor="#D9F0E1", edgecolor=C_QUEUE, linewidth=0.9,
))
ax.text(oq_cx, oq_cy + oq_h / 2 + 0.18, "FIFO",
ha="center", fontsize=7, color=C_DIM, style="italic")
# Two small flits inside (A and B) hinting at the in-flight contents
mini_w, mini_h = 0.22, 0.34
for i, (col, lab) in enumerate([(C_A, "A"), (C_B, "B")]):
mx = oq_cx - 0.30 + i * (mini_w + 0.06)
ax.add_patch(patches.Rectangle(
(mx, oq_cy - mini_h / 2), mini_w, mini_h,
facecolor=col, edgecolor="black", linewidth=0.3,
))
ax.text(mx + mini_w / 2, oq_cy, lab,
ha="center", va="center",
fontsize=5.5, color="white", weight="bold")
# Short feeder line (no arrowhead) from switch into out queue
ax.plot(
[sw_cx + sw_w / 2, oq_cx - oq_w / 2],
[sw_cy, oq_cy],
color=C_DIM, lw=0.7, zorder=1,
)
# Short feeder line from out queue to out port
out_x_internal = cx + w / 2 - 0.25
ax.plot(
[oq_cx + oq_w / 2, out_x_internal - 0.12],
[oq_cy, oq_cy],
color=C_DIM, lw=0.7, zorder=1,
)
# Output port circle on the right edge (a port marker, not a
# queue -- the queue is on the wire that follows, not the port).
ax.add_patch(patches.Circle(
(out_x_internal, oq_cy), 0.12,
facecolor="white", edgecolor=C_BORD, linewidth=0.9,
))
return (
(in_x_internal, in_y_top) if two_inputs else None,
(in_x_internal, in_y_bot),
(out_x_internal, oq_cy),
oq_cx, # also return the out-queue centre so callouts can target it
oq_cy,
)
# --- Box layout ---------------------------------------------------------
# Requester centres set to match Router 1's input-port y so that
# Edge 1A and Edge 1B are strictly horizontal. Box heights are kept
# small enough that a visible gap separates the two Requester boxes.
R1_TMP_CY = 7.0
ReqA = (2.8, R1_TMP_CY + 0.55) # = 7.55, matches in_y_top of R1
ReqB = (2.8, R1_TMP_CY - 0.55) # = 6.45, matches in_y_bot of R1
box(*ReqA, 3.0, 0.85, "Requester\nNode A", fs=10)
box(*ReqB, 3.0, 0.85, "Requester\nNode B", fs=10)
R_W, R_H = 3.0, 2.8
R1 = (10.5, 7.0)
R2 = (20.0, 7.0)
r1_in_top, r1_in_bot, r1_out, r1_oq_cx, r1_oq_cy = draw_router(
*R1, R_W, R_H, two_inputs=True,
)
_, r2_in_bot, r2_out, r2_oq_cx, r2_oq_cy = draw_router(
*R2, R_W, R_H, two_inputs=False,
)
# Destination Node (same height as Router; wide enough so queue +
# drain + processing-logic all fit on a single horizontal row).
Dst = (28.4, 7.0)
Dst_W, Dst_H = 6.0, R_H # match the Router height
box(*Dst, Dst_W, Dst_H, "")
ax.text(Dst[0], Dst[1] + Dst_H / 2 - 0.25, "Destination Node",
ha="center", fontsize=10, weight="bold")
# --- Edges: requester -> router 1 (Edge 1A & 1B, with flits) ------------
draw_wire_with_flits(
ReqA[0] + 1.6, ReqA[1],
r1_in_top[0] - 0.02, r1_in_top[1],
n_flits=4, color=C_A, label_char="A",
)
draw_wire_with_flits(
ReqB[0] + 1.6, ReqB[1],
r1_in_bot[0] - 0.02, r1_in_bot[1],
n_flits=4, color=C_B, label_char="B",
)
# Edge 2: router1 out -> router2 in (horizontal, interleaved)
labels_e2 = [("A", C_A) if i % 2 == 0 else ("B", C_B) for i in range(8)]
E2_y = r1_out[1]
draw_wire_with_flits(
r1_out[0] + 0.02, r1_out[1],
r2_in_bot[0] - 0.02, r2_in_bot[1],
n_flits=8, label_char=labels_e2,
)
# Edge 3: router2 out -> destination (horizontal, interleaved)
labels_e3 = [("A", C_A) if i % 2 == 0 else ("B", C_B) for i in range(4)]
draw_wire_with_flits(
r2_out[0] + 0.02, r2_out[1],
Dst[0] - Dst_W / 2, r2_out[1],
n_flits=4, label_char=labels_e3,
)
# --- Destination internals: queue -> drain -> processing logic ----------
# Light-green halo behind the queue area marks it as a queueing point.
C_QUEUE_BG = "#D9F0E1"
dy_main = Dst[1] - 0.10
Dst_left = Dst[0] - Dst_W / 2
# Queue (4 flits, FIFO; rightmost is the next to drain, so the
# order is set so the serve sequence after the in-flight A alternates
# A (in drain) -> B -> A -> B -> A -- giving a clean BABA queue.
qW, qH, qGap = 0.42, 0.62, 0.07
q_labels = ["A", "B", "A", "B"]
q_cols = [C_A, C_B, C_A, C_B]
q_total = len(q_labels) * qW + (len(q_labels) - 1) * qGap
q_x_start = Dst_left + 0.35
# Green halo behind the queue boxes (the queue itself is a queueing
# location -- mark it with the C_QUEUE colour family)
ax.add_patch(patches.FancyBboxPatch(
(q_x_start - 0.10, dy_main - qH / 2 - 0.08),
q_total + 0.20, qH + 0.16,
boxstyle="round,pad=0.02",
facecolor=C_QUEUE_BG, edgecolor=C_QUEUE,
linewidth=0.9, zorder=1.5,
))
for i, (lab, col) in enumerate(zip(q_labels, q_cols)):
qx = q_x_start + i * (qW + qGap)
ax.add_patch(patches.Rectangle(
(qx, dy_main - qH / 2), qW, qH,
facecolor=col, edgecolor="black", linewidth=0.4,
zorder=2,
))
ax.text(qx + qW / 2, dy_main, lab,
ha="center", va="center",
fontsize=8, color="white", weight="bold", zorder=3)
# Common y for the "queue" / "drain" labels. Matches the FIFO label
# spacing inside the Router (0.18 above the box top) and uses the same
# font size for visual consistency.
DST_LABEL_Y = dy_main + qH / 2 + 0.18
ax.text(q_x_start + q_total / 2, DST_LABEL_Y,
"queue",
ha="center", fontsize=7, style="italic", color=C_DIM)
# Drain slot (height matched to queue boxes; font matched to FIFO/queue
# labels for visual consistency)
dr_W, dr_H = 0.85, qH
dr_x = q_x_start + q_total + 0.35
dr_cx = dr_x + dr_W / 2
dr_cy = dy_main
ax.add_patch(patches.FancyBboxPatch(
(dr_x, dr_cy - dr_H / 2), dr_W, dr_H,
boxstyle="round,pad=0.03",
facecolor=C_A, edgecolor="black", linewidth=1.0,
))
ax.text(dr_cx, dr_cy, "A", ha="center", va="center",
fontsize=9, color="white", weight="bold")
ax.text(dr_cx, DST_LABEL_Y,
"drain",
ha="center", fontsize=7, style="italic", color=C_DIM)
# Drain-time double-arrow underneath the drain slot. Move the "drain
# time" text a touch further down so the descender does not collide
# with the arrowhead.
dt_y = dr_cy - dr_H / 2 - 0.28
ax.annotate(
"", xy=(dr_x + dr_W, dt_y), xytext=(dr_x, dt_y),
arrowprops=dict(arrowstyle="<->", lw=0.9, color=C_ANN),
)
ax.text(dr_cx, dt_y - 0.42, "drain time",
ha="center", fontsize=8.5, color=C_ANN, style="italic")
# Processing-logic block (after drain) — wider so the two-line text
# does not collide with the box edges
pl_W, pl_H = 1.75, 1.00
pl_x = dr_x + dr_W + 0.40
pl_cx = pl_x + pl_W / 2
pl_cy = dr_cy
ax.add_patch(patches.FancyBboxPatch(
(pl_x, pl_cy - pl_H / 2), pl_W, pl_H,
boxstyle="round,pad=0.03",
facecolor=C_PROC, edgecolor=C_BORD, linewidth=1.0,
))
ax.text(pl_cx, pl_cy, "processing\nlogic",
ha="center", va="center", fontsize=9)
# Small gray arrows along the queue -> drain -> processing chain
ax.annotate(
"", xy=(dr_x - 0.04, dr_cy),
xytext=(q_x_start + q_total + 0.13, dr_cy),
arrowprops=dict(arrowstyle="-|>", lw=0.7,
color=C_DIM, mutation_scale=7),
)
ax.annotate(
"", xy=(pl_x - 0.04, dr_cy),
xytext=(dr_x + dr_W + 0.04, dr_cy),
arrowprops=dict(arrowstyle="-|>", lw=0.7,
color=C_DIM, mutation_scale=7),
)
# --- Annotations (strictly vertical, no diagonals, no overlap) ---------
def vcallout(text, xy, x_text_top_y, fs=10, clearance=0.55,
marker_color=None):
"""Vertical dotted callout. Optional `marker_color` paints a small
coloured circle just left of the text -- used to tie the label to
a colour code in the diagram.
"""
text_y = x_text_top_y
if text_y < xy[1]:
line_top_y = text_y + clearance
else:
line_top_y = text_y - clearance
ax.plot([xy[0], xy[0]], [xy[1], line_top_y],
color=C_ANN, lw=0.9, linestyle=":", zorder=2)
if marker_color is not None:
# Text-width estimate at fs=10 (~0.18 data units per char) so
# the marker is placed clearly to the left of the text.
text_w = len(text) * 0.18
marker_x = xy[0] - text_w / 2 - 0.30
ax.add_patch(patches.Circle(
(marker_x, text_y), 0.14,
facecolor=marker_color, edgecolor=C_BORD, linewidth=0.6,
zorder=3,
))
ax.text(xy[0], text_y, text,
ha="center", va="center",
fontsize=fs, color=C_ANN, zorder=3)
# Transmission delay -> flit on Edge 2 (text BELOW the wire).
# Extra clearance ~ 2 x text height so the line stops well clear of
# the label.
e2_total = 8 * 0.66 + 7 * 0.10
e2_centre = (r1_out[0] + r2_in_bot[0]) / 2
e2_start = e2_centre - e2_total / 2
trans_idx = 5
trans_cx = e2_start + (trans_idx + 0.5) * (0.66 + 0.10)
vcallout(
"transmission delay = flit_size / BW",
xy=(trans_cx, E2_y - 0.55),
x_text_top_y=E2_y - 2.6,
clearance=0.45,
)
# Flit-level interleaving -> ABOVE the wire, anchored on a flit near
# the centre/right of the wire. Text y matches the queuing-delay
# callout so the two labels sit on the same horizontal row, separated
# horizontally so they don't collide.
int_idx = 4
int_cx = e2_start + (int_idx + 0.5) * (0.66 + 0.10)
vcallout(
"flit-level interleaving on wires",
xy=(int_cx, E2_y + 0.55),
x_text_top_y=R1[1] + R_H / 2 + 0.85,
)
# Queuing delay -> at Router 1 out-queue. Text positioned just above
# the Router-1 box.
vcallout(
"queuing delay",
xy=(r1_oq_cx, r1_oq_cy + 0.30),
x_text_top_y=R1[1] + R_H / 2 + 0.85,
marker_color=C_QUEUE,
)
# Drain -> below the drain slot (text is below).
# Extra clearance ~ 2.5 x text height so the line stops further from
# the label.
vcallout(
"drain = per-flit service occupancy",
xy=(dr_cx, dt_y - 0.66), # just *inside* the Destination Node
# bottom edge -- the dotted line then
# penetrates the box slightly rather
# than hanging below it
x_text_top_y=E2_y - 2.6,
clearance=0.60,
)
# Per-node overhead -> points at Router 1's switch block (where the
# component's fixed processing cost lives). Text on the same row as
# transmission delay and drain so the three "below the wire" callouts
# sit on one horizontal baseline.
R1_SW_CX = R1[0] - 0.55 # sw_cx for Router 1
R1_SW_BOTTOM_Y = R1[1] - 0.55 # bottom of sw box
vcallout(
"per-node overhead",
xy=(R1_SW_CX, R1_SW_BOTTOM_Y),
x_text_top_y=E2_y - 2.6,
clearance=0.60,
marker_color=C_PROC,
)
# --- Legend (close to the diagram) --------------------------------------
LX, LY = 1.0, 2.6
ax.add_patch(patches.Rectangle((LX, LY), 0.7, 0.45,
facecolor=C_A, edgecolor="black",
linewidth=0.4))
ax.text(LX + 0.95, LY + 0.22, "Transaction A flit",
va="center", fontsize=10)
ax.add_patch(patches.Rectangle((LX + 4.8, LY), 0.7, 0.45,
facecolor=C_B, edgecolor="black",
linewidth=0.4))
ax.text(LX + 5.75, LY + 0.22, "Transaction B flit",
va="center", fontsize=10)
# --- Save ----------------------------------------------------------------
fig.savefig(OUT, dpi=140, bbox_inches="tight", facecolor="white")
print(f"Wrote {OUT}")
scripts/paper/paper_plot_gemm_async_vs_composite.py
+261
View File
@@ -0,0 +1,261 @@
"""Composite vs async-loading GEMM TFLOPS comparison.
For each shape in the milestone sweep, run both benches:
- matmul-composite (load_ref variant — A pre-staged, B streamed by
scheduler inside one composite command)
- matmul-async (A and B both async-loaded via tl.load, then a single
tl.dot — no per-tile overlap of streaming-B with GEMM)
Compute per-PE TFLOPS = 2*M*K*N / pe_window_ns for each kernel and emit a
side-by-side bar chart PNG to
src/kernbench/benches/1H_milestone_output/gemm/gemm_composite_vs_async_tflops.png
Run from repo root:
python scripts/paper/paper_plot_gemm_async_vs_composite.py
"""
from __future__ import annotations
import json
import os
import sys
from pathlib import Path
REPO = Path(__file__).resolve().parents[2]
OUT_DIR = REPO / "src" / "kernbench" / "benches" / "1H_milestone_output" / "gemm"
OUT_PNG = OUT_DIR / "gemm_composite_vs_async_tflops.png"
OUT_JSON = OUT_DIR / "gemm_composite_vs_async.json"
TOPO_PATH = REPO / "topology.yaml"
# Same 7 shapes as the milestone composite sweep (skip square=512: B alone
# is 512 KiB; async needs all of A+B+out in TCM scratch and the 512×512
# square pushes total scratch use past the 1 MiB cap).
SHAPES: list[tuple[int, int, int]] = [
(32, 32, 32), # 1 tile, K<TILE_K (under-tile in K)
(32, 64, 32), # 1 tile
(32, 128, 32), # 2 tiles
(32, 128, 128), # 8 tiles
(32, 3072, 32), # 48 tiles (deep-K)
(8, 128, 128), # under-tile in M
(128, 8, 128), # under-tile in K
]
ENGINES = ("pe_dma", "pe_fetch_store", "pe_gemm", "pe_math")
STAGES = ("DMA_READ", "DMA_WRITE", "FETCH", "STORE", "GEMM", "MATH")
def _pe_records(op_log):
return [r for r in op_log
if any(r.component_id.endswith("." + e) for e in ENGINES)]
def _pe_window_ns(op_log) -> float:
pe = _pe_records(op_log)
if not pe:
return 0.0
return max(r.t_end for r in pe) - min(r.t_start for r in pe)
def _composite_window_ns(op_log) -> float:
"""For the composite kernel: window of records carrying a stage_type
set by the composite plan (DMA_READ, FETCH, GEMM, STORE, DMA_WRITE).
Excludes the initial up-front tl.load(A) because that record is an
atomic DmaReadCmd with no stage_type. Matches the existing
milestone_1h_gemm.py / gemm_per_pe_tflops.png methodology.
"""
stage_records = [r for r in op_log
if r.params.get("stage_type") in STAGES]
if not stage_records:
return 0.0
return max(r.t_end for r in stage_records) \
- min(r.t_start for r in stage_records)
def _async_engine_window_ns(op_log) -> float:
"""For the async kernel: engine pipeline window that excludes the
initial tl.load(A), paralleling composite_window's exclusion of the
up-front A pre-stage. The first DMA_READ record on pe_dma is the
tl.load(A); the window starts at the SECOND pe_dma record's t_start
(= tl.load(B)) and ends at the last engine record's t_end.
"""
pe = _pe_records(op_log)
if not pe:
return 0.0
dma = sorted(
(r for r in op_log if r.component_id.endswith(".pe_dma")),
key=lambda r: r.t_start,
)
if len(dma) < 2:
return _pe_window_ns(op_log)
window_start = dma[1].t_start
window_end = max(r.t_end for r in pe)
return window_end - window_start
def _run_one(bench_name: str, variant: str | None, M: int, K: int, N: int) -> dict:
os.environ["MATMUL_M"] = str(M)
os.environ["MATMUL_K"] = str(K)
os.environ["MATMUL_N"] = str(N)
if variant is not None:
os.environ["MATMUL_VARIANT"] = variant
elif "MATMUL_VARIANT" in os.environ:
del os.environ["MATMUL_VARIANT"]
from kernbench.benches.registry import resolve as resolve_bench
from kernbench.runtime_api.bench_runner import run_bench
from kernbench.runtime_api.types import resolve_device
from kernbench.sim_engine.engine import GraphEngine
from kernbench.topology.builder import resolve_topology
# Chunked async loads A in TILE_K-sized blocks at sequential offsets;
# those offsets do not match the row-major (M, K) layout of A, so the
# DataExecutor would fail on the resulting mismatched read. The
# simulator's *timing* path doesn't need correct data — only the
# number of bytes per DMA matters — so bypass DataExecutor for the
# chunked kernel. Composite/naive-async loads use the full (M, K)
# shape and remain data-correct.
if bench_name in ("matmul-async-chunked", "matmul-async-chunked-db"):
GraphEngine._flush_data_phase = lambda self: None
topo = resolve_topology(str(TOPO_PATH))
bench = resolve_bench(bench_name).run
device = resolve_device(None)
result = run_bench(
topology=topo, bench_fn=bench, device=device,
engine_factory=lambda t, d: GraphEngine(
getattr(t, "topology_obj", t), enable_data=True,
),
)
if not result.completion.ok:
raise RuntimeError(f"{bench_name} failed at {M}x{K}x{N}: {result.completion}")
log = result.engine.op_log
pe_window = _pe_window_ns(log)
if bench_name == "matmul-composite":
engine_window = _composite_window_ns(log)
elif bench_name in ("matmul-async-chunked", "matmul-async-chunked-db"):
# First N_chunks pe_dma records are A pre-stage; engine window
# starts at the (N_chunks+1)-th pe_dma record (= first B-chunk
# load). Fall back to the naive analog for K <= TILE_K (kernel
# collapses to single load+dot+store).
TILE_K = 64
n_chunks = max(K // TILE_K, 1)
if n_chunks <= 1:
engine_window = _async_engine_window_ns(log)
else:
dma = sorted(
(r for r in log if r.component_id.endswith(".pe_dma")),
key=lambda r: r.t_start,
)
pe = _pe_records(log)
if len(dma) > n_chunks and pe:
engine_window = max(r.t_end for r in pe) - dma[n_chunks].t_start
else:
engine_window = _async_engine_window_ns(log)
else:
engine_window = _async_engine_window_ns(log)
flops = 2 * M * K * N
# flops / ns * 1e-3 = TFLOP/s (since 1 flop/ns = 1 GFLOP/s)
return {
"M": M, "K": K, "N": N,
"bench": bench_name, "variant": variant,
"pe_window_ns": pe_window,
"engine_window_ns": engine_window,
"flops": flops,
"tflops": (flops / engine_window / 1000.0) if engine_window > 0 else 0.0,
"n_records": len(log),
}
def collect() -> list[dict]:
rows: list[dict] = []
for M, K, N in SHAPES:
print(f" shape M={M:4d} K={K:5d} N={N:4d} ...", flush=True)
comp = _run_one("matmul-composite", "load_ref", M, K, N)
asyn = _run_one("matmul-async", None, M, K, N)
chnk = _run_one("matmul-async-chunked", None, M, K, N)
chnkdb = _run_one("matmul-async-chunked-db", None, M, K, N)
print(f" composite: {comp['engine_window_ns']:8.1f} ns "
f"{comp['tflops']:6.3f} TFLOPS")
print(f" async naive: {asyn['engine_window_ns']:8.1f} ns "
f"{asyn['tflops']:6.3f} TFLOPS")
print(f" chunked all: {chnk['engine_window_ns']:8.1f} ns "
f"{chnk['tflops']:6.3f} TFLOPS")
print(f" chunked db=2: {chnkdb['engine_window_ns']:8.1f} ns "
f"{chnkdb['tflops']:6.3f} TFLOPS")
rows.append(comp)
rows.append(asyn)
rows.append(chnk)
rows.append(chnkdb)
return rows
def plot(rows: list[dict]) -> None:
import matplotlib.pyplot as plt
import numpy as np
shape_keys = []
for M, K, N in SHAPES:
shape_keys.append((M, K, N))
by_key: dict[tuple[int, int, int], dict[str, dict]] = {}
for r in rows:
key = (r["M"], r["K"], r["N"])
by_key.setdefault(key, {})[r["bench"]] = r
labels = [f"M={M}\nK={K}\nN={N}" for (M, K, N) in shape_keys]
comp_tflops = [by_key[k]["matmul-composite"]["tflops"] for k in shape_keys]
asyn_tflops = [by_key[k]["matmul-async"]["tflops"] for k in shape_keys]
chnk_tflops = [by_key[k]["matmul-async-chunked"]["tflops"] for k in shape_keys]
chnkdb_tflops = [by_key[k]["matmul-async-chunked-db"]["tflops"]
for k in shape_keys]
x = np.arange(len(shape_keys))
width = 0.20
fig, ax = plt.subplots(figsize=(14, 5.5))
b_c = ax.bar(x - 1.5*width, comp_tflops, width,
label="Composite (load_ref)", color="#10b981")
b_a = ax.bar(x - 0.5*width, asyn_tflops, width,
label="Async naive (tl.load full + tl.dot)",
color="#f59e0b")
b_db = ax.bar(x + 0.5*width, chnkdb_tflops, width,
label="Async chunked-prefetch, depth=2 "
"(TCM-bounded)", color="#a855f7")
b_k = ax.bar(x + 1.5*width, chnk_tflops, width,
label="Async chunked-prefetch, depth=$\\infty$ "
"(all B-tiles queued up front)", color="#3b82f6")
ax.axhline(8.0, linestyle="--", color="#94a3b8", linewidth=0.8,
label="Per-PE GEMM peak (8 TFLOP/s)")
ax.set_ylabel("Per-PE achieved TFLOP/s")
ax.set_title("Composite vs async-loading GEMM — per-PE throughput "
"(engine pipeline window, A pre-stage excluded)")
ax.set_xticks(x)
ax.set_xticklabels(labels, fontsize=8)
ax.grid(axis="y", alpha=0.25)
ax.legend(loc="upper left", fontsize=9)
for bar in (*b_c, *b_a, *b_db, *b_k):
h = bar.get_height()
ax.annotate(f"{h:.2f}", xy=(bar.get_x() + bar.get_width()/2, h),
xytext=(0, 2), textcoords="offset points",
ha="center", va="bottom", fontsize=7, color="#475569")
fig.tight_layout()
OUT_DIR.mkdir(parents=True, exist_ok=True)
fig.savefig(OUT_PNG, dpi=150)
print(f"wrote {OUT_PNG.relative_to(REPO)}")
plt.close(fig)
def main() -> int:
rows = collect()
OUT_DIR.mkdir(parents=True, exist_ok=True)
OUT_JSON.write_text(json.dumps(rows, indent=2))
print(f"wrote {OUT_JSON.relative_to(REPO)}")
plot(rows)
return 0
if __name__ == "__main__":
sys.exit(main())
scripts/paper/paper_plot_gqa.py
+2 -4
View File
@@ -24,10 +24,8 @@ _FIG_DIR = Path(__file__).resolve().parents[2] / "docs" / "report" / "1H-codesig
_IN_JSON = _FIG_DIR / "gqa_latency.json"
_LABELS = {
"single_user_prefill_gqa": "prefill\nC=1",
"multi_user_prefill_gqa": "prefill\nC=4 (Ring KV)",
"single_user_decode_gqa": "decode\nC=1, P=8",
"multi_user_decode_gqa": "decode\nC=4, P=8",
"single_kv_group_prefill_gqa_c8_p8":
"prefill\nC=8, P=8\n(single KV group)",
}
scripts/paper/paper_plot_gqa_4cases_summary.py
+638
View File
@@ -0,0 +1,638 @@
"""GQA 4-cases (+ 2 d_head-TP variants) — combined memory + comm summary.
Single PNG with two panels side-by-side:
(left) Per-PE KV memory at S_kv = 1 M tokens (across 80 layers).
(right) Per-PE communication per output token at S_kv = 1 M (decode).
================================================================
SYSTEM INPUTS (LLaMA-3.1-70B single-KV-head group)
================================================================
N_layers = 80 (transformer layers per token)
h_kv = 1 (per KV group)
h_q = 8 (query heads per KV group)
d_head = 128
d_model = 8192 (LLaMA-3.1-70B hidden dim — for Wo / FFN AR)
bytes = 2 (FP16)
C = 8 cubes per KV group P = 8 PEs per cube
HBM_per_PE = 6.0 GB
Attn weights = 1.76 GB → KV headroom = 4.24 GB / PE
B = 1 user S_q = 1 token (decode)
================================================================
(1) PER-PE KV MEMORY (left panel)
================================================================
KV bytes per token across all 80 layers, single KV group:
KV/tok = 2 (K+V) · h_kv · d_head · bytes · N_layers
= 2 · 1 · 128 · 2 · 80
= 40 KB / token
Per-PE share = KV/tok ÷ divisor, where divisor depends on the K/V
tensor placement (NOT on the compute-side label):
Case 1 Cube-SP × PE-replicate divisor = C = 8
Case 1' Cube-SP × PE-TP (on d_head) divisor = C·P = 64
Case 2 Cube-Repl × PE-replicate divisor = 1 = 1
Case 3 Cube-Repl × PE-SP divisor = P = 8
Case 3' Cube-TP × PE-SP divisor = C·P = 64
Case 4 Cube-SP × PE-SP ★ divisor = C·P = 64
Per-PE bytes @ S_kv = 1 M = KV/tok · 1 M / divisor:
Case 1 → 5.0 GB ✗ (exceeds 4.24 GB headroom)
Case 1' → 640 MB ✓
Case 2 → 40.0 GB ✗
Case 3 → 5.0 GB ✗
Case 3' → 640 MB ✓
Case 4 → 640 MB ✓
Max KV context per PE = 4.24 GB · 1024² / (KV/tok ÷ divisor):
Case 1, 3 : 889 K tokens
Case 1', 3', 4 : 7.11 M tokens ✓
Case 2 : 111 K tokens
================================================================
(2) PER-PE COMMUNICATION (right panel)
================================================================
(A) WEIGHT AllReduces (constant across all cases)
Wo AR ≈ 8 KB / layer / PE (Wq replicated → AR partial Y)
FFN AR ≈ 8 KB / layer / PE
× 80 layers = 1.25 MB / token / PE (the bottom blue stack on every bar)
(B) ATTENTION-TIME collective (the differentiator)
Cases that compute partial attention locally (PE-SP / PE-repl):
Q is replicated and each rank attends to its own complete KV slice
(full d_head) → only the small (m, ℓ, O) triple needs AllReducing.
Payload (h_q · S_q · d_head · 2) ≈ 4 KB per AR step, hierarchical ≈ 32 KB / layer.
CONSTANT in S_kv.
Case 1 : inter-cube AR on (m,ℓ,O) → ~32 KB / layer
Case 2 : nothing → 0
Case 3 : intra-cube AR on (m,ℓ,O) → ~32 KB / layer
Case 4 ★ : intra + inter-cube AR on (m,ℓ,O) → ~64 KB / layer
d_head-TP variants (1', 3'):
Each rank holds only d_head/divisor dims → Q·K^T produces only
partial sums → must AR PARTIAL SCORES before softmax. The score
tile per AR is (h_q · S_q · S_kv / slices_for_seq) · 2 bytes,
which SCALES with S_kv.
Case 1' : intra-cube partial-score AR
payload = h_q · S_q · (S_kv/C) · 2 = 2 MB / layer at 1 M
Case 3' : inter-cube partial-score AR (UCIe, slower than NoC)
payload = h_q · S_q · (S_kv/P) · 2 = 2 MB / layer at 1 M
Plus tiny (m,ℓ,O) AR for both 1' and 3' to combine partial attentions
across the remaining axis (~32 KB / layer).
(C) Total per output token per PE (sum over 80 layers, S_kv = 1 M):
Case 1 ≈ 4 MB (1.25 MB Wo+FFN + 2.5 MB inter-cube AR)
Case 1' ≈ 166 MB (1.25 MB + 160 MB partial-score AR + ~5 MB other)
Case 2 ≈ 1.2 MB (Wo+FFN only)
Case 3 ≈ 4 MB (1.25 MB + 2.5 MB intra-cube AR)
Case 3' ≈ 166 MB (1.25 MB + 160 MB partial-score AR + ~5 MB other)
Case 4 ★ ≈ 6 MB (1.25 MB + 5 MB 2-phase AR)
================================================================
KEY TAKEAWAYS
================================================================
- Memory winners (640 MB / PE @ 1M): Cases 1', 3', 4 (any 64-way sharding).
- Comm winner among those three: Case 4 (6 MB). Cases 1' and 3' both
pay ~160 MB because the d_head-sharded variant ARs a 2 MB partial-
score tile every layer — the score tensor scales with S_kv while
the (m, ℓ, O) triple does not.
- Cases 1 and 3 (8-way sharding) don't fit 1M context (5 GB vs 4.24 GB
headroom).
- Case 2 has the cheapest comm but the most memory (40 GB / PE).
- Case 4 ★ is the Pareto-best: fits 1M context AND lowest comm
among memory-feasible options.
Output PNG:
src/kernbench/benches/1H_milestone_output/gqa/long_ctx/
gqa_4cases_summary.png
"""
from __future__ import annotations
import json
import textwrap
from pathlib import Path
import matplotlib.patches as mpatches
import matplotlib.pyplot as plt
# ── System constants (LLaMA-3.1-70B single-KV-head group, decode) ───
_N_LAYERS = 80
_H_KV = 1
_H_Q = 8
_D_HEAD = 128
_D_MODEL = 8192
_BYTES_PER_ELEM = 2 # FP16
_C = 8 # cubes per KV group
_P = 8 # PEs per cube
_B = 1
_S_Q = 1 # decode
_HBM_PER_PE_GB = 6.0
_WEIGHTS_PER_PE_GB = 1.76
_HEADROOM_GB = _HBM_PER_PE_GB - _WEIGHTS_PER_PE_GB # 4.24 GB
_HEADLINE_S_KV = 1 << 20 # 1 Mi tokens
# Per-token KV bytes (single KV group, all 80 layers).
_KV_PER_TOK_BYTES = (
2 * _H_KV * _D_HEAD * _BYTES_PER_ELEM * _N_LAYERS
) # 40 KB
# Per-token Wo + FFN AR (constant across cases).
_WO_PER_LAYER_BYTES = 8 * 1024
_FFN_PER_LAYER_BYTES = 8 * 1024
# ── Per-PE attention weight breakdown (single KV-head group) ────────
#
# LLaMA-3.1-70B single-KV-head group dimensions:
# d_model = 8192
# h_q per group = 8 (8 query heads attend to 1 KV head per group)
# h_kv per group = 1
# d_head = 128
# FP16 (2 bytes)
#
# Per layer, per PE:
# Wq shape (d_model, h_q · d_head) REPLICATED across all 64 PEs of group
# = 8192 · 8·128 · 2 bytes = 16 MB / layer / PE
# Wk shape (d_model, h_kv · d_head) REPLICATED
# = 8192 · 1·128 · 2 bytes = 2 MB / layer / PE
# Wv shape (d_model, h_kv · d_head) REPLICATED
# = 8192 · 1·128 · 2 bytes = 2 MB / layer / PE
# Wo shape (h_q · d_head, d_model) ROW-SPLIT across C=8 cubes,
# replicated within cube
# = (1024/C) · 8192 · 2 bytes = 2 MB / layer / PE
# ────────────────────────────────────────────────────────────────
# total attn weights / layer / PE = 22 MB
# × 80 layers = 1.76 GB / PE (decimal GB)
#
# KV headroom per PE = HBM_PER_PE - weights = 6.0 - 1.76 = 4.24 GB.
# (FFN weights are accounted for in a separate budget, not in this 4.24 GB.)
# Slide-17 convention: per-layer values in binary MB (MiB), totals in
# "GB" formed by ×80 layers ÷ 1000 — giving the canonical 1.76 GB total
# and 4.24 GB headroom that the slide-17 chart reports.
_WQ_MB_PER_LAYER = (_D_MODEL * _H_Q * _D_HEAD * _BYTES_PER_ELEM) / (1024 ** 2) # 16
_WK_MB_PER_LAYER = (_D_MODEL * _H_KV * _D_HEAD * _BYTES_PER_ELEM) / (1024 ** 2) # 2
_WV_MB_PER_LAYER = (_D_MODEL * _H_KV * _D_HEAD * _BYTES_PER_ELEM) / (1024 ** 2) # 2
_WO_MB_PER_LAYER = ((_H_Q * _D_HEAD) // _C * _D_MODEL
* _BYTES_PER_ELEM) / (1024 ** 2) # 2
_WQ_GB = _WQ_MB_PER_LAYER * _N_LAYERS / 1000 # 1.28 GB
_WK_GB = _WK_MB_PER_LAYER * _N_LAYERS / 1000 # 0.16 GB
_WV_GB = _WV_MB_PER_LAYER * _N_LAYERS / 1000 # 0.16 GB
_WO_GB = _WO_MB_PER_LAYER * _N_LAYERS / 1000 # 0.16 GB
_WEIGHTS_GB = _WQ_GB + _WK_GB + _WV_GB + _WO_GB # 1.76 GB
# (m, ℓ, O) AR cost per layer, derived from the kernel topology.
#
# Per-PE payload T for one (m, ℓ, O) merge step:
# O — h_q · S_q · d_head · 2 bytes (FP16)
# m — h_q · S_q · 4 bytes (FP32)
# ℓ — h_q · S_q · 4 bytes (FP32)
# T = h_q · S_q · (d_head · 2 + 8) ≈ 2.1 KB
#
# Reduce algorithm: the decode kernels use hierarchical reduce-only
# (chain / tree) rather than ring all-reduce, because the merged result
# only needs to land on the cube that runs the downstream Wo gemm —
# not on every PE. For a chain of N participants the total traffic is
# (N-1)·T and the per-PE average is (N-1)/N · T.
#
# The previous version used a hard-coded 32 KB / layer placeholder which
# overestimated the per-PE cost by ~6× (and ~3× even under a hypothetical
# ring-AR assumption). See `gqa_long_ctx_6cases_measured_comm.json` for
# the measured numbers this matches against.
_MLO_PAYLOAD_BYTES = _H_Q * _S_Q * (_D_HEAD * _BYTES_PER_ELEM + 8)
def _reduce_chain_per_pe_bytes(n_participants: int) -> int:
"""Per-PE average bytes for a single-stage chain/tree reduce of T."""
if n_participants <= 1:
return 0
return _MLO_PAYLOAD_BYTES * (n_participants - 1) // n_participants
_MLO_INTRA_BYTES_PER_LAYER = _reduce_chain_per_pe_bytes(_P) # PE-axis
_MLO_INTER_BYTES_PER_LAYER = _reduce_chain_per_pe_bytes(_C) # cube-axis
# Cases — renumbered in MEMORY-DESCENDING ORDER (left to right):
# Case 1 (40 GB) : no sharding
# Cases 2, 3 ( 5 GB) : single-axis sharding (cube OR PE only)
# Cases 4, 5, 6 (640 MB) : two-axis sharding (cube AND PE)
# — Case 6 ★ is the Pareto-best (lowest comm)
_CASES = (1, 2, 3, 4, 5, 6)
_DIVISOR = {
1: 1, # Cube-Repl × PE-replicate — no sharding
2: _C, # Cube-SP × PE-replicate — cube-axis only
3: _P, # Cube-Repl × PE-SP — PE-axis (S_kv) only
4: _C * _P, # Cube-SP × PE-TP (d_head) — 64-way (d_head intra)
5: _C * _P, # Cube-TP × PE-SP — 64-way (d_head inter)
6: _C * _P, # Cube-SP × PE-SP — 64-way (S_kv both axes) ★
}
_CASE_LABEL = {
1: "Case 1\nCube-Repl\nPE-repl",
2: "Case 2\nCube-SP\nPE-repl",
3: "Case 3\nCube-Repl\nPE-SP",
4: "Case 4\nCube-SP\nPE-TP",
5: "Case 5\nCube-TP\nPE-SP",
6: "Case 6 ★\nCube-SP\nPE-SP",
}
_CASE_COLOR = {
1: "#C0504D", # red — worst memory (no sharding)
2: "#E0834A", # orange — single-axis sharded (cube)
3: "#EBA854", # tan — single-axis sharded (PE)
4: "#A6C2E0", # light blue — d_head-TP 64-way (PE)
5: "#C7D8A0", # light green — d_head-TP 64-way (cube)
6: "#8064A2", # purple — Pareto winner ★ (S_kv 64-way)
}
_ATTN_DESC = {
1: "none",
2: "online-softmax (m,ℓ,O) — inter-cube",
3: "online-softmax (m,ℓ,O) — intra-cube",
4: "partial scores + (m,ℓ,O) merge (d_head-TP)",
5: "partial scores + (m,ℓ,O) merge (d_head-TP)",
6: "online-softmax (m,ℓ,O) — intra + inter",
}
_WO_COLOR = "#9EC5E8"
_FFN_COLOR = "#4A78B8"
_ATTN_COLOR = "#E07A3F"
_OUT_DIR = (
Path(__file__).resolve().parents[2]
/ "src" / "kernbench" / "benches"
/ "1H_milestone_output" / "gqa" / "long_ctx"
)
_MEASURED_JSON = _OUT_DIR / "gqa_long_ctx_6cases_measured_comm.json"
def _load_measured() -> dict[int, float] | None:
"""Load measured per-PE comm bytes (already scaled to S_kv=1M).
Returns {case_id: per_token_total_bytes} or None if JSON missing.
Produced by scripts/paper/measure_gqa_decode_placement_comm.py.
"""
if not _MEASURED_JSON.exists():
return None
data = json.loads(_MEASURED_JSON.read_text())
return {
int(cid): info["per_pe_total_bytes_per_token_at_1M"]
for cid, info in data["cases"].items()
}
# ── Formulae ────────────────────────────────────────────────────────
def kv_per_pe_bytes(case, s_kv: int) -> int:
"""Per-PE KV bytes at the given S_kv."""
return _KV_PER_TOK_BYTES * s_kv // _DIVISOR[case]
def max_s_kv(case) -> int:
"""Max KV context per PE given the 4.24 GB headroom."""
headroom_bytes = int(_HEADROOM_GB * (1 << 30))
bytes_per_tok = _KV_PER_TOK_BYTES // _DIVISOR[case]
return headroom_bytes // bytes_per_tok
def _partial_score_bytes_per_layer(s_kv: int, slices_for_seq: int) -> int:
"""AR payload for partial scores in d_head-TP cases."""
return _H_Q * _S_Q * (s_kv // slices_for_seq) * _BYTES_PER_ELEM
def attn_comm_per_layer_bytes(case: int, s_kv: int) -> int:
"""Per-layer per-PE attention-time comm bytes (decode, B=1, S_q=1).
Case numbering follows the memory-descending order defined in
_CASES (1=no sharding, 6=Pareto winner).
"""
if case == 1: # Cube-Repl × PE-repl — no sharding
return 0
if case == 2: # Cube-SP × PE-repl — inter-cube AR on (m,ℓ,O)
return _MLO_INTER_BYTES_PER_LAYER
if case == 3: # Cube-Repl × PE-SP — intra-cube AR on (m,ℓ,O)
return _MLO_INTRA_BYTES_PER_LAYER
if case == 4: # Cube-SP × PE-TP(d_head) — partial-score AR
return (_partial_score_bytes_per_layer(s_kv, _C)
+ _MLO_INTRA_BYTES_PER_LAYER
+ _MLO_INTER_BYTES_PER_LAYER)
if case == 5: # Cube-TP(d_head) × PE-SP — partial-score AR
return (_partial_score_bytes_per_layer(s_kv, _P)
+ _MLO_INTRA_BYTES_PER_LAYER
+ _MLO_INTER_BYTES_PER_LAYER)
if case == 6: # Cube-SP × PE-SP — 2-phase (m,ℓ,O) AR ★
return _MLO_INTRA_BYTES_PER_LAYER + _MLO_INTER_BYTES_PER_LAYER
raise ValueError(f"unknown case {case}")
def per_token_bytes(case, s_kv: int) -> tuple[int, int, int]:
"""(Wo AR, FFN AR, Attn) bytes per output token per PE — × 80 layers."""
wo = _WO_PER_LAYER_BYTES * _N_LAYERS
ffn = _FFN_PER_LAYER_BYTES * _N_LAYERS
attn = attn_comm_per_layer_bytes(case, s_kv) * _N_LAYERS
return wo, ffn, attn
# ── Formatters ──────────────────────────────────────────────────────
def _fmt_bytes(b: float) -> str:
if b >= (1 << 30):
return f"{b / (1 << 30):.2f} GB"
if b >= (1 << 20):
return f"{b / (1 << 20):.1f} MB"
if b >= (1 << 10):
return f"{b / (1 << 10):.0f} KB"
return f"{b:.0f} B"
def _fmt_tokens(n: int) -> str:
if n >= 1_000_000:
return f"{n / 1_000_000:.2f} M"
if n >= 1_000:
return f"{n / 1_000:.0f} K"
return f"{n}"
# ── Panels ──────────────────────────────────────────────────────────
def _plot_budget(ax) -> None:
"""HBM budget per PE — stacked weights + KV headroom + HBM ceiling.
Wq is the dominant weight slice (~1.28 GB). Wk, Wv, Wo are each
small (~0.16 GB) so their slice labels would overlap on the bar —
they're shown in the legend only, and only Wq + KV-headroom get
on-bar annotations.
"""
components = [
("Wq (REPL)", _WQ_GB, "#7B9CC4"),
("Wk (REPL)", _WK_GB, "#A0BBD8"),
("Wv (REPL)", _WV_GB, "#C5D6E8"),
("Wo (cube-split)", _WO_GB, "#E2EAF3"),
("KV cache headroom",
_HBM_PER_PE_GB - _WEIGHTS_GB, "#9BBB59"),
]
bottom = 0.0
for label, val, color in components:
ax.bar(0, val, bottom=bottom, color=color, edgecolor="black",
width=0.7, label=f"{label} · {val:.2f} GB")
# Only annotate slices thick enough to fit text without overlap.
if val >= 0.50:
ax.text(0, bottom + val / 2, f"{label}\n{val:.2f} GB",
ha="center", va="center", fontsize=9, weight="bold")
bottom += val
ax.axhline(_HBM_PER_PE_GB, color="red", ls="--", lw=1.4,
label=f"HBM = {_HBM_PER_PE_GB} GB")
ax.set_xticks([0])
ax.set_xticklabels(["per-PE HBM"], fontsize=10)
ax.set_ylabel("GB per PE")
ax.set_ylim(0, _HBM_PER_PE_GB * 1.10)
ax.set_title(
f"Per-PE HBM budget\n"
f"weights {_WEIGHTS_GB:.2f} GB + KV = {_HBM_PER_PE_GB} GB",
fontsize=10,
)
ax.grid(axis="y", ls=":", alpha=0.5)
ax.legend(loc="upper right", fontsize=7.5, framealpha=0.92)
def _plot_memory(ax) -> None:
vals_gb = [kv_per_pe_bytes(c, _HEADLINE_S_KV) / (1 << 30) for c in _CASES]
labels = [_CASE_LABEL[c] for c in _CASES]
colors = [_CASE_COLOR[c] for c in _CASES]
x = list(range(len(_CASES)))
bars = ax.bar(x, vals_gb, color=colors, width=0.65)
ax.set_xticks(x)
ax.set_xticklabels(labels, fontsize=9)
ax.set_ylabel("KV bytes per PE (GB, log)")
ax.set_yscale("log")
ax.set_title(
f"Per-PE KV memory at S_kv = {_HEADLINE_S_KV:,} tokens "
f"(across {_N_LAYERS} layers, FP16)",
fontsize=11,
)
ax.grid(axis="y", ls=":", alpha=0.5, which="both")
ax.axhline(_HEADROOM_GB, color="red", ls="--", lw=1.4,
label=f"KV headroom = {_HEADROOM_GB} GB / PE")
for bar, v_gb in zip(bars, vals_gb):
v_bytes = v_gb * (1 << 30)
fits = v_gb <= _HEADROOM_GB
ax.text(bar.get_x() + bar.get_width() / 2,
v_gb * 1.10,
_fmt_bytes(v_bytes) + (" ✓" if fits else " ✗"),
ha="center", va="bottom", fontsize=9,
color="green" if fits else "red", weight="bold")
ax.legend(loc="upper right", fontsize=9)
def _plot_comm(ax, *, mode: str = "analytical") -> None:
"""Per-PE comm panel.
mode = "analytical": single solid bars from per_token_bytes formula.
mode = "paired" : analytical (solid) + simulator-measured
(hatched) side-by-side per case, when the
measurement JSON is available.
"""
labels = [_CASE_LABEL[c] for c in _CASES]
wo_mb_list: list[float] = []
ffn_mb_list: list[float] = []
attn_mb: list[float] = []
for c in _CASES:
wo, ffn, attn = per_token_bytes(c, _HEADLINE_S_KV)
wo_mb_list.append(wo / (1 << 20))
ffn_mb_list.append(ffn / (1 << 20))
attn_mb.append(attn / (1 << 20))
measured = _load_measured() if mode == "paired" else None
paired = measured is not None
source_tag = "analytical (solid) vs simulator-measured (hatched)" \
if paired else "analytical"
n_cases = len(_CASES)
x = list(range(n_cases))
bar_w = 0.36 if paired else 0.65
x_ana = [xi - bar_w / 2 for xi in x] if paired else x
x_meas = [xi + bar_w / 2 for xi in x] if paired else None
# Analytical bars (solid).
ax.bar(x_ana, wo_mb_list, width=bar_w,
color=_WO_COLOR, edgecolor="black")
ax.bar(x_ana, ffn_mb_list, width=bar_w, bottom=wo_mb_list,
color=_FFN_COLOR, edgecolor="black")
bottoms_attn = [w + f for w, f in zip(wo_mb_list, ffn_mb_list)]
ax.bar(x_ana, attn_mb, width=bar_w, bottom=bottoms_attn,
color=_ATTN_COLOR, edgecolor="black")
# Measured bars (hatched) — same Wo+FFN base, attn from op_log.
meas_attn_mb: list[float] = []
if paired:
for i, c in enumerate(_CASES):
meas_total = measured.get(c, 0) / (1 << 20)
meas_attn_mb.append(
max(meas_total - wo_mb_list[i] - ffn_mb_list[i], 0.0))
ax.bar(x_meas, wo_mb_list, width=bar_w,
color=_WO_COLOR, edgecolor="black",
hatch="///", alpha=0.85)
ax.bar(x_meas, ffn_mb_list, width=bar_w, bottom=wo_mb_list,
color=_FFN_COLOR, edgecolor="black",
hatch="///", alpha=0.85)
ax.bar(x_meas, meas_attn_mb, width=bar_w, bottom=bottoms_attn,
color=_ATTN_COLOR, edgecolor="black",
hatch="///", alpha=0.85)
ax.set_xticks(x)
ax.set_xticklabels(labels, fontsize=9)
ax.set_ylabel("Comm per token per PE (MB, log)")
ax.set_yscale("log")
title = (
f"Per-PE communication at S_kv = {_HEADLINE_S_KV:,} tokens "
f"(decode S_q=1, B=1; {_N_LAYERS} layers) — {source_tag}"
)
if paired:
title += "\n(simulator measured at S_kv = 8K; " \
"partial-score AR scaled ×128 to S_kv = 1M)"
ax.set_title(title, fontsize=10)
ax.grid(axis="y", ls=":", alpha=0.5, which="both")
totals_ana = [wo_mb_list[i] + ffn_mb_list[i] + attn_mb[i]
for i in range(n_cases)]
ymax = max(totals_ana)
if paired:
ymax = max(ymax, max(
(measured.get(c, 0) / (1 << 20)) for c in _CASES))
ax.set_ylim(top=ymax * 22)
for i, c in enumerate(_CASES):
ana_bytes = totals_ana[i] * (1 << 20)
if paired:
meas_bytes = measured.get(c, 0)
top_y = max(totals_ana[i], meas_bytes / (1 << 20))
label = (f"ana: {_fmt_bytes(ana_bytes)}\n"
f"sim: {_fmt_bytes(meas_bytes)}")
else:
top_y = totals_ana[i]
label = _fmt_bytes(ana_bytes)
ax.text(x[i], top_y * 1.5, label,
ha="center", va="bottom",
fontsize=8 if paired else 9,
weight="bold", linespacing=1.05)
# Attention-time AR descriptor — placed at the attention-segment
# midpoint, centered between the analytical and measured bars so
# it visually spans both. Coloured the same as the attention bar
# (no background box) so it lives within the case's bar zone.
# Wrapped narrow so each line fits inside the bar-pair width.
wrapped = textwrap.fill(_ATTN_DESC[c], width=14)
if attn_mb[i] > 0:
mid = bottoms_attn[i] + attn_mb[i] / 2
ax.text(x[i], mid, wrapped,
ha="center", va="center",
fontsize=7 if paired else 8,
color="black", weight="bold",
linespacing=1.0)
else:
ax.text(x[i], totals_ana[i] * 0.4, wrapped,
ha="center",
fontsize=7 if paired else 8,
color="grey", style="italic",
linespacing=1.0)
legend_handles = [
mpatches.Patch(facecolor=_WO_COLOR, edgecolor="black",
label=f"Wo AR (× {_N_LAYERS} layers)"),
mpatches.Patch(facecolor=_FFN_COLOR, edgecolor="black",
label=f"FFN AR (× {_N_LAYERS} layers)"),
mpatches.Patch(facecolor=_ATTN_COLOR, edgecolor="black",
label="Attn-time collective"),
]
if paired:
legend_handles.append(
mpatches.Patch(facecolor="white", edgecolor="black",
hatch="///", label="simulator-measured"))
ax.legend(handles=legend_handles, loc="upper right", fontsize=8,
framealpha=0.92)
def main() -> Path:
_OUT_DIR.mkdir(parents=True, exist_ok=True)
# (a) Per-PE HBM budget — standalone PNG.
fig_b, ax_b = plt.subplots(figsize=(4.0, 6.0))
_plot_budget(ax_b)
fig_b.tight_layout()
out_b = _OUT_DIR / "gqa_long_ctx_6cases_hbm_budget.png"
fig_b.savefig(out_b, dpi=150)
plt.close(fig_b)
print(f"wrote {out_b}")
# (b) Combined 3-panel summary — HBM budget + KV memory + comm.
fig = plt.figure(figsize=(22.0, 6.5))
gs = fig.add_gridspec(1, 3, width_ratios=[0.7, 1.6, 1.6], wspace=0.22)
ax_b2 = fig.add_subplot(gs[0, 0])
ax_m = fig.add_subplot(gs[0, 1])
ax_c = fig.add_subplot(gs[0, 2])
_plot_budget(ax_b2)
_plot_memory(ax_m)
_plot_comm(ax_c)
fig.tight_layout()
out = _OUT_DIR / "gqa_long_ctx_6cases_summary.png"
fig.savefig(out, dpi=150)
plt.close(fig)
print(f"wrote {out}")
# (c) 2-panel companion (analytical only).
fig2 = plt.figure(figsize=(18.0, 6.5))
gs2 = fig2.add_gridspec(1, 2, width_ratios=[1.0, 1.0], wspace=0.18)
ax_m2 = fig2.add_subplot(gs2[0, 0])
ax_c2 = fig2.add_subplot(gs2[0, 1])
_plot_memory(ax_m2)
_plot_comm(ax_c2, mode="analytical")
fig2.tight_layout()
out2 = _OUT_DIR / "gqa_long_ctx_6cases_memory_comm_analytical.png"
fig2.savefig(out2, dpi=150)
plt.close(fig2)
print(f"wrote {out2}")
# (d) 2-panel companion — analytical vs simulator-measured paired.
fig3 = plt.figure(figsize=(19.0, 6.5))
gs3 = fig3.add_gridspec(1, 2, width_ratios=[1.0, 1.0], wspace=0.18)
ax_m3 = fig3.add_subplot(gs3[0, 0])
ax_c3 = fig3.add_subplot(gs3[0, 1])
_plot_memory(ax_m3)
_plot_comm(ax_c3, mode="paired")
fig3.tight_layout()
out3 = _OUT_DIR / "gqa_long_ctx_6cases_memory_comm_paired.png"
fig3.savefig(out3, dpi=150)
plt.close(fig3)
print(f"wrote {out3}")
# Paper-ready table to stdout.
print()
print(f" {'Case':<27} {'KV/tok·PE':>12} {'KV @ 1M':>11} "
f"{'Max S_kv':>11} {'Comm @ 1M':>12}")
print(" " + "-" * 80)
for c in _CASES:
kv_per_tok_pe = _KV_PER_TOK_BYTES / _DIVISOR[c]
mem_1m = kv_per_pe_bytes(c, _HEADLINE_S_KV)
max_s = max_s_kv(c)
comm_1m = sum(per_token_bytes(c, _HEADLINE_S_KV))
label = _CASE_LABEL[c].replace(chr(10), " ")
print(f" {label:<27} "
f"{kv_per_tok_pe / 1024:>9.3f} KB "
f"{_fmt_bytes(mem_1m):>11} "
f"{_fmt_tokens(max_s):>11} "
f"{_fmt_bytes(comm_1m):>12}")
return out
if __name__ == "__main__":
main()
scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py
+248
View File
@@ -0,0 +1,248 @@
"""Comparative figures for milestone-gqa-decode-long-ctx-4cases.
Reads sweep_decode.json (emitted by the milestone-1h-gqa bench) and
writes four PNGs into the same bench-output dir
(src/kernbench/benches/1H_milestone_output/gqa/long_ctx/):
gqa_decode_long_ctx_6cases_latency.png end-to-end latency per case
gqa_decode_long_ctx_6cases_traffic.png ipcq/dma op-count breakdown
gqa_decode_long_ctx_6cases_memory.png per-PE KV bytes per case
gqa_decode_long_ctx_6cases_parallelism.png per-PE S_local (compute work)
Filename still says "4cases" for backwards compat, but the script now
covers all SIX kv-sharding strategies from the analytical chart
(`gqa_4cases_summary.png`) — the original 4 plus the two new
d_head-TP variants:
Case 1 Cube-Repl × PE-repl (PE-TP doesn't shard KV)
Case 2 Cube-SP × PE-repl
Case 3 Cube-Repl × PE-SP
Case 4 Cube-SP × PE-TP-d_head ← NEW
Case 5 Cube-TP-d_head × PE-SP ← NEW
Case 6 ★ Cube-SP × PE-SP (Pareto-best)
Run (after the bench):
GQA_DECODE_LONG_CTX_4CASES_RUN=1 python -m kernbench.cli.main run \\
--bench milestone-gqa-decode-long-ctx-4cases --topology topology.yaml
python scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py
"""
from __future__ import annotations
import json
from pathlib import Path
import matplotlib
matplotlib.use("Agg")
import matplotlib.pyplot as plt # noqa: E402
_REPO_ROOT = Path(__file__).resolve().parents[2]
# Sweep JSON + PNGs live together under the bench output dir.
_FIG_DIR = (
_REPO_ROOT / "src" / "kernbench" / "benches"
/ "1H_milestone_output" / "gqa" / "long_ctx"
)
_SWEEP_JSON = _FIG_DIR / "sweep_decode.json"
# Panel name → (short label, case ordinal, accent flag) using the
# analytical chart's memory-descending ordering. PE-TP doesn't shard
# KV memory, so the cube_repl_pe_tp panel maps to Case 1 (no
# sharding, KV-wise) and cube_sp_pe_tp panel maps to Case 2.
_NORMAL, _OVERFLOW, _PARETO = "normal", "overflow", "pareto"
_CASE_INFO = {
# panel name label ord flag
"single_kv_group_decode_long_ctx_gqa_cube_repl_pe_tp": ("Case 1\nCube-Repl × PE-repl", 1, _OVERFLOW),
"single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp": ("Case 2\nCube-SP × PE-repl", 2, _OVERFLOW),
"single_kv_group_decode_long_ctx_gqa_cube_repl_pe_sp": ("Case 3\nCube-Repl × PE-SP", 3, _OVERFLOW),
"single_kv_group_decode_long_ctx_gqa_cube_sp_pe_tp_dhead": ("Case 4\nCube-SP × PE-TP-d_head", 4, _NORMAL),
"single_kv_group_decode_long_ctx_gqa_cube_tp_dhead_pe_sp": ("Case 5\nCube-TP-d_head × PE-SP", 5, _NORMAL),
"single_kv_group_decode_long_ctx_gqa_cube_sp_pe_sp": ("Case 6 ★\nCube-SP × PE-SP", 6, _PARETO),
}
# Bar fill colour per flag (used by every panel).
_FLAG_COLOR = {
_NORMAL: "#888888", # neutral grey
_OVERFLOW: "#c0504d", # red — fails the per-PE HBM budget
_PARETO: "#3b6ea5", # blue — Pareto-best
}
def _load() -> list[dict]:
return json.loads(_SWEEP_JSON.read_text())["rows"]
def _sorted_by_case(rows: list[dict]) -> list[dict]:
return sorted(rows, key=lambda r: _CASE_INFO[r["panel"]][1])
def _bar_colors(rows: list[dict]) -> list[str]:
return [_FLAG_COLOR[_CASE_INFO[r["panel"]][2]] for r in rows]
def _plot_latency(rows: list[dict]) -> Path:
rows = _sorted_by_case(rows)
labels = [_CASE_INFO[r["panel"]][0] for r in rows]
lat_us = [r["latency_ns"] / 1e3 for r in rows]
fig, ax = plt.subplots(figsize=(12.0, 4.8))
bars = ax.bar(labels, lat_us, color=_bar_colors(rows), width=0.6)
ax.set_ylabel("end-to-end latency (µs)")
ax.set_title(
"Long-context decode 6-cases — end-to-end latency per case\n"
"LLaMA-3.1-70B single-KV-head group (8 cubes × 8 PEs)"
)
ax.bar_label(bars, fmt="%.1f", padding=3, fontsize=9)
ax.grid(axis="y", ls=":", alpha=0.5)
ax.set_ylim(0, max(lat_us) * 1.15)
fig.tight_layout()
out = _FIG_DIR / "gqa_decode_long_ctx_6cases_latency.png"
fig.savefig(out, dpi=150)
plt.close(fig)
return out
def _plot_traffic(rows: list[dict]) -> Path:
rows = _sorted_by_case(rows)
labels = [_CASE_INFO[r["panel"]][0] for r in rows]
x = list(range(len(rows)))
keys = ["ipcq_copy_count", "dma_read_count", "dma_write_count"]
disp = ["IPCQ copy", "DMA read", "DMA write"]
colors = ["#c0504d", "#9bbb59", "#8064a2"]
w = 0.25
fig, ax = plt.subplots(figsize=(11.0, 4.5))
for i, (k, d, c) in enumerate(zip(keys, disp, colors)):
vals = [r["op_log_summary"][k] for r in rows]
ax.bar([xi + (i - 1) * w for xi in x], vals, width=w, label=d, color=c)
ax.set_xticks(list(x))
ax.set_xticklabels(labels, fontsize=9)
ax.set_ylabel("op count")
ax.set_title("Long-context decode 6-cases — op-count breakdown per case")
ax.legend(fontsize=9)
ax.grid(axis="y", ls=":", alpha=0.5)
fig.tight_layout()
out = _FIG_DIR / "gqa_decode_long_ctx_6cases_traffic.png"
fig.savefig(out, dpi=150)
plt.close(fig)
return out
def _s_local_per_pe(panel: str, *, S_kv: int, C: int, P: int) -> int:
"""S_local (token count) each PE attends over locally.
cube_repl_pe_tp (Case 1): S_kv (no sharding, KV-wise)
cube_sp_pe_tp (Case 2): S_kv / C (cube splits S_kv, PEs replicate)
cube_repl_pe_sp (Case 3): S_kv / P
cube_sp_pe_tp_dhead (Case 4): S_kv / C (cube splits S_kv, PE splits d_head)
cube_tp_dhead_pe_sp (Case 5): S_kv / P (cube splits d_head, PE splits S_kv)
cube_sp_pe_sp (Case 6 ★): S_kv / (C·P)
"""
cube_splits_s = "cube_sp" in panel
pe_splits_s = "pe_sp" in panel
S_per_cube = S_kv // C if cube_splits_s else S_kv
return S_per_cube // P if pe_splits_s else S_per_cube
def _d_head_per_pe(panel: str, *, d_head: int, C: int, P: int) -> int:
"""d_head dims each PE owns (Cases 4 and 5 shard d_head)."""
if "cube_tp_dhead" in panel: # Case 5: cube shards d_head
return d_head // C
if "pe_tp_dhead" in panel: # Case 4: PE shards d_head
return d_head // P
return d_head # Cases 1, 2, 3, 6: full d_head per PE
def _active_pe_count(panel: str, *, C: int, P: int) -> int:
"""Number of PEs doing non-idle attention work.
cube_repl_pe_tp (Case 1): 1 (PE-TP idle for B=1; only one PE works)
cube_sp_pe_tp (Case 2): C (PE 0 of each cube; 7 PEs idle per cube)
cube_repl_pe_sp (Case 3): C·P (all PEs busy, cube-side redundant)
cube_sp_pe_tp_dhead (Case 4): C·P (PE shards d_head — all 64 active)
cube_tp_dhead_pe_sp (Case 5): C·P (PE shards S_kv — all active)
cube_sp_pe_sp (Case 6 ★): C·P (all 64 PEs doing unique work)
"""
if "cube_repl" in panel and "pe_tp" in panel and "dhead" not in panel:
return 1
if "cube_sp" in panel and "pe_tp" in panel and "dhead" not in panel:
return C
return C * P
def _kv_bytes_per_pe(panel: str, *, S_kv: int, h_kv: int,
d_head: int, C: int, P: int) -> int:
"""KV bytes a single PE references (K + V, f16, 2 B/elem)."""
s_local = _s_local_per_pe(panel, S_kv=S_kv, C=C, P=P)
d_local = _d_head_per_pe(panel, d_head=d_head, C=C, P=P)
return 2 * s_local * h_kv * d_local * 2
def _plot_memory(rows: list[dict]) -> Path:
"""Per-PE KV bytes — Case 6 ★ wins (64-way split)."""
rows = _sorted_by_case(rows)
labels = [_CASE_INFO[r["panel"]][0] for r in rows]
mib_per_pe = [
_kv_bytes_per_pe(
r["panel"], S_kv=r["S_kv"], h_kv=r["h_kv"],
d_head=r["d_head"], C=r["C"], P=r["P"],
) / (1024 * 1024)
for r in rows
]
fig, ax = plt.subplots(figsize=(12.0, 4.8))
bars = ax.bar(labels, mib_per_pe, color=_bar_colors(rows), width=0.6)
ax.set_ylabel("KV bytes per PE (MiB, K + V, f16)")
ax.set_title(
"Long-context decode 6-cases — KV memory per PE\n"
"(one KV-head group; per-layer, per-token state)"
)
ax.bar_label(bars, fmt="%.3f", padding=3, fontsize=9)
ax.grid(axis="y", ls=":", alpha=0.5)
ax.set_ylim(0, max(mib_per_pe) * 1.15)
fig.tight_layout()
out = _FIG_DIR / "gqa_decode_long_ctx_6cases_memory.png"
fig.savefig(out, dpi=150)
plt.close(fig)
return out
def _plot_parallelism(rows: list[dict]) -> Path:
"""Total active PE-token compute load — exposes redundant-work cases."""
rows = _sorted_by_case(rows)
labels = [_CASE_INFO[r["panel"]][0] for r in rows]
total_work = [
_active_pe_count(r["panel"], C=r["C"], P=r["P"])
* _s_local_per_pe(r["panel"], S_kv=r["S_kv"], C=r["C"], P=r["P"])
for r in rows
]
fig, ax = plt.subplots(figsize=(12.0, 4.8))
bars = ax.bar(labels, total_work, color=_bar_colors(rows), width=0.6)
ax.set_ylabel("active-PE × S_local (PE-tokens; lower ⇒ less wasted work)")
ax.set_title(
"Long-context decode 6-cases — total compute load across active PEs\n"
"(Case 3 replicates KV across 8 cubes → 8× wasted PE-tokens; "
"Case 6 ★ is fully parallel without replication)"
)
ax.bar_label(bars, fmt="%d", padding=3, fontsize=9)
ax.grid(axis="y", ls=":", alpha=0.5)
ax.set_ylim(0, max(total_work) * 1.15)
fig.tight_layout()
out = _FIG_DIR / "gqa_decode_long_ctx_6cases_parallelism.png"
fig.savefig(out, dpi=150)
plt.close(fig)
return out
def main() -> None:
rows = _load()
_FIG_DIR.mkdir(parents=True, exist_ok=True)
p1 = _plot_latency(rows)
p2 = _plot_traffic(rows)
p3 = _plot_memory(rows)
p4 = _plot_parallelism(rows)
print(f"wrote {p1}")
print(f"wrote {p2}")
print(f"wrote {p3}")
print(f"wrote {p4}")
if __name__ == "__main__":
main()
scripts/paper/paper_plot_gqa_decode_long_ctx_composite.py
+146
View File
@@ -0,0 +1,146 @@
"""Comparative figure for the Case-6 composite-command decode study.
Reads sweep_decode_composite.json (emitted by milestone-1h-gqa, sweep
``composite``) and writes one two-panel PNG into the bench-output dir:
gqa_decode_long_ctx_composite.png
Left — end-to-end decode latency (µs) vs context length, per command
form (primitive / composite / composite_extended).
Right — PE_CPU command count vs context length: the hand-tiled
primitive kernel issues O(n_tiles) commands (rises with
context), while the coarse composite forms issue O(1) and
*saturate* — PE_SCHEDULER absorbs the per-tile fan-out.
The x-axis is the global context length S_kv; each PE owns
S_local = S_kv/(C·P=64) tokens, so at the 1M production point every PE
runs Q·Kᵀ of (G·T_q, d_head)·(d_head, 16384) and P·V of
(G·T_q, 16384)·(16384, d_head).
Run (after the bench):
GQA_1H_RUN=1 GQA_1H_SWEEPS=composite python -m kernbench.cli.main run \\
--bench milestone-1h-gqa --topology topology.yaml
python scripts/paper/paper_plot_gqa_decode_long_ctx_composite.py
"""
from __future__ import annotations
import json
from pathlib import Path
import matplotlib
matplotlib.use("Agg")
import matplotlib.pyplot as plt # noqa: E402
_REPO_ROOT = Path(__file__).resolve().parents[2]
_FIG_DIR = (
_REPO_ROOT / "src" / "kernbench" / "benches"
/ "1H_milestone_output" / "gqa" / "long_ctx"
)
_SWEEP_JSON = _FIG_DIR / "sweep_decode_composite.json"
_PAPER_FIG_DIR = (
_REPO_ROOT / "docs" / "report" / "1H-codesign-paper" / "figures"
)
_N_RANKS = 64 # C·P for the Case-6 64-way split.
# variant key → (display label, colour, marker)
_VARIANT_STYLE = {
"primitive": ("primitive (tl.dot, hand-tiled)", "#c0504d", "o"),
"composite": ("composite GEMM", "#3b6ea5", "s"),
"composite_extended": ("composite + softmax_merge", "#4f8a4f", "^"),
}
_ORDER = ("primitive", "composite", "composite_extended")
def _load() -> dict:
return json.loads(_SWEEP_JSON.read_text())
def _series(rows: list[dict], variant: str, key: str):
"""Sorted (S_kv, value) series for a variant, skipping null values
(latency is only measured over the tractable S_kv subset)."""
pts = sorted(
((r["S_kv"], r[key]) for r in rows
if r["variant"] == variant and r.get(key) is not None),
key=lambda t: t[0],
)
return [p[0] for p in pts], [p[1] for p in pts]
def _xticklabels(s_kvs: list[int]) -> list[str]:
out = []
for s in s_kvs:
if s >= 1 << 20:
out.append(f"{s // (1 << 20)}M")
else:
out.append(f"{s // 1024}K")
return out
def main() -> None:
sweep = _load()
rows = sweep["rows"]
s_kv_op = sweep["s_kv_opcount"]
s_kv_lat = sweep["s_kv_latency"]
fig, (ax_lat, ax_cmd) = plt.subplots(1, 2, figsize=(13.0, 4.8))
for v in _ORDER:
label, color, marker = _VARIANT_STYLE[v]
xs, lat = _series(rows, v, "latency_ns")
ax_lat.plot(xs, [y / 1e3 for y in lat], marker=marker,
color=color, label=label, lw=2)
xs, cmds = _series(rows, v, "pe_cpu_cmd_count")
ax_cmd.plot(xs, cmds, marker=marker, color=color, label=label, lw=2)
ax_lat.set_xticks(s_kv_lat)
ax_lat.set_xticklabels(_xticklabels(s_kv_lat))
ax_cmd.set_xticks(s_kv_op)
ax_cmd.set_xticklabels(_xticklabels(s_kv_op), fontsize=8)
for ax in (ax_lat, ax_cmd):
ax.set_xscale("log", base=2)
ax.set_xlabel(
r"context length $S_{kv}$ "
r"($S_{\mathrm{local}}=S_{kv}/64$ per PE)"
)
ax.grid(True, ls=":", alpha=0.5)
ax.legend(fontsize=9)
ax_lat.set_ylabel("end-to-end decode latency (µs)")
ax_lat.set_title(
"Case-6 decode latency per command form\n"
"(memory-bound: command form does not move the critical path)"
)
ax_cmd.set_ylabel("PE_CPU commands issued")
ax_cmd.set_title(
"PE_CPU command count per command form\n"
"(primitive O(n$_\\mathrm{tiles}$) rises; composite O(1) saturates)"
)
ax_cmd.axvline(1 << 20, color="#888", ls="--", lw=1, alpha=0.7)
ax_cmd.annotate("1M production\ncontext", xy=(1 << 20, 0),
xytext=(1 << 18, 0.6), fontsize=8,
textcoords=("data", "axes fraction"), ha="right",
color="#555")
fig.suptitle(
"Case-6 (Cube-SP × PE-SP) long-context decode — use of composite "
"commands\nLLaMA-3.1-70B single-KV-head group (8 cubes × 8 PEs), "
"$T_q{=}1$",
fontsize=11,
)
fig.tight_layout(rect=(0, 0, 1, 0.94))
out = _FIG_DIR / "gqa_decode_long_ctx_composite.png"
fig.savefig(out, dpi=150)
plt.close(fig)
print(f"wrote {out}")
# Mirror into the paper figures dir (derived artifact).
if _PAPER_FIG_DIR.is_dir():
dst = _PAPER_FIG_DIR / out.name
dst.write_bytes(out.read_bytes())
print(f"copied {dst}")
if __name__ == "__main__":
main()
scripts/paper/paper_plot_gqa_decode_streaming.py
+124
View File
@@ -0,0 +1,124 @@
"""Comparative figure for the memory-bound decode-streaming composite study.
Reads sweep_decode_streaming.json (emitted by milestone-1h-gqa, sweep
``decode_streaming``) and writes one two-panel PNG:
gqa_decode_streaming.png
Left — end-to-end single-rank decode latency (µs) vs per-rank context.
Right — achieved HBM bandwidth (GB/s) vs context, against the
256 GB/s per-rank roofline.
The memory-bound mirror of the compute-bound prefill figure. With T_q=1
the GEMMs are skinny (M=8) and the kernel is bound by streaming the KV
cache. Isolating a single rank (no inter-CUBE reduce) reveals what the
64-way Case-6 decode masks: the composite command still wins, not by
feeding the MAC array but by keeping the DMA pipeline full — its
scheduler-streamed concurrent tile DMAs extract ~230 GB/s (near the
256 GB/s roofline) while the primitive kernel's blocking tl.dot serializes
one tile DMA at a time and plateaus at ~166 GB/s. That bandwidth gap is a
~25-28 % latency win that grows nowhere near prefill's compute-bound
margin but is decidedly not zero.
Run (after the bench):
GQA_1H_RUN=1 GQA_1H_SWEEPS=decode_streaming python -m kernbench.cli.main \\
run --bench milestone-1h-gqa --topology topology.yaml
python scripts/paper/paper_plot_gqa_decode_streaming.py
"""
from __future__ import annotations
import json
from pathlib import Path
import matplotlib
matplotlib.use("Agg")
import matplotlib.pyplot as plt # noqa: E402
_REPO_ROOT = Path(__file__).resolve().parents[2]
_FIG_DIR = (
_REPO_ROOT / "src" / "kernbench" / "benches"
/ "1H_milestone_output" / "gqa" / "long_ctx"
)
_SWEEP_JSON = _FIG_DIR / "sweep_decode_streaming.json"
_PAPER_FIG_DIR = (
_REPO_ROOT / "docs" / "report" / "1H-codesign-paper" / "figures"
)
# Per-rank HBM roofline: 8 pseudo-channels × 32 GB/s (topology.yaml
# hbm_ctrl.num_pcs / pc_bw_gbs; = pe_dma_to_noc_bw_gbs).
_PEAK_HBM_GBS = 256.0
_VARIANT_STYLE = {
"primitive": ("primitive (tl.dot, hand-tiled)", "#c0504d", "o"),
"composite": ("composite GEMM", "#3b6ea5", "s"),
"composite_extended": ("composite + softmax_merge", "#4f8a4f", "^"),
}
_ORDER = ("primitive", "composite", "composite_extended")
def _ctx_label(c: int) -> str:
return f"{c // 1024}K" if c >= 1024 else str(c)
def _series(rows, variant, key):
pts = sorted(((r["s_kv"], r[key]) for r in rows
if r["variant"] == variant), key=lambda t: t[0])
return [p[0] for p in pts], [p[1] for p in pts]
def main() -> None:
sweep = json.loads(_SWEEP_JSON.read_text())
rows = sweep["rows"]
ctxs = sweep["s_kv_points"]
fig, (ax_lat, ax_bw) = plt.subplots(1, 2, figsize=(13.0, 4.8))
for v in _ORDER:
label, color, marker = _VARIANT_STYLE[v]
xs, lat = _series(rows, v, "latency_ns")
ax_lat.plot(xs, [y / 1e3 for y in lat], marker=marker,
color=color, label=label, lw=2)
xs, bw = _series(rows, v, "achieved_bw_gbs")
ax_bw.plot(xs, bw, marker=marker, color=color, label=label, lw=2)
for ax in (ax_lat, ax_bw):
ax.set_xscale("log", base=2)
ax.set_xticks(ctxs)
ax.set_xticklabels([_ctx_label(c) for c in ctxs])
ax.set_xlabel(r"per-rank context length $S_{kv}$ ($T_q{=}1$)")
ax.grid(True, ls=":", alpha=0.5)
ax.legend(fontsize=9)
ax_lat.set_ylabel("end-to-end decode latency (µs)")
ax_lat.set_title("Single-rank memory-bound decode latency per command form")
ax_bw.set_ylabel("achieved HBM bandwidth (GB/s)")
ax_bw.set_title(
"HBM bandwidth — composite keeps the DMA pipe full; primitive plateaus"
)
ax_bw.axhline(_PEAK_HBM_GBS, color="#888", ls="--", lw=1, alpha=0.7)
ax_bw.text(ctxs[0], _PEAK_HBM_GBS - 8, "256 GB/s roofline",
fontsize=8, color="#555", va="top")
ax_bw.set_ylim(0, _PEAK_HBM_GBS * 1.08)
fig.suptitle(
"Memory-bound decode streaming — use of composite commands\n"
"single-rank, GQA single-KV-head group ($h_q{=}8$, $d_{\\text{head}}"
"{=}128$); $M{=}8$ skinny, KV-streaming-bound",
fontsize=11,
)
fig.tight_layout(rect=(0, 0, 1, 0.92))
out = _FIG_DIR / "gqa_decode_streaming.png"
fig.savefig(out, dpi=150)
plt.close(fig)
print(f"wrote {out}")
if _PAPER_FIG_DIR.is_dir():
dst = _PAPER_FIG_DIR / out.name
dst.write_bytes(out.read_bytes())
print(f"copied {dst}")
if __name__ == "__main__":
main()
scripts/paper/paper_plot_gqa_kv_sharding_diagram.py
+324
View File
@@ -0,0 +1,324 @@
"""6-case KV-sharding tensor diagram (the slide-13 PNG export).
Flat 2-D rectangles, one per sharding case, with:
Y axis = S_kv (vertical) — Cube-SP / PE-SP slice it
X axis = d_head (horizontal) — Cube-TP-d_head / PE-TP-d_head slice it
Drops the batch axis entirely (decode: B = 1, T_q = 1). Same case set
and visual encoding as slide 13 of GQA_full_deck.pptx; matplotlib
renders it cleanly so the PNG sits next to the other GQA summary
artifacts in 1H_milestone_output/gqa/long_ctx/.
"""
from __future__ import annotations
from pathlib import Path
import matplotlib.patches as mpatches
import matplotlib.pyplot as plt
_C = 8
_P = 8
_GROUP_FILLS = [
"#A5D8FF", "#B2F2BB", "#FFD8A8", "#FFC9C9",
"#D0BFFF", "#99E9F2", "#FCC2D7", "#FFEC99",
]
_ACC = {
"red": "#E03131",
"orange": "#FD7E14",
"blue": "#1C7ED6",
"green": "#37B24D",
}
# (label, accent, kv, comm, overflow, encoding-flags, axis-spec)
# y_split = 8 horizontal Y bands (Cube-SP on S_kv)
# x_split = 8 vertical X bands (Cube-TP-d_head)
# pe_y = 7 fine horizontal dividers within each Y band
# pe_x = 7 fine vertical dividers within each X band
# axes = small annotation under the chip naming the axes
# that the cube/PE actually shard, so the reader can
# parse Case 5 (where cube colour fills run X instead
# of Y, breaking the visual symmetry of the rest).
_CASES = [
dict(label="Case 1\nCube-Repl / PE-repl", accent=_ACC["red"],
kv="40 GB", comm="1.2 MB", overflow=True,
y_split=False, x_split=False, pe_y=False, pe_x=False,
axes="Cube: replicated PE: replicated"),
dict(label="Case 2\nCube-SP / PE-repl", accent=_ACC["orange"],
kv="5 GB", comm="3.8 MB", overflow=True,
y_split=True, x_split=False, pe_y=False, pe_x=False,
axes="Cube → Y (S_kv) PE: replicated"),
dict(label="Case 3\nCube-Repl / PE-SP", accent=_ACC["orange"],
kv="5 GB", comm="3.8 MB", overflow=True,
y_split=False, x_split=False, pe_y=True, pe_x=False,
axes="Cube: replicated PE → Y (S_kv)"),
dict(label="Case 4\nCube-SP / PE-TP-d_head", accent=_ACC["blue"],
kv="640 MB", comm="166 MB", overflow=False,
y_split=True, x_split=False, pe_y=False, pe_x=True,
axes="Cube → Y (S_kv) PE → X (d_head)"),
dict(label="Case 5\nCube-TP-d_head / PE-SP", accent=_ACC["blue"],
kv="640 MB", comm="166 MB", overflow=False,
y_split=False, x_split=True, pe_y=True, pe_x=False,
axes="Cube → X (d_head) PE → Y (S_kv)"),
dict(label="Case 6 ★\nCube-SP / PE-SP", accent=_ACC["green"],
kv="640 MB", comm="6.2 MB", overflow=False,
y_split=True, x_split=False, pe_y=True, pe_x=False,
axes="Cube → Y (S_kv) PE → Y (S_kv)"),
]
_OUT_DIR = (
Path(__file__).resolve().parents[2]
/ "src" / "kernbench" / "benches"
/ "1H_milestone_output" / "gqa" / "long_ctx"
)
def _draw_panel(ax, cfg):
"""Draw one case's 2-D KV-tensor rectangle into a panel ax."""
ax.set_xlim(0, 1)
ax.set_ylim(1, 0) # Y points down (S_kv ↓)
ax.set_aspect("auto")
ax.set_xticks([])
ax.set_yticks([])
cube_repl = not cfg["y_split"] and not cfg["x_split"]
pe_repl = not cfg["pe_y"] and not cfg["pe_x"]
# Cube-level colour fill.
if cfg["y_split"] and not cfg["x_split"]:
# 8 horizontal Y bands.
for c in range(_C):
ax.add_patch(mpatches.Rectangle(
(0, c / _C), 1, 1 / _C,
facecolor=_GROUP_FILLS[c], edgecolor="black", linewidth=0.6))
ax.text(0.04, c / _C + 0.5 / _C, f"C{c}",
ha="left", va="center", fontsize=8,
fontweight="bold", color="#333")
elif cfg["x_split"] and not cfg["y_split"]:
# 8 vertical X bands.
for c in range(_C):
ax.add_patch(mpatches.Rectangle(
(c / _C, 0), 1 / _C, 1,
facecolor=_GROUP_FILLS[c], edgecolor="black", linewidth=0.6))
ax.text(c / _C + 0.5 / _C, 0.04, f"C{c}",
ha="center", va="top", fontsize=8,
fontweight="bold", color="#333")
else:
ax.add_patch(mpatches.Rectangle(
(0, 0), 1, 1,
facecolor="#F5F5F5", edgecolor="black", linewidth=0.8))
ax.text(0.5, 0.5, "× 8 cubes\nfull KV",
ha="center", va="center",
fontsize=10, fontweight="bold",
fontstyle="italic", color="#666")
# PE-level fine dividers — distinguished from cube boundaries by
# using a dashed style + slightly stronger contrast. This is what
# makes Case 5's PE-SP (horizontal lines across vertical cube
# bands) read as "different axis from the cubes" at a glance.
if cfg["pe_y"]:
outer = _C if cfg["y_split"] else 1
band = 1 / outer
for o in range(outer):
for p in range(1, _P):
y = o * band + band * p / _P
ax.axhline(y, color="#222", linewidth=0.8,
linestyle=(0, (3, 2)), alpha=0.75)
if cfg["pe_x"]:
outer = _C if cfg["x_split"] else 1
band = 1 / outer
for o in range(outer):
for p in range(1, _P):
x = o * band + band * p / _P
ax.axvline(x, color="#222", linewidth=0.8,
linestyle=(0, (3, 2)), alpha=0.75)
# Heavy outline on top.
ax.add_patch(mpatches.Rectangle(
(0, 0), 1, 1, facecolor="none",
edgecolor="black", linewidth=1.2))
# Replication badges — small text-only badges in the corners of
# the rectangle, no ghost-card stacking (which mis-reads as a
# larger enclosing tensor).
badges: list[str] = []
if cube_repl:
badges.append("× 8 cube copies")
if pe_repl and (cfg["y_split"] or cfg["x_split"]):
# Cube is sharded but PEs in each cube replicate that shard.
badges.append("× 8 PEs / cube replicate")
elif pe_repl and cube_repl:
# Both replicated — PE replication adds to the cube one.
badges.append("× 8 PEs / cube replicate")
if badges:
ax.text(0.98, 0.02, "\n".join(badges),
ha="right", va="top", fontsize=7,
fontweight="bold", color="#444",
fontstyle="italic",
bbox=dict(facecolor="white", edgecolor="#888",
boxstyle="round,pad=0.20", linewidth=0.5))
def _make_table_png() -> Path:
"""Slide-14 companion table: per-PE memory + comm for all 6 cases."""
headers = ["Case", "Sharding", "KV / PE", "Fit",
"Comm/tok\n(analytical)", "Notes"]
rows = [
("Case 1", "Cube-Repl · PE-repl", "40 GB", "✗",
"1.2 MB",
"no sharding —\nfull KV on every PE"),
("Case 2", "Cube-SP · PE-repl", "5 GB", "✗",
"3.8 MB",
"cube-axis\nsharded only"),
("Case 3", "Cube-Repl · PE-SP", "5 GB", "✗",
"3.8 MB",
"PE-axis\nsharded only"),
("Case 4", "Cube-SP · PE-TP-d_head", "640 MB", "✓",
"166 MB",
"d_head split intra-cube\npartial-score AR ∝ S_kv"),
("Case 5", "Cube-TP-d_head · PE-SP", "640 MB", "✓",
"166 MB",
"d_head split inter-cube\npartial-score AR on UCIe"),
("Case 6 ★", "Cube-SP · PE-SP", "640 MB", "✓",
"6.2 MB",
"S_kv split both axes\n(m,ℓ,O) AR only"),
]
accents = [_ACC["red"], _ACC["orange"], _ACC["orange"],
_ACC["blue"], _ACC["blue"], _ACC["green"]]
fig, ax = plt.subplots(figsize=(15.0, 5.0))
ax.set_axis_off()
cell_data = [headers] + [list(r) for r in rows]
tbl = ax.table(cellText=cell_data,
colWidths=[0.07, 0.20, 0.09, 0.05, 0.14, 0.28],
cellLoc="center", loc="center")
tbl.auto_set_font_size(False)
tbl.set_fontsize(10.5)
tbl.scale(1.0, 2.4)
n_cols = len(headers)
n_rows = len(rows) + 1 # +1 header
# Header styling.
for ci in range(n_cols):
cell = tbl[(0, ci)]
cell.set_facecolor("#1F4E79")
cell.set_text_props(color="white", weight="bold")
cell.set_edgecolor("#1F4E79")
# Body styling.
for ri, row in enumerate(rows, start=1):
is_pareto = row[0].endswith("★")
row_fill = "#E8F5E9" if is_pareto else (
"white" if ri % 2 == 1 else "#F5F5F7")
# Case-name cell uses accent.
case_cell = tbl[(ri, 0)]
case_cell.set_facecolor(accents[ri - 1])
case_cell.set_text_props(color="white", weight="bold")
# Remaining cells.
for ci in range(1, n_cols):
cell = tbl[(ri, ci)]
cell.set_facecolor(row_fill)
txt_kwargs = {"weight": "bold" if is_pareto else "normal",
"color": "#333"}
if ci == 2: # KV / PE
txt_kwargs["color"] = (
"#C62828" if row[3] == "✗" else "#2E7D32")
txt_kwargs["weight"] = "bold"
if ci == 3: # Fit
txt_kwargs["color"] = (
"#C62828" if row[3] == "✗" else "#2E7D32")
txt_kwargs["weight"] = "bold"
cell.set_text_props(**txt_kwargs)
# Last-column (Notes) cells left-aligned for readability.
tbl[(ri, n_cols - 1)].get_text().set_ha("left")
# Force left-align on the Notes header too.
tbl[(0, n_cols - 1)].get_text().set_ha("left")
fig.suptitle(
"GQA decode KV-sharding — per-PE memory & communication\n"
"(LLaMA 70B GQA single KV-head group · S_kv = 1 M, FP16, "
"80 layers)",
fontsize=11.5, y=0.94,
)
out = _OUT_DIR / "gqa_long_ctx_6cases_kv_sharding_table.png"
fig.savefig(out, dpi=150, bbox_inches="tight")
plt.close(fig)
print(f"wrote {out}")
return out
def main() -> Path:
_OUT_DIR.mkdir(parents=True, exist_ok=True)
n = len(_CASES)
fig = plt.figure(figsize=(20.0, 7.0))
# Three rows per column: case chip · axis-spec annotation · rectangle.
gs = fig.add_gridspec(3, n,
height_ratios=[0.55, 0.32, 8.5],
hspace=0.05, wspace=0.20,
left=0.04, right=0.99,
top=0.93, bottom=0.06)
for i, cfg in enumerate(_CASES):
# Top: case chip header.
ax_chip = fig.add_subplot(gs[0, i])
ax_chip.set_xticks([])
ax_chip.set_yticks([])
for spine in ax_chip.spines.values():
spine.set_visible(False)
ax_chip.add_patch(mpatches.Rectangle(
(0, 0), 1, 1, transform=ax_chip.transAxes,
facecolor=cfg["accent"], edgecolor=cfg["accent"]))
ax_chip.text(0.5, 0.5, cfg["label"],
ha="center", va="center",
fontsize=10, fontweight="bold",
color="white")
# Middle: axis-spec annotation — names which axis the cube
# shards on and which axis the PE shards on (essential for
# parsing Case 5 where the cube colour fills run X instead
# of Y, breaking the visual symmetry of the rest).
ax_axes = fig.add_subplot(gs[1, i])
ax_axes.set_xticks([])
ax_axes.set_yticks([])
for spine in ax_axes.spines.values():
spine.set_visible(False)
ax_axes.add_patch(mpatches.Rectangle(
(0, 0), 1, 1, transform=ax_axes.transAxes,
facecolor="#F5F5F7", edgecolor="#CCCCCC",
linewidth=0.6))
ax_axes.text(0.5, 0.5, cfg["axes"],
ha="center", va="center",
fontsize=8.5, fontweight="bold",
color="#1F4E79")
# Bottom: the tensor rectangle.
ax = fig.add_subplot(gs[2, i])
_draw_panel(ax, cfg)
ax.set_xlabel("X : d_head = 128 →",
fontsize=9, fontweight="bold",
fontstyle="italic", color="#1F4E79")
ax.set_ylabel("Y : S_kv = 1 M ↓",
fontsize=9, fontweight="bold",
fontstyle="italic", color="#1F4E79")
fig.suptitle(
"GQA decode KV-tensor sharding — 6 cases · "
"LLaMA 70B GQA single KV-head group · "
"C = 8 cubes × P = 8 PEs · S_kv = 1 M, FP16, 80 layers",
fontsize=12, y=0.99,
)
out = _OUT_DIR / "gqa_long_ctx_6cases_kv_sharding_diagram.png"
fig.savefig(out, dpi=150, bbox_inches="tight")
plt.close(fig)
print(f"wrote {out}")
# Companion table PNG (slide-14 export).
_make_table_png()
return out
if __name__ == "__main__":
main()
scripts/paper/paper_plot_gqa_prefill_compute_bound.py
+111
View File
@@ -0,0 +1,111 @@
"""Comparative figure for the compute-bound prefill composite study.
Reads sweep_prefill_compute_bound.json (emitted by milestone-1h-gqa,
sweep ``prefill_cb``) and writes one two-panel PNG:
gqa_prefill_compute_bound.png
Left — end-to-end prefill latency (µs) vs context length.
Right — MAC utilization (achieved / 8 TFLOP·s⁻¹ per-PE peak) vs context.
Unlike memory-bound decode (where command form is latency-neutral), in
compute-bound prefill the composite command keeps the MAC array fed by
streaming DMA↔compute per HW tile, so it wins on both latency and
utilization — and the margin grows with context (deeper P·V reduction =
more tiles to pipeline).
Run (after the bench):
GQA_1H_RUN=1 GQA_1H_SWEEPS=prefill_cb python -m kernbench.cli.main run \\
--bench milestone-1h-gqa --topology topology.yaml
python scripts/paper/paper_plot_gqa_prefill_compute_bound.py
"""
from __future__ import annotations
import json
from pathlib import Path
import matplotlib
matplotlib.use("Agg")
import matplotlib.pyplot as plt # noqa: E402
_REPO_ROOT = Path(__file__).resolve().parents[2]
_FIG_DIR = (
_REPO_ROOT / "src" / "kernbench" / "benches"
/ "1H_milestone_output" / "gqa" / "long_ctx"
)
_SWEEP_JSON = _FIG_DIR / "sweep_prefill_compute_bound.json"
_PAPER_FIG_DIR = (
_REPO_ROOT / "docs" / "report" / "1H-codesign-paper" / "figures"
)
_VARIANT_STYLE = {
"primitive": ("primitive (tl.dot, hand-tiled)", "#c0504d", "o"),
"composite": ("composite GEMM", "#3b6ea5", "s"),
"composite_extended": ("composite + softmax_merge", "#4f8a4f", "^"),
}
_ORDER = ("primitive", "composite", "composite_extended")
def _ctx_label(c: int) -> str:
return f"{c // 1024}K" if c >= 1024 else str(c)
def _series(rows, variant, key):
pts = sorted(((r["ctx_len"], r[key]) for r in rows
if r["variant"] == variant), key=lambda t: t[0])
return [p[0] for p in pts], [p[1] for p in pts]
def main() -> None:
sweep = json.loads(_SWEEP_JSON.read_text())
rows = sweep["rows"]
ctxs = sweep["ctx_points"]
fig, (ax_lat, ax_util) = plt.subplots(1, 2, figsize=(13.0, 4.8))
for v in _ORDER:
label, color, marker = _VARIANT_STYLE[v]
xs, lat = _series(rows, v, "latency_ns")
ax_lat.plot(xs, [y / 1e3 for y in lat], marker=marker,
color=color, label=label, lw=2)
xs, util = _series(rows, v, "mac_util")
ax_util.plot(xs, [u * 100 for u in util], marker=marker,
color=color, label=label, lw=2)
for ax in (ax_lat, ax_util):
ax.set_xscale("log", base=2)
ax.set_xticks(ctxs)
ax.set_xticklabels([_ctx_label(c) for c in ctxs])
ax.set_xlabel(r"context length (= $T_q$ = $S_{kv}$)")
ax.grid(True, ls=":", alpha=0.5)
ax.legend(fontsize=9)
ax_lat.set_ylabel("end-to-end prefill latency (µs)")
ax_lat.set_title("Compute-bound prefill latency per command form")
ax_util.set_ylabel("MAC utilization (% of 8 TFLOP·s⁻¹ peak)")
ax_util.set_title(
"MAC utilization — composite keeps the array fed; primitive starves"
)
ax_util.axhline(100, color="#888", ls="--", lw=1, alpha=0.6)
fig.suptitle(
"Compute-bound prefill attention — use of composite commands\n"
"single-rank, GQA single-KV-head group ($h_q{=}8$, $d_{\\text{head}}"
"{=}128$); $M{=}8T_q$ tile-filling",
fontsize=11,
)
fig.tight_layout(rect=(0, 0, 1, 0.92))
out = _FIG_DIR / "gqa_prefill_compute_bound.png"
fig.savefig(out, dpi=150)
plt.close(fig)
print(f"wrote {out}")
if _PAPER_FIG_DIR.is_dir():
dst = _PAPER_FIG_DIR / out.name
dst.write_bytes(out.read_bytes())
print(f"copied {dst}")
if __name__ == "__main__":
main()
scripts/paper/paper_plot_gqa_prefill_long_ctx_4cases.py
+253
View File
@@ -0,0 +1,253 @@
"""Comparative figures for milestone-gqa-prefill-long-ctx-4cases.
Mirror of paper_plot_gqa_decode_long_ctx_4cases but for the prefill
variant. Reads sweep.json (emitted by ``kernbench run --bench
milestone-gqa-prefill-long-ctx-4cases``) and writes four PNGs into
``docs/report/1H-codesign-paper/figures/``:
gqa_prefill_long_ctx_4cases_latency.png end-to-end latency per case
gqa_prefill_long_ctx_4cases_traffic.png ipcq/dma op-count breakdown
gqa_prefill_long_ctx_4cases_memory.png per-PE KV bytes per case
gqa_prefill_long_ctx_4cases_parallelism.png active-PE × S_local work load
Run (after the bench):
GQA_PREFILL_LONG_CTX_4CASES_RUN=1 python -m kernbench.cli.main run \\
--bench milestone-gqa-prefill-long-ctx-4cases --topology topology.yaml
python scripts/paper/paper_plot_gqa_prefill_long_ctx_4cases.py
"""
from __future__ import annotations
import json
from pathlib import Path
import matplotlib
matplotlib.use("Agg")
import matplotlib.pyplot as plt # noqa: E402
_REPO_ROOT = Path(__file__).resolve().parents[2]
# Sweep JSON + PNGs live together under the bench output dir.
_FIG_DIR = (
_REPO_ROOT / "src" / "kernbench" / "benches"
/ "1H_milestone_output" / "gqa" / "long_ctx"
)
_SWEEP_JSON = _FIG_DIR / "sweep_prefill.json"
# Panel name → (short label, case ordinal for left-to-right plot order).
_CASE_INFO = {
"single_kv_group_prefill_long_ctx_gqa_cube_sp_pe_tp": (
"Case 1\nCube-SP × PE-TP", 1),
"single_kv_group_prefill_long_ctx_gqa_cube_repl_pe_tp": (
"Case 2\nCube-Repl × PE-TP", 2),
"single_kv_group_prefill_long_ctx_gqa_cube_repl_pe_sp": (
"Case 3\nCube-Repl × PE-SP", 3),
"single_kv_group_prefill_long_ctx_gqa_cube_sp_pe_sp": (
"Case 4 ★\nCube-SP × PE-SP", 4),
}
def _load() -> list[dict]:
return json.loads(_SWEEP_JSON.read_text())["rows"]
def _sorted_by_case(rows: list[dict]) -> list[dict]:
return sorted(rows, key=lambda r: _CASE_INFO[r["panel"]][1])
def _plot_latency(rows: list[dict]) -> Path:
rows = _sorted_by_case(rows)
labels = [_CASE_INFO[r["panel"]][0] for r in rows]
lat_us = [r["latency_ns"] / 1e3 for r in rows]
colors = ["#888", "#888", "#888", "#3b6ea5"] # Case 4 highlighted
fig, ax = plt.subplots(figsize=(8.0, 4.5))
bars = ax.bar(labels, lat_us, color=colors, width=0.6)
ax.set_ylabel("end-to-end latency (µs)")
ax.set_title(
"Long-context prefill 4-cases — end-to-end latency per case\n"
"LLaMA-3.1-70B single-KV-head group (8 cubes × 8 PEs)"
)
ax.bar_label(bars, fmt="%.1f", padding=3, fontsize=9)
ax.grid(axis="y", ls=":", alpha=0.5)
ax.set_ylim(0, max(lat_us) * 1.15)
fig.tight_layout()
out = _FIG_DIR / "gqa_prefill_long_ctx_4cases_latency.png"
fig.savefig(out, dpi=150)
plt.close(fig)
return out
def _plot_traffic(rows: list[dict]) -> Path:
rows = _sorted_by_case(rows)
labels = [_CASE_INFO[r["panel"]][0] for r in rows]
x = list(range(len(rows)))
keys = ["ipcq_copy_count", "dma_read_count", "dma_write_count"]
disp = ["IPCQ copy", "DMA read", "DMA write"]
colors = ["#c0504d", "#9bbb59", "#8064a2"]
w = 0.25
fig, ax = plt.subplots(figsize=(9.0, 4.5))
for i, (k, d, c) in enumerate(zip(keys, disp, colors)):
vals = [r["op_log_summary"][k] for r in rows]
ax.bar([xi + (i - 1) * w for xi in x], vals, width=w, label=d, color=c)
ax.set_xticks(list(x))
ax.set_xticklabels(labels, fontsize=9)
ax.set_ylabel("op count")
ax.set_title("Long-context prefill 4-cases — op-count breakdown per case")
ax.legend(fontsize=9)
ax.grid(axis="y", ls=":", alpha=0.5)
fig.tight_layout()
out = _FIG_DIR / "gqa_prefill_long_ctx_4cases_traffic.png"
fig.savefig(out, dpi=150)
plt.close(fig)
return out
def _s_local_per_pe(panel: str, *, S_kv: int, C: int, P: int) -> int:
"""S_local (token count) each PE attends over locally.
Encodes the cube/pe sharding axes from the panel name:
cube_sp_pe_tp (Case 1): S_kv / C (pe=replicate within cube)
cube_repl_pe_tp (Case 2): S_kv (full S_kv per active PE)
cube_repl_pe_sp (Case 3): S_kv / P (pe=row_wise within cube)
cube_sp_pe_sp (Case 4): S_kv / (C·P) (★ 64-way split)
"""
S_per_cube = S_kv if "cube_repl" in panel else S_kv // C
return S_per_cube // P if "pe_sp" in panel else S_per_cube
def _active_pe_count(panel: str, *, C: int, P: int) -> int:
"""Number of PEs doing non-idle attention work.
Prefill T_q≫1 means PE-TP is *useful* (not wasted as in decode):
cube_sp_pe_tp (Case 1): C·P (T_q sharded across all 64 ranks)
cube_repl_pe_tp (Case 2): P (CUBE 0 only; T_q sharded across its P PEs)
cube_repl_pe_sp (Case 3): C·P (all PEs busy but cubes redundant)
cube_sp_pe_sp (Case 4): C·P (all 64 PEs doing unique work)
"""
if "cube_repl" in panel and "pe_tp" in panel:
return P
return C * P
def _kv_bytes_per_pe(panel: str, *, S_kv: int, h_kv: int,
d_head: int, C: int, P: int) -> int:
"""KV bytes a single PE references (K + V, f16, 2 B/elem)."""
s_local = _s_local_per_pe(panel, S_kv=S_kv, C=C, P=P)
return 2 * s_local * h_kv * d_head * 2
def _plot_memory(rows: list[dict]) -> Path:
"""Per-PE KV bytes — Case 4 wins (64-way split)."""
rows = _sorted_by_case(rows)
labels = [_CASE_INFO[r["panel"]][0] for r in rows]
mib_per_pe = [
_kv_bytes_per_pe(
r["panel"], S_kv=r["S_kv"], h_kv=r["h_kv"],
d_head=r["d_head"], C=r["C"], P=r["P"],
) / (1024 * 1024)
for r in rows
]
colors = ["#888", "#c0504d", "#888", "#3b6ea5"]
fig, ax = plt.subplots(figsize=(8.0, 4.5))
bars = ax.bar(labels, mib_per_pe, color=colors, width=0.6)
ax.set_ylabel("KV bytes per PE (MiB, K + V, f16)")
ax.set_title(
"Long-context prefill 4-cases — KV memory per PE\n"
"(one KV-head group; per-layer, full S_kv state)"
)
ax.bar_label(bars, fmt="%.3f", padding=3, fontsize=9)
ax.grid(axis="y", ls=":", alpha=0.5)
ax.set_ylim(0, max(mib_per_pe) * 1.15)
fig.tight_layout()
out = _FIG_DIR / "gqa_prefill_long_ctx_4cases_memory.png"
fig.savefig(out, dpi=150)
plt.close(fig)
return out
def _t_q_per_pe(panel: str, *, T_q: int, C: int, P: int) -> int:
"""T_q row count each active PE computes attention for.
cube_sp_pe_tp (Case 1): T_q / (C·P) (T_q sharded across all 64 ranks)
cube_repl_pe_tp (Case 2): T_q / P (T_q sharded across CUBE 0's P PEs)
cube_repl_pe_sp (Case 3): T_q (Q replicated on every PE)
cube_sp_pe_sp (Case 4): T_q / C (Q sharded by cube, replicated within)
"""
if "cube_sp" in panel and "pe_tp" in panel:
return T_q // (C * P)
if "cube_repl" in panel and "pe_tp" in panel:
return T_q // P
if "cube_repl" in panel and "pe_sp" in panel:
return T_q
return T_q // C # cube_sp_pe_sp
def _s_kv_processed_per_pe(panel: str, *, S_kv: int, C: int, P: int) -> int:
"""S_kv tokens each PE actually processes attention over.
Differs from ``_s_local_per_pe`` (OWNED KV bytes): for cases with
a Ring (Case 1, 4) each PE sees C ring steps so processes more
tokens than it locally owns.
cube_sp_pe_tp (Case 1): S_kv (Ring + pe=replicate within cube)
cube_repl_pe_tp (Case 2): S_kv (full KV per PE)
cube_repl_pe_sp (Case 3): S_kv / P (pe=row_wise; no Ring)
cube_sp_pe_sp (Case 4): S_kv / P (Ring restores full S_kv/P per PE)
"""
if "cube_sp" in panel and "pe_tp" in panel:
return S_kv
if "cube_repl" in panel and "pe_tp" in panel:
return S_kv
return S_kv // P # both pe_sp cases
def _plot_parallelism(rows: list[dict]) -> Path:
"""Total compute work (PE × T_q × S_kv token-pairs) — exposes Case 3's
redundancy. Cases 1, 2, 4 all do the same total work (correct
attention over T_q × S_kv); Case 3 does C× more (cubes redundantly
repeat the same compute because K/V is cube-replicated).
"""
rows = _sorted_by_case(rows)
labels = [_CASE_INFO[r["panel"]][0] for r in rows]
total_work = [
_active_pe_count(r["panel"], C=r["C"], P=r["P"])
* _t_q_per_pe(r["panel"], T_q=r["T_q"], C=r["C"], P=r["P"])
* _s_kv_processed_per_pe(
r["panel"], S_kv=r["S_kv"], C=r["C"], P=r["P"],
)
for r in rows
]
colors = ["#888", "#888", "#c0504d", "#3b6ea5"] # Case 3 red, Case 4 highlighted
fig, ax = plt.subplots(figsize=(8.0, 4.5))
bars = ax.bar(labels, total_work, color=colors, width=0.6)
ax.set_ylabel(
"total compute (PE × T_q × S_kv token-pairs; lower ⇒ less wasted work)"
)
ax.set_title(
"Long-context prefill 4-cases — total compute work across active PEs\n"
"(Case 3 replicates K/V across 8 cubes ⇒ 8× redundant compute)"
)
ax.bar_label(bars, fmt="%d", padding=3, fontsize=9)
ax.grid(axis="y", ls=":", alpha=0.5)
ax.set_ylim(0, max(total_work) * 1.15)
fig.tight_layout()
out = _FIG_DIR / "gqa_prefill_long_ctx_4cases_parallelism.png"
fig.savefig(out, dpi=150)
plt.close(fig)
return out
def main() -> None:
rows = _load()
_FIG_DIR.mkdir(parents=True, exist_ok=True)
p1 = _plot_latency(rows)
p2 = _plot_traffic(rows)
p3 = _plot_memory(rows)
p4 = _plot_parallelism(rows)
print(f"wrote {p1}")
print(f"wrote {p2}")
print(f"wrote {p3}")
print(f"wrote {p4}")
if __name__ == "__main__":
main()
scripts/paper/paper_plot_ipcq_alternatives_comparison.py
+1075
View File
File diff suppressed because it is too large Load Diff
scripts/paper/paper_sip_architecture.py
+91
View File
@@ -0,0 +1,91 @@
"""Re-emit sip_view.svg in an academic palette (white background, black
strokes) and convert it to PDF for the 1H-codesign-paper Figure 1.
Source of truth: docs/diagrams/sip_view.svg (generated by
src/kernbench/topology/visualizer.py, dark-ish theme).
Treatment per user request: keep the layout intact, just force a white
canvas and turn every stroke black; promote faint label text to black so
all annotations stay legible on white.
"""
from __future__ import annotations
import re
import shutil
import subprocess
from pathlib import Path
REPO = Path(__file__).resolve().parents[2]
SRC_SVG = REPO / "docs" / "diagrams" / "sip_view.svg"
OUT_DIR = REPO / "docs" / "report" / "1H-codesign-paper" / "figures"
OUT_SVG = OUT_DIR / "sip_architecture.svg"
OUT_PDF = OUT_DIR / "sip_architecture.pdf"
# Drop the "SIP VIEW" title; the figure caption already names the level.
TITLE_REMOVE: list[tuple[str, str]] = [
(' <text x="324" y="18" text-anchor="middle" font-family="monospace" '
'font-size="14" font-weight="bold" fill="#1e293b">SIP VIEW</text>\n',
''),
]
COLOR_MAP: list[tuple[str, str]] = [
# canvas background slate-50 -> pure white
('fill="#f8fafc"', 'fill="#ffffff"'),
# all strokes -> black
('stroke="#3b82f6"', 'stroke="#000000"'), # UCIe mesh blue lines
('stroke="#475569"', 'stroke="#000000"'), # cube block borders
('stroke="#0ea5e9"', 'stroke="#000000"'), # I/O sky-blue lines
# link annotation text slate-500 -> black for readability
('fill="#64748b"', 'fill="#000000"'),
]
# Font-size bumps so labels survive LaTeX \linewidth scaling at the
# half-text-width subfigure. CUBE block labels overflow the 48px block
# rects, which is acceptable here.
FONT_MAP: dict[str, str] = {
"7": "10",
"14": "17",
}
# Tighten whitespace: cube grid occupies x=[84,564], y=[128,520]; IO
# chiplet sits around y~50. With the title removed, crop top to y=40 so
# the IO chiplet keeps a small headroom. Crop 70px on each side and 113
# px from bottom.
LAYOUT_FIXUP: list[tuple[str, str]] = [
('<svg xmlns="http://www.w3.org/2000/svg" width="648" height="648" '
'viewBox="0 0 648 648">',
'<svg xmlns="http://www.w3.org/2000/svg" width="508" height="495" '
'viewBox="70 40 508 495">'),
]
def _bump_font(m: re.Match) -> str:
return f'font-size="{FONT_MAP.get(m.group(1), m.group(1))}"'
def main() -> None:
if not SRC_SVG.exists():
raise SystemExit(f"source SVG missing: {SRC_SVG}")
rsvg = shutil.which("rsvg-convert")
if rsvg is None:
raise SystemExit("rsvg-convert not found (brew install librsvg)")
svg = SRC_SVG.read_text(encoding="utf-8")
for old, new in TITLE_REMOVE + COLOR_MAP + LAYOUT_FIXUP:
if old not in svg:
print(f"warn: pattern not present in source SVG: {old}")
svg = svg.replace(old, new)
svg = re.sub(r'font-size="(\d+)"', _bump_font, svg)
OUT_DIR.mkdir(parents=True, exist_ok=True)
OUT_SVG.write_text(svg, encoding="utf-8")
subprocess.run(
[rsvg, "-f", "pdf", "-o", str(OUT_PDF), str(OUT_SVG)],
check=True,
)
print(f"wrote {OUT_SVG.relative_to(REPO)}")
print(f"wrote {OUT_PDF.relative_to(REPO)}")
if __name__ == "__main__":
main()
scripts/verify_case4_dhead_tp.py
+148
View File
@@ -0,0 +1,148 @@
"""Phase 1a smoke test for the Case 4 d_head-TP decode kernel.
Runs `gqa_attention_decode_long_ctx_cube_sp_pe_tp_dhead_kernel` at small
S_kv (2K) to verify:
1. it imports cleanly,
2. it runs without error under the same harness as Case 6,
3. captures op_log_summary (gemm/ipcq/dma counts),
4. compares vs analytical predictions and vs Case 6 (same memory tier).
Usage:
python scripts/verify_case4_dhead_tp.py
"""
from __future__ import annotations
import os
import sys
from pathlib import Path
from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_sp_pe_sp import (
gqa_attention_decode_long_ctx_cube_sp_pe_sp_kernel as _case6_kernel,
)
from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_sp_pe_tp_dhead import (
gqa_attention_decode_long_ctx_cube_sp_pe_tp_dhead_kernel as _case4_kernel,
)
from kernbench.benches.gqa_helpers.long_ctx._gqa_attention_decode_long_ctx_cube_tp_dhead_pe_sp import (
gqa_attention_decode_long_ctx_cube_tp_dhead_pe_sp_kernel as _case5_kernel,
)
from kernbench.benches.gqa_helpers.shared._gqa_panel_helpers import (
_ccl_cfg, _summarize_op_log,
)
from kernbench.ccl.sfr_config import configure_sfr_intercube_multisip
from kernbench.policy.placement.dp import DPPolicy
from kernbench.runtime_api.bench_runner import run_bench
from kernbench.runtime_api.types import resolve_device
from kernbench.sim_engine.engine import GraphEngine
from kernbench.topology.builder import resolve_topology
# Small smoke-test params (S_kv=2K is enough to exercise tile-loop + AR).
_PARAMS = dict(C=8, P=8, T_q=1, S_kv=2_048,
d_head=128, h_q=8, h_kv=1)
def _bench_fn_case4(ctx):
p = _PARAMS
configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
dp_full = DPPolicy(cube="replicate", pe="column_wise",
num_cubes=p["C"], num_pes=p["P"])
dp_kv = DPPolicy(cube="row_wise", pe="column_wise",
num_cubes=p["C"], num_pes=p["P"])
q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="q_c4")
k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="k_c4")
v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="v_c4")
o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="o_c4")
ctx.launch("case4_dhead_tp", _case4_kernel,
q, k, v, o,
p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
p["d_head"], p["C"], p["P"],
_auto_dim_remap=False)
def _bench_fn_case5(ctx):
p = _PARAMS
configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
dp_q = DPPolicy(cube="column_wise", pe="replicate",
num_cubes=p["C"], num_pes=p["P"])
dp_kv = DPPolicy(cube="column_wise", pe="row_wise",
num_cubes=p["C"], num_pes=p["P"])
q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_q, name="q_c5")
k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="k_c5")
v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="v_c5")
o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_q, name="o_c5")
ctx.launch("case5_dhead_tp_inter", _case5_kernel,
q, k, v, o,
p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
p["d_head"], p["C"], p["P"],
_auto_dim_remap=False)
def _bench_fn_case6(ctx):
p = _PARAMS
configure_sfr_intercube_multisip(ctx.engine, ctx.spec, _ccl_cfg())
dp_full = DPPolicy(cube="replicate", pe="replicate",
num_cubes=p["C"], num_pes=p["P"])
dp_kv = DPPolicy(cube="row_wise", pe="row_wise",
num_cubes=p["C"], num_pes=p["P"])
q = ctx.zeros((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="q_c6")
k = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="k_c6")
v = ctx.zeros((p["S_kv"], p["h_kv"] * p["d_head"]),
dtype="f16", dp=dp_kv, name="v_c6")
o = ctx.empty((p["T_q"], p["h_q"] * p["d_head"]),
dtype="f16", dp=dp_full, name="o_c6")
ctx.launch("case6_sp_sp", _case6_kernel,
q, k, v, o,
p["T_q"], p["S_kv"], p["h_q"], p["h_kv"],
p["d_head"], p["C"], p["P"],
_auto_dim_remap=False)
def main() -> int:
topology = os.environ.get("GQA_1H_TOPOLOGY", "topology.yaml")
topo = resolve_topology(topology)
print(f"Smoke params: {_PARAMS}")
print()
for label, bench_fn in (
("Case 4 (Cube-SP × PE-TP d_head)", _bench_fn_case4),
("Case 5 (Cube-TP d_head × PE-SP)", _bench_fn_case5),
("Case 6 (Cube-SP × PE-SP S_kv)", _bench_fn_case6),
):
try:
res = run_bench(
topology=topo, bench_fn=bench_fn,
device=resolve_device(None),
engine_factory=lambda t, d: GraphEngine(
getattr(t, "topology_obj", t), enable_data=True,
),
)
except Exception as e:
print(f" {label:<42} FAIL: {type(e).__name__}: {e}")
continue
if not res.completion.ok:
print(f" {label:<42} ENGINE FAIL: {res.completion}")
continue
s = _summarize_op_log(res.engine.op_log)
lat = (res.engine.op_log[-1].t_end if res.engine.op_log else 0.0)
print(f" {label:<42} "
f"gemm={s['gemm_count']:>4} "
f"ipcq={s['ipcq_copy_count']:>4} "
f"dma_r={s['dma_read_count']:>4} "
f"dma_w={s['dma_write_count']:>3} "
f"latency={lat:.1f} ns "
f"(n_ops={len(res.engine.op_log)})")
return 0
if __name__ == "__main__":
sys.exit(main())
src/kernbench/benches/1H_milestone_output/IPCQ/ipcq_alternatives_architecture.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 252 KiB

src/kernbench/benches/1H_milestone_output/IPCQ/ipcq_alternatives_architecture_flow.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 305 KiB

src/kernbench/benches/1H_milestone_output/IPCQ/ipcq_alternatives_architecture_stacked.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 381 KiB

src/kernbench/benches/1H_milestone_output/IPCQ/ipcq_alternatives_control_data_overlay.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 153 KiB

src/kernbench/benches/1H_milestone_output/IPCQ/ipcq_alternatives_decision_matrix.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 189 KiB

src/kernbench/benches/1H_milestone_output/IPCQ/ipcq_alternatives_latency.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 160 KiB

src/kernbench/benches/1H_milestone_output/IPCQ/ipcq_alternatives_opcounts.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 86 KiB

src/kernbench/benches/1H_milestone_output/IPCQ/ipcq_alternatives_sequence.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 209 KiB

src/kernbench/benches/1H_milestone_output/ccl/AllReduce_LRAB_2DMesh_6SiP_2x3.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 38 KiB

After

Width:  |  Height:  |  Size: 38 KiB

src/kernbench/benches/1H_milestone_output/ccl/AllReduce_LRAB_2Dtorus_6SiP_2x3.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 36 KiB

After

Width:  |  Height:  |  Size: 36 KiB

src/kernbench/benches/1H_milestone_output/ccl/AllReduce_LRAB_2Dtorus_6SiP_2x3_with_TCM_SRAM_HBM.csv
+13 -13
View File
@@ -1,13 +1,13 @@
buffer_kind,sip_topology,n_sips,n_elem,bytes_per_pe,latency_ns
hbm,torus_2d,6,128,256,2120.040000000012
hbm,torus_2d,6,1024,2048,2717.2783333333473
hbm,torus_2d,6,8192,16384,7315.184999999989
hbm,torus_2d,6,32768,65536,23081.26500000037
sram,torus_2d,6,128,256,2060.040000000012
sram,torus_2d,6,1024,2048,2909.2783333333473
sram,torus_2d,6,8192,16384,9523.184999999869
sram,torus_2d,6,32768,65536,32201.265000000385
tcm,torus_2d,6,128,256,1964.040000000012
tcm,torus_2d,6,1024,2048,2477.2783333333473
tcm,torus_2d,6,8192,16384,6403.185000000109
tcm,torus_2d,6,32768,65536,19865.265000000378
buffer_kind,sip_topology,n_sips,n_elem,bytes_per_pe,latency_ns
hbm,torus_2d,6,128,256,3113.040000000012
hbm,torus_2d,6,1024,2048,3710.2783333333527
hbm,torus_2d,6,8192,16384,8308.184999999929
hbm,torus_2d,6,32768,65536,24074.26500000037
sram,torus_2d,6,128,256,3053.040000000012
sram,torus_2d,6,1024,2048,3902.2783333333573
sram,torus_2d,6,8192,16384,10516.18499999987
sram,torus_2d,6,32768,65536,33194.265000000385
tcm,torus_2d,6,128,256,2957.040000000012
tcm,torus_2d,6,1024,2048,3470.2783333333527
tcm,torus_2d,6,8192,16384,7396.18499999999
tcm,torus_2d,6,32768,65536,20858.265000000378
1 buffer_kind sip_topology n_sips n_elem bytes_per_pe latency_ns
2 hbm torus_2d 6 128 256 2120.040000000012 3113.040000000012
3 hbm torus_2d 6 1024 2048 2717.2783333333473 3710.2783333333527
4 hbm torus_2d 6 8192 16384 7315.184999999989 8308.184999999929
5 hbm torus_2d 6 32768 65536 23081.26500000037 24074.26500000037
6 sram torus_2d 6 128 256 2060.040000000012 3053.040000000012
7 sram torus_2d 6 1024 2048 2909.2783333333473 3902.2783333333573
8 sram torus_2d 6 8192 16384 9523.184999999869 10516.18499999987
9 sram torus_2d 6 32768 65536 32201.265000000385 33194.265000000385
10 tcm torus_2d 6 128 256 1964.040000000012 2957.040000000012
11 tcm torus_2d 6 1024 2048 2477.2783333333473 3470.2783333333527
12 tcm torus_2d 6 8192 16384 6403.185000000109 7396.18499999999
13 tcm torus_2d 6 32768 65536 19865.265000000378 20858.265000000378
src/kernbench/benches/1H_milestone_output/ccl/AllReduce_LRAB_2Dtorus_6SiP_2x3_with_TCM_SRAM_HBM.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 75 KiB

After

Width:  |  Height:  |  Size: 77 KiB

src/kernbench/benches/1H_milestone_output/ccl/AllReduce_LRAB_Ring1D_6SiP_1x6.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 37 KiB

After

Width:  |  Height:  |  Size: 37 KiB

src/kernbench/benches/1H_milestone_output/ccl/comparison_mesh_vs_ring_vs_2DTorus_vs_theoretical_vs_fsim.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 86 KiB

After

Width:  |  Height:  |  Size: 86 KiB

src/kernbench/benches/1H_milestone_output/ccl/summary.csv
+37 -37
View File
@@ -1,37 +1,37 @@
algorithm,sip_topology,n_sips,n_elem,bytes_per_pe,bytes_per_sip,latency_ns
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8,16,256,2666.552500000015
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32,64,1024,2747.7400000000152
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,64,128,2048,2855.990000000018
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,128,256,4096,3072.490000000019
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,512,1024,16384,3337.1133333333582
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,1024,2048,32768,3708.0333333333692
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,2048,4096,65536,4449.873333333393
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,4096,8192,131072,5933.020000000124
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8192,16384,262144,8900.379999999863
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,16384,32768,524288,14835.099999999224
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32768,65536,1048576,26704.540000000765
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,49152,98304,1572864,38573.97999999701
lrab_hierarchical_allreduce,ring_1d,6,8,16,256,2365.255833333347
lrab_hierarchical_allreduce,ring_1d,6,32,64,1024,2436.9433333333473
lrab_hierarchical_allreduce,ring_1d,6,64,128,2048,2532.526666666683
lrab_hierarchical_allreduce,ring_1d,6,128,256,4096,2723.693333333349
lrab_hierarchical_allreduce,ring_1d,6,512,1024,16384,3048.635000000021
lrab_hierarchical_allreduce,ring_1d,6,1024,2048,32768,3393.4016666666957
lrab_hierarchical_allreduce,ring_1d,6,2048,4096,65536,4082.401666666714
lrab_hierarchical_allreduce,ring_1d,6,4096,8192,131072,5458.80166666677
lrab_hierarchical_allreduce,ring_1d,6,8192,16384,262144,8216.934999999943
lrab_hierarchical_allreduce,ring_1d,6,16384,32768,524288,13733.201666665835
lrab_hierarchical_allreduce,ring_1d,6,32768,65536,1048576,24765.73500000064
lrab_hierarchical_allreduce,ring_1d,6,49152,98304,1572864,35798.268333331536
lrab_hierarchical_allreduce,torus_2d,6,8,16,256,1700.6025000000095
lrab_hierarchical_allreduce,torus_2d,6,32,64,1024,1753.2900000000102
lrab_hierarchical_allreduce,torus_2d,6,64,128,2048,1823.540000000012
lrab_hierarchical_allreduce,torus_2d,6,128,256,4096,1964.040000000012
lrab_hierarchical_allreduce,torus_2d,6,512,1024,16384,2196.8183333333463
lrab_hierarchical_allreduce,torus_2d,6,1024,2048,32768,2477.2783333333473
lrab_hierarchical_allreduce,torus_2d,6,2048,4096,65536,3038.1983333333583
lrab_hierarchical_allreduce,torus_2d,6,4096,8192,131072,4159.5050000000665
lrab_hierarchical_allreduce,torus_2d,6,8192,16384,262144,6403.185000000109
lrab_hierarchical_allreduce,torus_2d,6,16384,32768,524288,10890.5449999995
lrab_hierarchical_allreduce,torus_2d,6,32768,65536,1048576,19865.265000000378
lrab_hierarchical_allreduce,torus_2d,6,49152,98304,1572864,28839.98500000059
algorithm,sip_topology,n_sips,n_elem,bytes_per_pe,bytes_per_sip,latency_ns
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8,16,256,3782.5525000000202
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32,64,1024,3863.7400000000207
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,64,128,2048,3971.9900000000216
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,128,256,4096,4188.4900000000225
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,512,1024,16384,4453.113333333365
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,1024,2048,32768,4824.033333333375
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,2048,4096,65536,5565.873333333401
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,4096,8192,131072,7049.020000000062
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,8192,16384,262144,10016.379999999745
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,16384,32768,524288,15951.099999999256
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,32768,65536,1048576,27820.540000000765
lrab_hierarchical_allreduce,mesh_2d_no_wrap,6,49152,98304,1572864,39689.97999999693
lrab_hierarchical_allreduce,ring_1d,6,8,16,256,3524.2558333333504
lrab_hierarchical_allreduce,ring_1d,6,32,64,1024,3595.943333333351
lrab_hierarchical_allreduce,ring_1d,6,64,128,2048,3691.5266666666857
lrab_hierarchical_allreduce,ring_1d,6,128,256,4096,3882.6933333333527
lrab_hierarchical_allreduce,ring_1d,6,512,1024,16384,4207.101666666693
lrab_hierarchical_allreduce,ring_1d,6,1024,2048,32768,4550.801666666699
lrab_hierarchical_allreduce,ring_1d,6,2048,4096,65536,5240.335000000055
lrab_hierarchical_allreduce,ring_1d,6,4096,8192,131072,6617.801666666763
lrab_hierarchical_allreduce,ring_1d,6,8192,16384,262144,9375.934999999827
lrab_hierarchical_allreduce,ring_1d,6,16384,32768,524288,14892.201666666035
lrab_hierarchical_allreduce,ring_1d,6,32768,65536,1048576,25924.735000000648
lrab_hierarchical_allreduce,ring_1d,6,49152,98304,1572864,36957.268333330976
lrab_hierarchical_allreduce,torus_2d,6,8,16,256,2693.6025000000113
lrab_hierarchical_allreduce,torus_2d,6,32,64,1024,2746.290000000012
lrab_hierarchical_allreduce,torus_2d,6,64,128,2048,2816.5400000000127
lrab_hierarchical_allreduce,torus_2d,6,128,256,4096,2957.040000000012
lrab_hierarchical_allreduce,torus_2d,6,512,1024,16384,3189.81833333335
lrab_hierarchical_allreduce,torus_2d,6,1024,2048,32768,3470.2783333333527
lrab_hierarchical_allreduce,torus_2d,6,2048,4096,65536,4031.1983333333665
lrab_hierarchical_allreduce,torus_2d,6,4096,8192,131072,5152.5050000000665
lrab_hierarchical_allreduce,torus_2d,6,8192,16384,262144,7396.18499999999
lrab_hierarchical_allreduce,torus_2d,6,16384,32768,524288,11883.544999999496
lrab_hierarchical_allreduce,torus_2d,6,32768,65536,1048576,20858.265000000378
lrab_hierarchical_allreduce,torus_2d,6,49152,98304,1572864,29832.98500000003
1 algorithm sip_topology n_sips n_elem bytes_per_pe bytes_per_sip latency_ns
2 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 8 16 256 2666.552500000015 3782.5525000000202
3 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 32 64 1024 2747.7400000000152 3863.7400000000207
4 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 64 128 2048 2855.990000000018 3971.9900000000216
5 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 128 256 4096 3072.490000000019 4188.4900000000225
6 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 512 1024 16384 3337.1133333333582 4453.113333333365
7 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 1024 2048 32768 3708.0333333333692 4824.033333333375
8 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 2048 4096 65536 4449.873333333393 5565.873333333401
9 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 4096 8192 131072 5933.020000000124 7049.020000000062
10 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 8192 16384 262144 8900.379999999863 10016.379999999745
11 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 16384 32768 524288 14835.099999999224 15951.099999999256
12 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 32768 65536 1048576 26704.540000000765 27820.540000000765
13 lrab_hierarchical_allreduce mesh_2d_no_wrap 6 49152 98304 1572864 38573.97999999701 39689.97999999693
14 lrab_hierarchical_allreduce ring_1d 6 8 16 256 2365.255833333347 3524.2558333333504
15 lrab_hierarchical_allreduce ring_1d 6 32 64 1024 2436.9433333333473 3595.943333333351
16 lrab_hierarchical_allreduce ring_1d 6 64 128 2048 2532.526666666683 3691.5266666666857
17 lrab_hierarchical_allreduce ring_1d 6 128 256 4096 2723.693333333349 3882.6933333333527
18 lrab_hierarchical_allreduce ring_1d 6 512 1024 16384 3048.635000000021 4207.101666666693
19 lrab_hierarchical_allreduce ring_1d 6 1024 2048 32768 3393.4016666666957 4550.801666666699
20 lrab_hierarchical_allreduce ring_1d 6 2048 4096 65536 4082.401666666714 5240.335000000055
21 lrab_hierarchical_allreduce ring_1d 6 4096 8192 131072 5458.80166666677 6617.801666666763
22 lrab_hierarchical_allreduce ring_1d 6 8192 16384 262144 8216.934999999943 9375.934999999827
23 lrab_hierarchical_allreduce ring_1d 6 16384 32768 524288 13733.201666665835 14892.201666666035
24 lrab_hierarchical_allreduce ring_1d 6 32768 65536 1048576 24765.73500000064 25924.735000000648
25 lrab_hierarchical_allreduce ring_1d 6 49152 98304 1572864 35798.268333331536 36957.268333330976
26 lrab_hierarchical_allreduce torus_2d 6 8 16 256 1700.6025000000095 2693.6025000000113
27 lrab_hierarchical_allreduce torus_2d 6 32 64 1024 1753.2900000000102 2746.290000000012
28 lrab_hierarchical_allreduce torus_2d 6 64 128 2048 1823.540000000012 2816.5400000000127
29 lrab_hierarchical_allreduce torus_2d 6 128 256 4096 1964.040000000012 2957.040000000012
30 lrab_hierarchical_allreduce torus_2d 6 512 1024 16384 2196.8183333333463 3189.81833333335
31 lrab_hierarchical_allreduce torus_2d 6 1024 2048 32768 2477.2783333333473 3470.2783333333527
32 lrab_hierarchical_allreduce torus_2d 6 2048 4096 65536 3038.1983333333583 4031.1983333333665
33 lrab_hierarchical_allreduce torus_2d 6 4096 8192 131072 4159.5050000000665 5152.5050000000665
34 lrab_hierarchical_allreduce torus_2d 6 8192 16384 262144 6403.185000000109 7396.18499999999
35 lrab_hierarchical_allreduce torus_2d 6 16384 32768 524288 10890.5449999995 11883.544999999496
36 lrab_hierarchical_allreduce torus_2d 6 32768 65536 1048576 19865.265000000378 20858.265000000378
37 lrab_hierarchical_allreduce torus_2d 6 49152 98304 1572864 28839.98500000059 29832.98500000003
src/kernbench/benches/1H_milestone_output/ccl/topology.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 194 KiB

After

Width:  |  Height:  |  Size: 194 KiB

src/kernbench/benches/1H_milestone_output/gemm/gemm_composite_vs_async.json
+338
View File
@@ -0,0 +1,338 @@
[
{
"M": 32,
"K": 32,
"N": 32,
"bench": "matmul-composite",
"variant": "load_ref",
"pe_window_ns": 171.394,
"engine_window_ns": 106.38400000000001,
"flops": 65536,
"tflops": 0.6160324860881335,
"n_records": 7
},
{
"M": 32,
"K": 32,
"N": 32,
"bench": "matmul-async",
"variant": null,
"pe_window_ns": 103.22199999999998,
"engine_window_ns": 71.71199999999999,
"flops": 65536,
"tflops": 0.9138777331548418,
"n_records": 4
},
{
"M": 32,
"K": 32,
"N": 32,
"bench": "matmul-async-chunked",
"variant": null,
"pe_window_ns": 103.22199999999998,
"engine_window_ns": 71.71199999999999,
"flops": 65536,
"tflops": 0.9138777331548418,
"n_records": 4
},
{
"M": 32,
"K": 32,
"N": 32,
"bench": "matmul-async-chunked-db",
"variant": null,
"pe_window_ns": 103.22199999999998,
"engine_window_ns": 71.71199999999999,
"flops": 65536,
"tflops": 0.9138777331548418,
"n_records": 4
},
{
"M": 32,
"K": 64,
"N": 32,
"bench": "matmul-composite",
"variant": "load_ref",
"pe_window_ns": 179.394,
"engine_window_ns": 106.38400000000001,
"flops": 131072,
"tflops": 1.232064972176267,
"n_records": 7
},
{
"M": 32,
"K": 64,
"N": 32,
"bench": "matmul-async",
"variant": null,
"pe_window_ns": 127.41399999999999,
"engine_window_ns": 87.904,
"flops": 131072,
"tflops": 1.4910811794685113,
"n_records": 4
},
{
"M": 32,
"K": 64,
"N": 32,
"bench": "matmul-async-chunked",
"variant": null,
"pe_window_ns": 127.41399999999999,
"engine_window_ns": 87.904,
"flops": 131072,
"tflops": 1.4910811794685113,
"n_records": 4
},
{
"M": 32,
"K": 64,
"N": 32,
"bench": "matmul-async-chunked-db",
"variant": null,
"pe_window_ns": 127.41399999999999,
"engine_window_ns": 87.904,
"flops": 131072,
"tflops": 1.4910811794685113,
"n_records": 4
},
{
"M": 32,
"K": 128,
"N": 32,
"bench": "matmul-composite",
"variant": "load_ref",
"pe_window_ns": 211.77800000000002,
"engine_window_ns": 122.76800000000003,
"flops": 262144,
"tflops": 2.135279551674703,
"n_records": 10
},
{
"M": 32,
"K": 128,
"N": 32,
"bench": "matmul-async",
"variant": null,
"pe_window_ns": 175.798,
"engine_window_ns": 120.28800000000001,
"flops": 262144,
"tflops": 2.1793030061186482,
"n_records": 4
},
{
"M": 32,
"K": 128,
"N": 32,
"bench": "matmul-async-chunked",
"variant": null,
"pe_window_ns": 185.81799999999998,
"engine_window_ns": 130.308,
"flops": 262144,
"tflops": 2.011726064401265,
"n_records": 8
},
{
"M": 32,
"K": 128,
"N": 32,
"bench": "matmul-async-chunked-db",
"variant": null,
"pe_window_ns": 185.81799999999998,
"engine_window_ns": 130.308,
"flops": 262144,
"tflops": 2.011726064401265,
"n_records": 8
},
{
"M": 32,
"K": 128,
"N": 128,
"bench": "matmul-composite",
"variant": "load_ref",
"pe_window_ns": 315.394,
"engine_window_ns": 226.38400000000001,
"flops": 1048576,
"tflops": 4.631846773623577,
"n_records": 34
},
{
"M": 32,
"K": 128,
"N": 128,
"bench": "matmul-async",
"variant": null,
"pe_window_ns": 394.102,
"engine_window_ns": 338.592,
"flops": 1048576,
"tflops": 3.096871751252245,
"n_records": 4
},
{
"M": 32,
"K": 128,
"N": 128,
"bench": "matmul-async-chunked",
"variant": null,
"pe_window_ns": 368.12200000000007,
"engine_window_ns": 312.6120000000001,
"flops": 1048576,
"tflops": 3.354241040011259,
"n_records": 8
},
{
"M": 32,
"K": 128,
"N": 128,
"bench": "matmul-async-chunked-db",
"variant": null,
"pe_window_ns": 368.12200000000007,
"engine_window_ns": 312.6120000000001,
"flops": 1048576,
"tflops": 3.354241040011259,
"n_records": 8
},
{
"M": 32,
"K": 3072,
"N": 32,
"bench": "matmul-composite",
"variant": "load_ref",
"pe_window_ns": 1701.4420000000007,
"engine_window_ns": 876.4320000000005,
"flops": 6291456,
"tflops": 7.178487321320988,
"n_records": 148
},
{
"M": 32,
"K": 3072,
"N": 32,
"bench": "matmul-async",
"variant": null,
"pe_window_ns": 2401.4620000000004,
"engine_window_ns": 1609.952,
"flops": 6291456,
"tflops": 3.9078531533859397,
"n_records": 4
},
{
"M": 32,
"K": 3072,
"N": 32,
"bench": "matmul-async-chunked",
"variant": null,
"pe_window_ns": 3264.8920000000217,
"engine_window_ns": 2473.3820000000214,
"flops": 6291456,
"tflops": 2.5436653133240017,
"n_records": 192
},
{
"M": 32,
"K": 3072,
"N": 32,
"bench": "matmul-async-chunked-db",
"variant": null,
"pe_window_ns": 3286.402000000022,
"engine_window_ns": 2494.8920000000217,
"flops": 6291456,
"tflops": 2.5217348085608298,
"n_records": 192
},
{
"M": 8,
"K": 128,
"N": 128,
"bench": "matmul-composite",
"variant": "load_ref",
"pe_window_ns": 291.394,
"engine_window_ns": 226.38400000000001,
"flops": 262144,
"tflops": 1.1579616934058943,
"n_records": 34
},
{
"M": 8,
"K": 128,
"N": 128,
"bench": "matmul-async",
"variant": null,
"pe_window_ns": 247.798,
"engine_window_ns": 216.288,
"flops": 262144,
"tflops": 1.212013611480988,
"n_records": 4
},
{
"M": 8,
"K": 128,
"N": 128,
"bench": "matmul-async-chunked",
"variant": null,
"pe_window_ns": 251.42399999999998,
"engine_window_ns": 214.914,
"flops": 262144,
"tflops": 1.2197623235340647,
"n_records": 8
},
{
"M": 8,
"K": 128,
"N": 128,
"bench": "matmul-async-chunked-db",
"variant": null,
"pe_window_ns": 251.42399999999998,
"engine_window_ns": 214.914,
"flops": 262144,
"tflops": 1.2197623235340647,
"n_records": 8
},
{
"M": 128,
"K": 8,
"N": 128,
"bench": "matmul-composite",
"variant": "load_ref",
"pe_window_ns": 462.01,
"engine_window_ns": 397.0,
"flops": 262144,
"tflops": 0.6603123425692695,
"n_records": 82
},
{
"M": 128,
"K": 8,
"N": 128,
"bench": "matmul-async",
"variant": null,
"pe_window_ns": 247.798,
"engine_window_ns": 216.288,
"flops": 262144,
"tflops": 1.212013611480988,
"n_records": 4
},
{
"M": 128,
"K": 8,
"N": 128,
"bench": "matmul-async-chunked",
"variant": null,
"pe_window_ns": 247.798,
"engine_window_ns": 216.288,
"flops": 262144,
"tflops": 1.212013611480988,
"n_records": 4
},
{
"M": 128,
"K": 8,
"N": 128,
"bench": "matmul-async-chunked-db",
"variant": null,
"pe_window_ns": 247.798,
"engine_window_ns": 216.288,
"flops": 262144,
"tflops": 1.212013611480988,
"n_records": 4
}
]
src/kernbench/benches/1H_milestone_output/gemm/gemm_composite_vs_async_tflops.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 75 KiB

src/kernbench/benches/1H_milestone_output/gemm/gemm_hbm_bw_util.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 32 KiB

src/kernbench/benches/1H_milestone_output/gemm/gemm_mac_utilization_measured.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 40 KiB

After

Width:  |  Height:  |  Size: 29 KiB

src/kernbench/benches/1H_milestone_output/gemm/gemm_mac_utilization_theoretical_vs_measured.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 45 KiB

After

Width:  |  Height:  |  Size: 32 KiB

Some files were not shown because too many files have changed in this diff Show More