diff --git a/docs/ClusterFinderCUDA_benchmark_results.md b/docs/ClusterFinderCUDA_benchmark_results.md index ccfb8b71..58e884d5 100644 --- a/docs/ClusterFinderCUDA_benchmark_results.md +++ b/docs/ClusterFinderCUDA_benchmark_results.md @@ -21,10 +21,14 @@ The measurement harness is [`python/tests/perf/`](../python/tests/perf/) (see it | tag | directory | contents | |---|---|---| -| **`[f64]`** | [`perf/results/2026-08-18_f64/`](../python/tests/perf/results/2026-08-18_f64/) | `ladder_3x3.csv`, `ladder_9x9.csv`, `probes.csv`, 4 × `.nsys-rep`/`.sqlite` — `DEVICE_PED_TYPE=double`. **Acts I and II.** | -| **`[f32]`** | [`perf/results/2026-08-18_f32/`](../python/tests/perf/results/2026-08-18_f32/) | same, `DEVICE_PED_TYPE=float`. **Act III.** | +| **`[f64]`** | [`perf/results/2026-08-18_f64/`](../python/tests/perf/results/2026-08-18_f64/) | `ladder_3x3.csv`, `ladder_9x9.csv`, `probes.csv` — `DEVICE_PED_TYPE=double`, **cap 1500 at 9×9**. **3×3 for Acts I and II.** | +| **`[f32]`** | [`perf/results/2026-08-18_f32/`](../python/tests/perf/results/2026-08-18_f32/) | same, `DEVICE_PED_TYPE=float`. **3×3 for Act III.** | +| **`[f64]` 9×9** | [`perf/results/2026-08-20_f64_cap1700/`](../python/tests/perf/results/2026-08-20_f64_cap1700/) | `ladder_9x9.csv` at the **lossless cap 1700** — the re-take every 9×9 end-to-end number in this document comes from. | +| **`[f32]` 9×9** | [`perf/results/2026-08-20_f32_cap1700/`](../python/tests/perf/results/2026-08-20_f32_cap1700/) | same, `float`. | +| **cap A/B** | [`perf/results/2026-08-20_{f64,f32}_capAB/`](../python/tests/perf/results/) | `probes.csv` with **both caps probed in one session**, s1 and s4 — the source of every 9×9 engine number in §4, §11 and §14. | -Both arms are git rev `7177f00` on `bench/opt2-pipeline`. Reproduce either with +All five directories are git rev `7177f00` on `bench/opt2-pipeline`, same driver and GPU; +`env.json` records each. Reproduce either arm with `./run_campaign.sh f64` (or `f32`); the arm is selected by one line in `include/aare/clusterfinder_kernel.cuh` and nothing else differs. The `step` column in the CSVs carries the harness's internal labels — §15 maps them to the step names used @@ -36,13 +40,20 @@ the configuration: | | 3×3 | 9×9 | |---|--:|--:| | `N` | 100 000 | 20 000 | -| `max_clusters_per_frame` | 3 000 | 1 500 | +| `max_clusters_per_frame` | 3 000 | **1 700** | | `n_streams` | 4 | 4 | | `BATCH_SIZE` | 2 000 | 2 000 | | pedestal frames / `n_sigma` | 1 000 / 5 | 1 000 / 5 | | reps | 5 | 5 | | nsys probe frames | 20 000 | 20 000 | +The 9×9 cap was **1 500 in the 2026-08-18 campaign and 1 700 from 2026-08-20 on**: 1 500 +sat below the observed per-frame maximum of 1 633 and silently truncated 0.0095 % of +clusters (§3.5). Every 9×9 number in this document is at 1 700 unless the text says +otherwise; §12.1–§12.2 keep their cap-1500 figures deliberately, because the page-fault +diagnosis was made at that cap and the slot size is stated alongside. The cap is not a +free parameter at 9×9 — it sets the D2H bar directly (§4.2). + 9×9 is held at N = 20 000 because its result heap is ~5× larger per frame (1422 × 328 B = 466 kB vs 2330 × 40 B = 93 kB); 100 k would need 46.6 GB to retain against 98 GB free with no swap. Probes are 20 000 frames because shorter ones do not @@ -422,6 +433,45 @@ Three things this pins down, each of which was got wrong at some point: | 9×9 `[f64]` | 32.66 µs | 30.01 µs | **30.01 µs → 33 323 FPS** (kernel) | sustained | | 9×9 `[f32]` | 25.24 µs | 25.14 µs | **25.14 µs → 39 775 FPS** (**D2H**) | sustained | +#### The two columns are two different runs, made by two different tools + +This is the single easiest thing to get wrong, so it is worth stating flatly with 9×9 +`[f64]` as the worked case: + +| | **32.66 µs/frame** | **30.01 µs/frame** | +|---|---|---| +| file | `probes.csv`, `2026-08-20_f64_capAB` | `ladder_9x9.csv`, `2026-08-20_f64_cap1700` | +| harness | `run_probes.py` → `nsys_kernel_probe.py` | `run_ladder.py` → `ladder.py`, step `opt8` | +| **profiler** | **under nsys** | **none** | +| column | `kernel_us_per_frame` = `kernel_busy_ms / n_frames` | `wall_s / n_frames` = `0.60019 / 20 000` | +| what it measures | **one engine's** union occupancy | the **whole pipeline** end to end | +| scope | 653.140 ms of merged kernel intervals | 600.19 ms of wall clock | + +**30.01 does not come from nsys.** It cannot: the *same profiled run* that yields 32.66 +reports its own wall clock as `window_us_per_frame = 69.54 µs/frame` — 2.3× slower than +30.01 — because API tracing inflates it. That is exactly why the two harnesses are +separate, per `run_probes.py`: *"a profiled run cannot produce a throughput number, and an +unprofiled run cannot produce a per-engine breakdown. Two tools, two questions."* + +The divisor is **`n_frames` = 20 000 in both cases**. `batch = 2000` is the chunk size +*inside* each run, matched between the two harnesses so the submission pattern — and +therefore the overlap — is comparable; it is never a denominator. + +**Why they conflict, and which one loses.** 32.66 claims one engine needs 32.66 µs of +occupancy per frame; 30.01 says the whole pipeline emits a frame every 30.01 µs. A +pipeline cannot outrun its busiest engine, so one of them is wrong — and it is the +profiled one. Under nsys, submission is sparser, kernels overlap each other *less*, and a +less-overlapped interval set has a *larger* union. The measured 1.32× is therefore a lower +bound on the real overlap, and the true per-frame occupancy in the unprofiled run is below +32.66. Hence: the floor is the **lower** of the two, and a sustained rate outranks an +estimate. + +**"Peak" here and "floor" in the deck are the same quantity**, read in opposite units: a +floor in µs/frame, its reciprocal a peak in FPS. Both are the *lower* of probe estimate +and best sustained rate, so a slide saying "at the floor" and a table saying "100 % of +peak" are the same statement. What neither means is the raw engine max — 32.66 µs at 9×9 +`[f64]` is the profiled `max(H2D, kernel, D2H)`, and the floor is 30.01. + The 3×3 `[f64]` row is what keeps this from being circular: there the pipeline stopped 5.4 % **short** of the probe estimate, so the probe binds and opt6 reads 95 %, not 100 %. Where the sustained rate does win, the probe still corroborates it — to 1.9 %, 6.8 % and @@ -452,9 +502,13 @@ packets travel upstream against D2H's posted writes: | | 1 stream | 4 streams | penalty | |---|--:|--:|--:| | 3×3 H2D | 13.15 µs | 16.63 µs | **+26 %** | -| 9×9 H2D | 13.24 µs | 19.74 µs | **+49 %** | +| 9×9 H2D | 13.22 µs | 20.54 µs | **+55 %** | -The 1-stream figures are identical across cluster sizes (13.15 / 13.24 µs) and builds — +Both rows are `[f32]`; the 9×9 row is at **cap 1700** — at the old 1 500 the same +comparison read 13.24 → 19.74 µs (+49 %), because a smaller D2H slot leaves more of the +link for H2D. The penalty is a function of the cap, not only of the stream count. + +The 1-stream figures are identical across cluster sizes (13.15 / 13.22 µs) and builds — the payload is the same 320 000 B frame — which proves the H2D path itself is unchanged and the difference is contention alone. @@ -473,13 +527,22 @@ kernel). **Source**: `probes.csv`, columns `*_duty_pct`, `*_overlap`. | 9×9, 4 str | f64 | **47.0 %** | 29.9 % | 36.3 % | **1.32×** | | 9×9, 4 str | f32 | 34.2 % | 29.4 % | **36.1 %** | 1.02× | -Two readings: +Three readings: 1. **At 3×3 the copy engine is the saturated one** (~70 % H2D in both builds) — the GPU spends three quarters of its time being fed. At 9×9 no single engine exceeds 47 % because all three are comparable and interleave; the kernel is tallest but not dominant. -2. **`overlap` quantifies cross-stream kernel concurrency.** At f32/9×9 it is 1.02× — +2. **Only the kernel row ever overlaps.** `H2D_overlap` and `D2H_overlap` are **1.000 in + every row of `probes.csv`**, s1 and s4 alike: there is one copy engine per direction, + so same-direction transfers queue no matter how many streams issue them. For those two + engines `busy == sum`, and the `s4` per-frame figure is still a true mean duration — + which is why H2D *rises* 13.20 → 20.77 µs under load (each copy genuinely slows) while + the kernel *falls*. The concurrency four streams actually buy is H2D ∥ kernel ∥ D2H, + plus kernel-with-kernel; never copy-with-copy in the same direction. The stream + timelines in the deck (`fig_streams`, `fig_opt2_timeline`) draw same-direction copies + overlapping and carry a note saying so. +3. **`overlap` quantifies cross-stream kernel concurrency.** At f32/9×9 it is 1.02× — one 9×9 kernel nearly fills the GPU, so extra streams buy transfer overlap, not kernel co-execution. On f64 it rises to **1.32×**, because the longer f64 kernels leave more opportunity to interleave. This is why the f64 `s4` kernel column @@ -522,7 +585,7 @@ not the pipeline). **opt4 is the largest single step in Act I (1.32×), and §4 says why**: at 3×3 the H2D bar is the tallest, so pinning the input attacks the bar that actually binds. The same step -is worth only 1.02× at 9×9 (§6), where H2D is the *shortest* bar. This is the first +is worth only 1.03× at 9×9 (§6), where H2D is the *shortest* bar. This is the first confirmation of the rule. **Where Act I ends: 62 % of peak.** The remaining 9.8 µs/frame is host-side, and @@ -538,7 +601,12 @@ so every speedup divides by a cold CPU and reads ~9 % generous. State the conven **Source**: `ladder_9x9.csv` in `2026-08-20_f64_cap1700/`. N = 20 000, **cap 1700**, 4 streams, 5 reps. The earlier cap of 1500 sat below the per-frame maximum of 1633 and silently truncated 0.0095 % of clusters (§3.5); every 9×9 number here is the lossless -re-take. opt1/opt2 are absent: `ClusterFinderCUDAOpt2` is registered for 3×3 only, so +re-take. **The CPU row is the exception**: it is the swept optimum from +`2026-08-19_cpu_threads/` (32 threads, 665.22 µs), not the ladder's own CPU row in that +directory (688.28 µs at whatever thread count `ladder.py` defaulted to). The CPU finder +ignores `max_clusters_per_frame` entirely, so the cap difference does not affect it — +but the thread count does, by 24 % (§3.5), which is why the swept value is the one +quoted. opt1/opt2 are absent: `ClusterFinderCUDAOpt2` is registered for 3×3 only, so the 9×9 ladder starts at opt3. Peak: **30.01 µs / 33 323 FPS** — the best sustained rate; the probe's engine-occupancy estimate is 32.66 µs / 30 621 FPS, 8.1 % lower (§4). @@ -550,15 +618,15 @@ rate; the probe's engine-occupancy estimate is 32.66 µs / 30 621 FPS, 8.1 % low Two things are different here and both matter: -1. **opt4 is worth almost nothing (1.02×)** — exactly as §4 predicts. At 9×9, H2D is +1. **opt4 is worth almost nothing (1.03×)** — exactly as §4 predicts. At 9×9, H2D is 13.20 µs against a 39.86 µs kernel; making the shortest bar shorter changes nothing. 2. **The spread is 14 %.** At 3×3 the same steps vary 2–3 %. This is not noise, it is the result heap: at 9×9 each pass allocates ~9.3 GB, above glibc's mmap threshold, so it is - `munmap`ed and re-faulted every pass and never plateaus (§12.2). **A 24 % spread means + `munmap`ed and re-faulted every pass and never plateaus (§12.2). **A 14 % spread means the number depends on which run you quote** — and it is the first symptom of the problem Act II solves. -**Where Act I ends: 37 % of peak.** The GPU delivers a frame every 30 µs; the system +**Where Act I ends: 38 % of peak.** The GPU delivers a frame every 30 µs; the system delivers one every 80. **Three quarters of the time is host-side**, and none of the remaining device-side work can be reached until that is fixed. @@ -872,7 +940,7 @@ rewrite of §11.3 to be correct at all. | **9×9, 1 stream** | **39.86 µs** | **23.70 µs** | **−40.5 %** | | 3×3, 1 stream | 14.72 µs | 4.32 µs | −70.6 % | -> ### **opt7 result: the 9×9 kernel drops 40.7 %**, nsys-verified on both builds at 20 000 frames. +> ### **opt7 result: the 9×9 kernel drops 40.5 %**, nsys-verified on both builds at 20 000 frames. ### End-to-end, at the roofline @@ -977,8 +1045,8 @@ t = 0 and accumulate centered `Y = X − X0`. After the rewrite, full-f32 matche | H2D | 20.77 | 20.54 | | At 3×3 the act achieves its goal decisively — the kernel ends 11.1 µs *below* H2D. At 9×9 -it very nearly does: kernel and D2H end **6 % apart**. Further kernel work at 9×9 is -therefore capped at ~6 % before D2H binds instead, and the next lever there is a smaller +it very nearly does: kernel and D2H end **5 % apart** (23.94 vs 25.24 µs). Further kernel +work at 9×9 is therefore capped at ~5 % before D2H binds instead, and the next lever there is a smaller `cap` (§4.2) or coarser transfer granularity, not a faster kernel. > ### **The arc, in one line: the bottleneck has been walked from the host, to the GPU, to the wire.** @@ -1128,7 +1196,7 @@ Three conclusions: it. That is **292 k × 0.7 µs ≈ 204 ms ≈ 10 µs/frame of the 68 µs, permanently**, and no number of re-runs removes it. It is a floor `collect()` cannot get under. -Point 3 is an independent argument for opt6, and it explains the 24 % spread at the end of +Point 3 is an independent argument for opt6, and it explains the 14 % spread at the end of Act I (§6): `collect_view()` allocates nothing per frame, so it removes this floor outright rather than amortizing it. @@ -1225,11 +1293,12 @@ opt1/opt2 stop over-counting extended charge-shared events: | **ladder harness** | `perf/run_ladder.py` + `ladder.py` | Acts I–III end-to-end. One process per step; one CSV row per (step, rep) | | **probe sweep** | `perf/run_probes.py` + `nsys_kernel_probe.py` | §4. Four configs × 20 000 frames → `probes.csv` | | **duty cycles** | `perf/gpu_span.py` | interval-union analysis of an nsys SQLite export — the only way to separate "engine busy" from "engine idle" | -| **per-operation times** | `nsys stats` on a committed `.sqlite` | §4, §11. Kernel/memcpy durations without re-running anything — see below | +| **per-operation times** | `nsys stats` on a `.sqlite` export | §4, §11. Kernel/memcpy durations without re-running the ladder. **The `.nsys-rep`/`.sqlite` pair is deliberately untracked** (`results/.gitignore`) — 0.5–2 MB each, ~40 MB per campaign — so `probes.csv` is the durable record and re-deriving means re-running `run_probes.py` | | **shared plumbing** | `perf/common.py` | dataset, pedestal, fault bracketing, env capture, idle-GPU guard, CSV/manifest format | | **registers / occupancy** | `perf/kernel_resources.py` | §11.2. `cuobjdump -res-usage` on the built extension + the sm_89 occupancy arithmetic — no rebuild needed | -| f64 results (Acts I–II) | `perf/results/2026-08-18_f64/` | `ladder_3x3.csv`, `ladder_9x9.csv`, `probes.csv`, 4 × `.nsys-rep`/`.sqlite` | -| f32 results (Act III) | `perf/results/2026-08-18_f32/` | same, `DEVICE_PED_TYPE=float` | +| 3×3 results | `perf/results/2026-08-18_{f64,f32}/` | `ladder_3x3.csv`, `ladder_9x9.csv`, `probes.csv` — 9×9 rows here are **cap 1500**, superseded | +| 9×9 results | `perf/results/2026-08-20_{f64,f32}_cap1700/` | `ladder_9x9.csv` at the lossless cap — every 9×9 end-to-end number | +| 9×9 engine times | `perf/results/2026-08-20_{f64,f32}_capAB/` | `probes.csv`, both caps × {s1, s4} in one session | | opt7 build axis | `include/aare/clusterfinder_kernel.cuh` lines 16–17 | `COMPUTE_TYPE` / `DEVICE_PED_TYPE` | | opt1/opt2 class | `include/aare/ClusterFinderCUDAOpt2.hpp` | frozen pre-refactor pipeline; gained `time_kernels` for comparability | | opt3 – opt7 | `include/aare/ClusterFinderCUDA.hpp`, `clusterfinder_kernel.cuh` | current pipeline | @@ -1237,7 +1306,7 @@ opt1/opt2 stop over-counting extended charge-shared events: | exploratory notebook | `python/tests/ClusterFinderCUDA_perf.ipynb` | **stores only the last run** — not a record. Archive a copy per cluster size if used | | correctness notebook | `python/tests/ClusterFinderFrozen_vs_CUDA.ipynb` | CPU↔CUDA agreement analysis | | precision study | `docs/pedestal_precision_f32_cancellation.md` | B1 derivation | -| deck | `docs/cf_cuda_fused.pptx` + `docs/deck/build_fused_deck.py` | 29 slides in the same three acts; figures from `docs/deck/make_figs.py` | +| deck | `docs/cf_cuda_fused.pptx` + `docs/deck/build_fused_deck.py` | 34 slides in the same three acts plus a 15-slide annex (A1–A5); figures from `docs/deck/make_figs.py`. The `.pptx` is untracked — rebuild it from the script | ### CSV step labels @@ -1281,10 +1350,13 @@ Three things to know: here). It is the divisor and is not stored in the profile. 3. **`nsys stats` cannot produce a roofline.** It reports per-operation *sums and averages*, which cannot distinguish "the engine was busy" from "the engine was idle - waiting". It answers *how long is one kernel* — the `s1` column and the opt7 −40.7 % + waiting". It answers *how long is one kernel* — the `s1` column and the opt7 −40.5 % — while duty cycle, `overlap` and the roofline need the interval-union pass of `gpu_span.py` (§4.1). This is why the same f64 9×9 `s4` run reads 32.66 µs of kernel - occupancy while `nsys stats` reports ~39.9 µs per kernel instance. + occupancy while `nsys stats` reports **43.2 µs** per kernel instance (863.51 ms of + summed duration over 20 000 launches): under contention each kernel is *slower* than + its 39.86 µs `s1` duration, and only the 1.32× self-overlap pulls the per-frame + occupancy below either figure. Other useful reports: `cuda_gpu_trace` (every operation with timestamps), `cuda_gpu_mem_size_sum` (transfer bytes). `$NSYS stats --help-reports` lists them all; @@ -1328,10 +1400,10 @@ Every number here is from `[f32]`/`[f64]` (§0) and reproducible with 4. **3×3: 1.00× → 2.34× → 3.66× → 4.32× → 5.69×** over the best CPU configuration (24 threads), at identical correctness (4 × 10⁻⁶). The kernel never changed; what shrank was everything around it, 148 → 26 µs/frame. -5. **opt4 (pinned input) is worth 1.32× at 3×3 and 1.02× at 9×9** — the rule's first +5. **opt4 (pinned input) is worth 1.32× at 3×3 and 1.03× at 9×9** — the rule's first confirmation. Pinning attacks H2D, which is the tallest bar at 3×3 and the shortest at 9×9. -6. **Act I ends at 62 % of peak at 3×3 and 37 % at 9×9.** The GPU is idle much of the +6. **Act I ends at 62 % of peak at 3×3 and 38 % at 9×9.** The GPU is idle much of the time and the whole remainder is host-side. ### Act II — getting the results back `[f64]` diff --git a/docs/cf_cuda_fused.pptx b/docs/cf_cuda_fused.pptx new file mode 100644 index 00000000..800444fc Binary files /dev/null and b/docs/cf_cuda_fused.pptx differ diff --git a/docs/cf_cuda_kernel.pptx b/docs/cf_cuda_kernel.pptx new file mode 100644 index 00000000..230953b1 Binary files /dev/null and b/docs/cf_cuda_kernel.pptx differ diff --git a/docs/deck/CHANGELOG_2026-08-24.md b/docs/deck/CHANGELOG_2026-08-24.md new file mode 100644 index 00000000..41ee14d3 --- /dev/null +++ b/docs/deck/CHANGELOG_2026-08-24.md @@ -0,0 +1,154 @@ +# 2026-08-24 — deck and report corrections + +Session log for `docs/cf_cuda_fused.pptx` (via `docs/deck/build_fused_deck.py` + +`make_figs.py`) and `docs/ClusterFinderCUDA_benchmark_results.md`. +Deck rebuilt at **49 slides** (34 + 15-slide annex A1–A5). + +--- + +## 1. "Floor" meant two things — resolved + +- Root cause: `max(H2D, kernel, D2H)` and *best sustained rate* were both called the + floor. They coincide at 3×3 (16.17) and diverge at 9×9 f64 (**32.66** profiled max vs + **30.01** sustained), so the deck read as self-contradictory only at 9×9. +- Vocabulary unified on **"floor"** across ~9 slides: `% of peak` → `% of floor`, + `Peak · the GPU floor` → `The GPU floor`, and three `Peak is the…` captions rewritten. +- **A1** now carries two separate rows — `engine max [s4]` **and** + `FLOOR = lower of max, sustained` — so 32.66 and 30.01 both appear, each labelled. +- **12/34** (where the term is introduced) now says the max is a *profiled estimate* and + the floor is the lower of it and the sustained rate. +- **25/34** caption defined the floor as the engine max while `fig_overhead` plotted + 30.01; caption corrected and now names which arm sets it at each cluster size. + +## 2. Cap normalisation to 1700 + +- **19/34** and **34/34**: kernel claim `−40.7 %` → **`−40.5 %`** (39.86 → 23.70 µs, + cap 1700 · s1). +- **A5·3** is now the only cap-1500 slide, flagged with ⚠ and the reason no re-run is + needed (the RUNTIME column exists only in the trace; host-call cost is cap-independent). + +## 3. s1 vs s4 — the 19/34 ↔ 22/34 confusion + +- **Bug found**: `fig_f32_kernel`'s right panel was hardcoded to **cap-1500 s1** + (`39.93 / 19.44 / 13.24`), contradicting the callout above it, and it showed the kernel + binding in *both* arms — refuting the "D2H takes over" line on the same slide. +- Figure rebuilt as **two panels, s1 and s4**, same three engines, shared axis: + + | cap 1700 | kernel | D2H | H2D | binds | + |---|--:|--:|--:|---| + | s1 f64 | **39.86** | 21.97 | 13.20 | kernel | + | s1 f32 | **23.70** | 21.95 | 13.22 | kernel | + | s4 f64 | **32.66** | 25.25 | 20.77 | kernel | + | s4 f32 | 23.94 | **25.24** | 20.54 | **D2H** | + +- Binding engine named inline (`◀ binds`); dashed wall at 25.24; in-plot note that s4 is + *union occupancy*, not duration (f64 kernel: 43.2 µs mean duration, 1.32× self-overlap + → 32.66 µs/frame). +- **19/34** keeps the duration claim + payoff (`30.0 → 25.1 µs [s4]`); caption states the + rule — *quote s1 for how long an engine takes, s4 for which one sets the floor*. +- **22/34** gets the kernel → D2H conclusion in a second amber callout, as the end of the + Act III arc. +- Surfaced: the **s4** kernel gain is **−26.7 %**, not −40.5 % — self-overlap had already + hidden part of the win. + +## 4. Copy engines never overlap same-direction — timelines rescheduled + +- `H2D_overlap` and `D2H_overlap` are **1.000 in every row of `probes.csv`** (one copy + engine per direction). Only `kernel_overlap` exceeds 1. +- `fig_opt2_timeline` and `fig_streams` drew same-direction copies overlapping, which the + hardware cannot do. **Both redrawn from a real schedule** (`_schedule()` in + `make_figs.py`): H2D and D2H are single FIFO resources, kernels may overlap, a frame's + stages stay ordered, and a stream waits on its own previous D2H. +- Bar proportions changed from an unsourced `12 : 22 : 12` to **3×3 `[f64 · s1]`, + `13 : 15 : 5`** (13.14 / 14.72 / 5.31 µs) — these slides are 3×3-only. The H2D lane now + fills solid *because H2D is the tallest bar at 3×3*, rather than that being asserted. +- opt2's barrier drain is no longer a drawn-in gap: it falls out of the schedule as the + interval where the H2D engine starves waiting for the round to finish. +- Schedule reproduces the real step gain: opt2 18 units/frame → opt3 13, ratio **1.21×** + against a measured ×1.18. +- **A false claim fell out**: the old opt2 figure said "all four in flight". With one copy + engine the frame span (33) is only 2.5× the H2D stagger (13), so stream 0 has retired + before stream 3 gets the engine. The figure now *computes* the max concurrency off the + schedule and reports **3 frames in flight**. + +## 4b. 32.66 vs 30.01 — two runs, two tools + +- New §4 subsection *"The two columns are two different runs, made by two different + tools"*, because this was the easiest thing in the document to misread. +- **32.66 µs/frame** = `probes.csv` `kernel_us_per_frame` = `kernel_busy_ms / 20 000` + (653.140 ms of merged kernel intervals) — **one engine's union occupancy, under nsys**. +- **30.01 µs/frame** = `ladder_9x9.csv` opt8 `wall_s / 20 000` (0.60019 s) — the **whole + pipeline end to end, unprofiled**. It does **not** come from nsys. +- Proof it cannot: the same profiled run reports `window_us_per_frame = 69.54 µs/frame`, + 2.3× slower than 30.01. +- `batch = 2000` is the chunk size *inside* each run (matched across harnesses so the + overlap picture is comparable), never a denominator — both divide by `n_frames = 20 000`. +- Why the conflict resolves toward 30.01: under nsys submission is sparser, kernels + overlap *less*, and a less-overlapped interval set has a *larger* union. The 1.32× is a + lower bound on the real overlap. + +## 5. `fig_cancellation` right panel redrawn + +- Was: `axhline(42)` labelled "f32 error floor", damage shaded to rms 6.5 — and + 6.5² = 42.25, i.e. the line and the shading were each other's source. ~14× too large in + variance, and inconsistent with the same figure's left panel (which draws 3). +- Now sourced from `pedestal_precision_f32_cancellation.md` §5–§6: floor **±3–4 ADU²**, + axes rescaled to `(0, 6) × (0, 40)`, two zones (**rms < 2** clamped to 0; **rms 2–5** + corrupted but not clamped), and the doc's three table rows marked on the curve. +- Now consistent with the slide's own bullet and with the rail's "~1–2 % of the sensor". + +## 6. Layout fixes + +- **19/34**: bottom-right callout ran to 7.38 and covered the page number (footer 7.14) → + now 6.20 → 7.08. +- **13/34**: timeline picture overlapped its callout by 0.08" → figure 4.77 → 6.41, + callout 6.48 → 7.08. +- **13/34** and **14/34** re-laid out for the taller annotated timeline figures. + +--- + +## Report — `ClusterFinderCUDA_benchmark_results.md` + +Audited against the CSVs in `python/tests/perf/results/`. + +**Cap-1500 residues corrected** + +- §0 `max_clusters_per_frame` 9×9: `1 500` → **`1 700`**, plus a paragraph on which + sections deliberately keep 1500 (§12.1–§12.2, the page-fault diagnosis). +- §0 provenance table listed only the two `2026-08-18` dirs → now all five, including + `2026-08-20_{f64,f32}_cap1700` and `_capAB`, which §4 and §6 already cite. +- §4 9×9 H2D interference: `13.24 → 19.74 µs (+49 %)` → **`13.22 → 20.54 (+55 %)`**, old + reading kept as an explicit cap comparison. +- §6 / §11.2 / §16: opt4 at 9×9 `1.02×` → **`1.03×`** (3 places). +- §6 / §12.2: "a 24 % spread" → **`14 %`** (24 % was the cap-1500 spread). +- §6 / §16: "37 % of peak" → **`38 %`**. +- §11 / §15: opt7 kernel `−40.7 %` → **`−40.5 %`**. +- §11.4: kernel and D2H "6 % apart" → **`5 %`** (23.94 vs 25.24). + +**Other corrections** + +- §15 claimed the `.sqlite` files are *committed* — they are gitignored by design + (`results/.gitignore`) and absent from the tree. Now says so; `probes.csv` named as the + durable record. +- §15 said the deck is "29 slides" → **34 + 15-slide annex**; noted the `.pptx` is + untracked and rebuilt from the script. +- §15 read "the same f64 9×9 **s4** run … ~39.9 µs per kernel instance" — 39.9 is the + cap-1500 **s1** duration. For that s4 run the per-instance mean is **43.2 µs** + (863.51 ms / 20 000). Corrected, with the point spelled out. +- §6: the CPU baseline row (665.22 µs) comes from `2026-08-19_cpu_threads/`, not the + cap-1700 ladder (688.28 µs). Provenance footnote added. + +**Additions** + +- §4.1: new reading — only the kernel row ever overlaps; H2D/D2H stay true durations and + *rise* under contention while the kernel *falls*. +- §4: note that the report's "peak" (FPS) and the deck's "floor" (µs/frame) are the same + quantity in reciprocal units, and that neither is the raw engine max. + +--- + +## Still open + +- `fig_bottleneck` is generated by `make_figs.py` but referenced by no slide since 22/34 + moved to `fig_f32_absolute`. Dead weight, harmless. +- Slide 29/34's `[f64 ped]` validation rows are not backed by committed JSON. diff --git a/docs/deck/build_fused_deck.py b/docs/deck/build_fused_deck.py index 3505dd00..930d4790 100644 --- a/docs/deck/build_fused_deck.py +++ b/docs/deck/build_fused_deck.py @@ -166,23 +166,31 @@ def chrome(s, idx, eyebrow, title, title_size=27): run(para(tf, True, align=PP_ALIGN.RIGHT), f"{idx} / {n}", 8.5, MUTED) -N_ANNEX = 6 +N_ANNEX = 5 # annex GROUPS, not slides -def annex_chrome(s, idx, eyebrow, title, title_size=27): - """Same chrome, amber, on its own progress track — the annex is not part of - the 29-slide arc and should not look like it is.""" - rect(s, M, 0.60, 0.35, 0.035, AMBER) - tf = tb(s, 1.17, 0.50, 10.33, 0.32) - run(para(tf, True), f"ANNEX · {eyebrow}".upper(), 9, MUTED, bold=True, spc=1.6) +def annex_chrome(s, grp, eyebrow, title, part=None, nparts=None, title_size=27): + """Same chrome, amber, on its own progress track. + + The annex is GROUPED, not enumerated: slides that belong together carry the + same tag (A2 is both CUDA-Graphs slides, A5 all three artefacts), because + what a reader needs at a glance is which slides form one argument. The badge + is rendered in the header, not only in the corner. + """ + tf = tb(s, M, 0.44, 1.0, 0.34) + run(para(tf, True), f"A{grp}", 15, AMBER, bold=True) + tf = tb(s, 1.42, 0.50, 10.1, 0.32) + tag = f"ANNEX · {eyebrow}" + (f" · {part} of {nparts}" if nparts else "") + run(para(tf, True), tag.upper(), 9, MUTED, bold=True, spc=1.6) tf = tb(s, M, 0.86, 11.9, 1.0) run(para(tf, True, line=1.05), title, title_size, PALE, bold=True) pitch = 11.0 / N_ANNEX for i in range(N_ANNEX): rect(s, M + i * pitch, 7.28, pitch * 0.90, 0.045, - AMBER if i <= idx - 1 else RULE) - tf = tb(s, 12.0, 7.14, 0.9, 0.3) - run(para(tf, True, align=PP_ALIGN.RIGHT), f"A{idx} / {N_ANNEX}", 8.5, MUTED) + AMBER if i <= grp - 1 else RULE) + tf = tb(s, 11.7, 7.14, 1.2, 0.3) + foot = f"A{grp}" + (f" · {part}/{nparts}" if nparts else "") + run(para(tf, True, align=PP_ALIGN.RIGHT), foot, 8.5, MUTED) def table(s, x, y, w, header, rows, colw, size=9.5, rowh=0.62): @@ -480,26 +488,32 @@ run(para(tf, True), "AARE · HYBRID PIXEL DETECTORS · CUDA CLUSTERFINDER", tf = tb(s, M + 0.3, 1.28, 11.4, 1.45) run(para(tf, True, line=1.02), "The kernel was never the bottleneck — feeding it was", 42, PALE, bold=True) -tf = tb(s, M + 0.3, 2.80, 11.4, 0.8) +tf = tb(s, M + 0.3, 2.80, 11.4, 0.45) run(para(tf, True, line=1.05), "One kernel, one thread per pixel, and seven steps to keep it fed", 22, ACCENT) -tf = tb(s, M + 0.3, 3.40, 10.2, 0.9) +tf = tb(s, M + 0.3, 3.40, 11.4, 0.9) run(para(tf, True, line=1.3), - "The stencil was fast almost immediately. At 3×3 the kernel needs 5.5 µs per " - "frame; moving that frame across PCIe costs 13.2 µs at best, and 16.3 as the " - "pipeline runs it. Six of the seven optimization steps never touch the " - "arithmetic, they get data in, get results back, and measure honestly.", + "The stencil was fast almost immediately: at 3×3 the kernel needs 5.5 µs per " + "frame while getting that frame across PCIe costs 16.6 — 13.2 even " + "uncontended. The frame is gated by the wire, not the arithmetic. Six of the " + "seven steps get data in, get results back, and measure honestly.", 12.5, TEXT2) +tf = tb(s, M + 0.3, 4.40, 11.4, 0.26) +run(para(tf, True), + "WHERE THIS ENDS UP — THE SHIPPED f32 BUILD AFTER ALL SEVEN STEPS · " + "ENGINE TIMES [f32 · s4] · RECONCILED IN A1", + 8.5, AMBER, bold=True, spc=1.4) + stats = [("×9.1", "VS 24-THREAD CPU", ACCENT), ("61,312", "FRAMES / SECOND", PALE), ("16.3 µs", "PER FRAME, END TO END", PALE), ("6 / 23 M", "CLUSTER MISMATCH VS CPU TWIN", AMBER)] for i, (v, l, c) in enumerate(stats): x = M + 0.3 + i * 2.85 - rect(s, x, 4.35, 0.035, 0.95, c) - tf = tb(s, x + 0.22, 4.35, 2.5, 0.55) + rect(s, x, 4.78, 0.035, 0.95, c) + tf = tb(s, x + 0.22, 4.78, 2.5, 0.55) run(para(tf, True), v, 30, c, bold=True) - tf = tb(s, x + 0.22, 4.98, 2.5, 0.3) + tf = tb(s, x + 0.22, 5.41, 2.5, 0.3) run(para(tf, True), l, 8.5, MUTED, spc=1.2) rect(s, M + 0.3, 6.05, 11.0, 0.012, RULE) @@ -513,7 +527,7 @@ section("Context · what the code does", "What the kernel does, and what limits it", "No optimizations yet, only what the hardware has to do.", [("3–4", "The algorithm"), - ("5–6", "**The two machines**"), + ("5–6", "The two machines"), ("7–8", "The CUDA kernel"), ("9–10", "The hardware limit")], rng=(3, 10)) @@ -553,16 +567,19 @@ bullets(s, M, 1.95, COL, [ "**nσ · rms**, cut a 3×3 cluster around each local maximum.", "400×400 = 160 k pixels, **312.5 kB per frame**; Cu data yields ~2 330 clusters " "per frame at 3×3.", - "The pedestal is **updated by every non-photon pixel**, every frame, so the " - "arithmetic and the data movement are coupled.", + "The pedestal is **updated by every pixel that sees no photon** — about 80 % of " + "them — every frame, so the arithmetic and the data movement are coupled.", ]) code(s, M, 3.55, COL, [ "// the whole algorithm, per pixel", - "v = frame[i] - pedestal_mean[i]", - "rms = sqrt(pedestal_sum2[i]/n - pedestal_mean[i]^2)", - "if (v > «nSigma» * rms && v == max(3x3 window)) -> emit cluster", - "else -> update pedestal", -], title="THE KERNEL IN FIVE LINES") + "v = frame[i] - pedestal_mean[i]", + "rms = pedestal rms at i", + "m = max(v) over the 3x3 window at i", + "if (m > «nSigma» * rms) // a photon is within reach", + " if (v == m) -> emit cluster // ... and I am its peak", + " else -> «nothing» // ... I am in its shadow", + "else -> update pedestal // I saw nothing", +], title="THE WHOLE ALGORITHM · THREE OUTCOMES, NOT TWO") callout(s, M, 5.55, COL, "**Thesis of this talk:** the compute was fast almost immediately. " "Six of the seven steps are about feeding it.") @@ -709,22 +726,26 @@ caption(s, M, 5.80, 7.5, "handed to 128 SMs. Nothing about the launch depends on how many clusters " "the frame happens to contain, which is what makes the work uniform.") code(s, 8.5, 2.85, 4.1, [ - "if (!is_photon) {", - " // this pixel feeds the running", - " // pedestal for the NEXT frame", - " return;", + "// ClusterFinder.hpp — the serial CPU", + "if (max > nSigma * rms) {", + " if (value < max)", + " continue; // Not max go to the", + " // next pixel — but", + " // also no pedestal", + " // update", + "} else {", + " pedestal.«push_fast»(iy, ix, ...);", "}", - "uint32_t i = «atomicAdd»(d_count, 1u);", - "if (i < max_clusters)", - " d_clusters[i] = cluster;", -], size=8, title="THE CRITICAL BRANCH") -callout(s, 8.5, 4.62, 4.1, - "One counter per frame, bumped atomically. **The counter is the write " - "index**, a bump allocator, reset by a memset before each frame.", +], size=8, title="THREE OUTCOMES, NOT TWO") +callout(s, 8.5, 4.80, 4.1, + "That comment is **verbatim from the CPU source**. A pixel in a photon's " + "shadow is **neither recorded nor fed back** — and the CUDA kernel " + "reproduces it exactly.", h=1.15, size=10) -callout(s, 8.5, 5.95, 4.1, - "~99.9 % of threads take the **other** branch, which is why the pedestal " - "update, not the cluster write, dominates the kernel.", +callout(s, 8.5, 6.10, 4.1, + "**~80 % of threads update the pedestal**: every pixel with no photon in " + "its window. ~1.5 % are peaks, ~18 % are shadow. The update, not the " + "cluster write, dominates.", h=1.1, size=10, color=AMBER) # =========================================================== 6 · TILING @@ -839,15 +860,15 @@ s = new_slide() chrome(s, 11, "Roadmap", "Three acts, ordered by which bar is tallest") rows = [ - ("act", "ACT I · feed the GPU", "the host cannot submit work fast enough", "[f64]", ACCENT), + ("act", "ACT I · feed the GPU", "the host cannot submit work fast enough", "at 3×3 [f64]", ACCENT), ("opt1", "First CUDA port", "1 stream, one launch per frame", "×2.34", ACCENT), ("opt2", "Streams + batching", "4 streams, 2 000-frame batches", "×3.66", ACCENT), ("opt3", "Pipeline rework", "sync barriers removed", "×4.32", ACCENT), ("opt4", "Pinned memory", "DMA-speed host transfers", "×5.69", ACCENT), - ("act", "ACT II · get the results back", "the host copy is now the tallest bar", "[f64]", PALE), + ("act", "ACT II · get the results back", "the host copy is now the tallest bar", "at 3×3 [f64]", PALE), ("opt5", "Host↔GPU overlap", "chunked submit / collect", "×7.45", PALE), ("opt6", "Zero-copy collection", "read in place, never copy", "×8.65", PALE), - ("act", "ACT III · the kernel", "only now is the kernel the tallest bar", "[f32]", AMBER), + ("act", "ACT III · the kernel", "the kernel is the tallest bar — at 9×9", "at 9×9 [f32]", AMBER), ("opt7", "FP32 pedestal + variance rewrite", "the only kernel change in the deck", "kernel −41%", AMBER), ] @@ -881,7 +902,8 @@ caption(s, M, 6.62, 11.9, # --------------------------------------------------------- divider · ACT I section("Act I of III · feed the GPU", "The host cannot submit work fast enough", - "The kernel was fast almost immediately. This act is about the host.", + "The kernel was fast almost immediately. This act is about the host, and it " + "is told at 3×3, where the wire is the floor.", [(12, "opt1 · first port"), (13, "opt2 · streams + batching"), (14, "opt3 · no barriers"), @@ -902,22 +924,23 @@ bullets(s, M, 1.95, COL, [ "One cudaMemcpy in, one kernel, one cudaMemcpy out; **the host blocks " "on every frame**.", ]) -code(s, M, 3.62, COL, [ +code(s, M, 3.18, COL, [ "// one frame at a time, the host waits at every step", "cudaMemcpy(d_frame, h_frame, bytes, cudaMemcpyHostToDevice);", "find_clusters_in_single_frame", " <<>>(d_frame, d_pd_mean, ...);", "cudaMemcpy(h_out, d_out, out_bytes, cudaMemcpyDeviceToHost);", ], title="ClusterFinderCUDAOpt2.hpp · find_clusters()") -callout(s, M, 4.98, COL, +callout(s, M, 4.50, COL, "**PCIe is full-duplex**: H2D, kernel and D2H run on independent engines and " "overlap, so the floor is **max(H2D, kernel, D2H)**, never the sum. Each term " "is that engine's **busy time per frame at 4 streams**, the union of its " - "intervals, not how long one operation takes. At 3×3: max(**16.17**, 15.17, " - "7.69) = **16.2 µs → 61 859 FPS**.", h=1.30, size=10.5) -callout(s, M, 6.40, COL, - "opt1 delivers **15 807 FPS**, or **26 % of that peak**. Three quarters of every " - "frame is the host standing still.", h=0.70, size=10.5, color=AMBER) + "intervals, not how long one operation takes. At 3×3 **[f64 · s4]**: " + "max(**16.17**, 15.17, 7.69) = **16.2 µs → 61 859 FPS**. That max is a " + "**profiled estimate**, so the deck's floor is the **lower of it and the best " + "sustained rate** — the same 16.17 here; **A1** shows where the two part at 9×9.", + h=1.30, size=10.5) +figure(s, "fig_opt1_timeline", M, 5.85, COL) rail(s, [ ("label", "opt1 · 3×3 · 100 k frames · f64"), ("gap", 0.10), @@ -925,8 +948,8 @@ rail(s, [ ("stat", "vs 24-thread CPU", "×2.34", ACCENT), ("gap", 0.05), ("row", "Per frame", "63.3 µs", TEXT2), - ("row", "Peak · the GPU floor", "61 859 FPS", ACCENT), - ("row", "% of peak", "26 %", AMBER), + ("row", "The GPU floor", "61 859 FPS", ACCENT), + ("row", "% of floor", "26 %", AMBER), ]) # =========================================================== 12 · OPT2 @@ -940,7 +963,7 @@ bullets(s, M, 1.95, COL, [ "and pedestal. Frames are handed out **round-robin**.", "The host now submits **2 000 frames per call** instead of one.", ]) -code(s, M, 3.62, COL, [ +code(s, M, 3.30, COL, [ "struct StreamContext {", " cudaStream_t stream;", " FRAME_TYPE *d_frame; ClusterType *d_clusters;", @@ -948,18 +971,19 @@ code(s, M, 3.62, COL, [ "};", "auto &sc = v_sc[frame_idx % «n_streams»]; // round-robin", ], title="ClusterFinderCUDA.hpp · per-stream state") -callout(s, M, 5.72, COL, +figure(s, "fig_opt2_timeline", M, 4.77, COL) +callout(s, M, 6.48, COL, "**Scaffolding, not yet the payoff.** The streams exist, but the host still " - "synchronises after every round: see opt3.") + "synchronises after every round: see opt3.", h=0.60) rail(s, [ - ("label", "opt2 · 4 streams · batch 2 000"), + ("label", "opt2 · 3×3 · 4 streams · batch 2 000"), ("gap", 0.10), ("stat", "Throughput", "24 726 FPS", PALE), ("stat", "vs 24-thread CPU", "×3.66", ACCENT), ("gap", 0.05), ("row", "Per frame", "40.4 µs", TEXT2), ("row", "Step gain over opt1", "×1.56", ACCENT), - ("row", "% of peak · 61 859 FPS", "40 % (was 26)", AMBER), + ("row", "% of floor · 61 859 FPS", "40 % (was 26)", AMBER), ]) # =========================================================== 13 · OPT3 @@ -972,7 +996,7 @@ bullets(s, M, 1.95, 7.4, [ "opt3 submits every frame's H2D → kernel → D2H **asynchronously**, then " "synchronises **once at the end of the batch**.", ], size=10.5) -figure(s, "fig_streams", M, 3.05, 6.90) +figure(s, "fig_streams", M, 3.05, 6.55) code(s, 8.35, 1.95, 4.25, [ "// opt2: barrier after every round", "for (round) {", @@ -990,10 +1014,12 @@ code(s, 8.35, 1.95, 4.25, [ " «cudaStreamSynchronize»(sc.stream);", ], size=8, title="THE ONE-LINE IDEA") callout(s, 8.35, 5.05, 4.25, - "**29 188 FPS · ×4.32**\n34.3 µs/frame · 47 % of peak (was 40)", h=0.86, size=11) + "**29 188 FPS · ×4.32**\n34.3 µs/frame · 47 % of floor (was 40)", h=0.86, size=11) caption(s, 8.35, 6.15, 4.25, "Each lane is one stream. Removing the barrier lets a stream start its next " - "frame while its neighbours are still copying.") + "frame while its neighbours are still copying. The three panels are scheduled, " + "not sketched: H2D and D2H are one FIFO engine each, so a stream waits for the " + "copy engine, never for another stream's copy to finish overlapping it.") # =========================================================== 14 · OPT4 s = new_slide() @@ -1017,12 +1043,13 @@ code(s, 8.5, 3.15, 4.1, [ "«cudaHostUnregister»(ptr);", ], size=8, title="ClusterFinderCUDA.hpp") callout(s, 8.5, 4.72, 4.1, - "**38 486 FPS · ×5.69**\n26.0 µs/frame · 62 % of peak, the largest step in Act I", + "**38 486 FPS · ×5.69**\n26.0 µs/frame · 62 % of floor, the largest step in Act I", h=0.86, size=11) caption(s, 8.5, 5.82, 4.1, - "Measured H2D: one 400×400 uint16 frame (312.5 KiB = 320 000 B) in 13.2 µs " - "= 24.2 GB/s, 77% of PCIe 4.0 ×16 theoretical, i.e. true DMA speed. " - "Pageable staging runs ~15 GB/s.") + "Measured H2D [s1, uncontended]: one 400×400 uint16 frame (312.5 KiB = " + "320 000 B) in 13.2 µs = 24.2 GB/s, 77 % of PCIe 4.0 ×16 theoretical, i.e. " + "true DMA speed. In the shipped pipeline it reads 16.6 [s4], +26 % of " + "H2D↔D2H contention (A1). Pageable staging runs ~15 GB/s.") callout(s, M, 6.45, 7.6, "**The rule, first sighting: ×1.32 at 3×3 but only ×1.02 at 9×9.** Pinning " "attacks H2D, the tallest bar at 3×3, the shortest at 9×9.", @@ -1031,7 +1058,8 @@ callout(s, M, 6.45, 7.6, # -------------------------------------------------------- divider · ACT II section("Act II of III · get the results back", "The host copy is now the tallest bar", - "Frames go in at DMA speed. The results still come back slowly.", + "Frames go in at DMA speed. The results still come back slowly. Still 3×3 — " + "but 9×9 is where this act pays most.", [(16, "opt5 · host↔GPU overlap"), (17, "opt6 · zero-copy"), (18, "two rejected routes")], @@ -1054,7 +1082,7 @@ figure(s, "fig_overlap", M - 0.15, 3.28, COL + 0.30) callout(s, M, 5.86, COL, "The time hidden is **min(GPU, host)**, so this pays most when the two are " "comparable, and cannot rescue a host term that is simply larger.") -notes(s, """opt5 — host<->GPU overlap. Code and chunk sizing are on annex A6. +notes(s, """opt5 — host<->GPU overlap. Code and chunk sizing are on annex A3. tok = cf.submit_batch(data[a0:b0], first_frame=a0) for a, b in bounds[1:]: @@ -1077,13 +1105,13 @@ twice the GPU term, so overlap hides only the smaller one and opt5 is worth x1.20 - which is exactly the diagnosis that motivates opt6: you cannot overlap your way out of a host term that is simply larger. Report SS8.2.""") rail(s, [ - ("label", "opt5 · no CUDA work at all"), + ("label", "opt5 · 3×3 and 9×9 · no CUDA work at all"), ("gap", 0.10), ("stat", "3×3 throughput", "50 410 FPS", PALE), - ("row", "per frame · step · peak", "19.8 µs · ×1.31 · 81 %", ACCENT), + ("row", "per frame · step · of floor", "19.8 µs · ×1.31 · 81 %", ACCENT), ("gap", 0.14), ("stat", "9×9 throughput", "15 063 FPS", PALE), - ("row", "per frame · step · peak", "66.4 µs · ×1.20 · 45 %", AMBER), + ("row", "per frame · step · of floor", "66.4 µs · ×1.20 · 45 %", AMBER), ("gap", 0.15), ("note", "For clusters that must outlive the finder, this is the endpoint at " "3×3: opt6 lends, it does not give."), @@ -1110,17 +1138,18 @@ caption(s, M, 6.78, COL, "The two bars are the competing costs, not the two steps; the step times are " "on the right.", size=8.5) rail(s, [ - ("label", "opt6 · collect_view() · zero-copy"), + ("label", "opt6 · 3×3 and 9×9 · collect_view()"), ("gap", 0.10), ("stat", "3×3 throughput", "58 495 FPS", PALE), - ("row", "per frame · step · peak", "17.1 µs · ×1.16 · 95 %", ACCENT), + ("row", "per frame · step · of floor", "17.1 µs · ×1.16 · 95 %", ACCENT), ("gap", 0.14), ("stat", "9×9 throughput", "33 323 FPS", PALE), - ("row", "per frame · step · peak", "30.0 µs · ×2.21 · 100 %", AMBER), - ("gap", 0.15), - ("note", "End to end, opt5 → opt6: 19.8 → 17.1 µs and 66.4 → 30.0 µs. " - "Bit-identical results. Spread over 5 reps drops to 0.2 % at both " - "sizes, with zero warm page faults."), + ("row", "per frame · step · of floor", "30.0 µs · ×2.21 · 100 %", AMBER), + ("gap", 0.12), + ("note", "opt5 → opt6: 19.8 → 17.1 and 66.4 → 30.0 µs, bit-identical, " + "0.2 % spread, zero warm faults. FLOOR = lower of the s4 engine max " + "and the best sustained rate: at 3×3 the 16.17 µs max sets it; at 9×9 " + "f64 the run BEATS the 32.66 µs max, so the 30.0 sustained sets it [A1]."), ]) # =========================================================== 20 · ACT II rejected @@ -1163,7 +1192,8 @@ callout(s, M, 6.38, 11.9, # ------------------------------------------------------- divider · ACT III section("Act III of III · the kernel", "Only now is the kernel the tallest bar", - "The first change to the arithmetic, and the last step in the ladder.", + "The story moves to 9×9. At 3×3 the wire was always the floor and the kernel " + "never mattered; at 9×9 it is the tallest bar, which is why this act exists.", [(19, "opt7 · FP32 pedestal"), (20, "the correctness trap"), (21, "the variance rewrite"), @@ -1177,8 +1207,10 @@ s = new_slide() chrome(s, 19, "Act III · opt7 · FP32 device pedestal", "FP32 halves pedestal traffic: −41 % kernel time") bullets(s, M, 1.95, 7.5, [ - "~99.9% of pixels take the **pedestal-update** branch, which reads and writes " - "mean, sum and sum². In FP64 that is **32 bytes per pixel**; in FP32, 16.", + "**~80 % of pixels** take the **pedestal-update** branch: it reads off, sum, sum² " + "and writes back sum, sum², mean. All four pedestal arrays are DEVICE_PED_TYPE, " + "so one typedef halves every one of those six accesses — **48 bytes per updating " + "pixel in FP64**, **24 bytes in FP32**.", "The kernel is **bandwidth-bound**, so halving that traffic nearly halves the time.", "Second effect: on a GeForce part, **FP64 arithmetic runs at 1/64 of FP32**. " "The pedestal update was paying that tax on every pixel.", @@ -1193,14 +1225,20 @@ code(s, 8.5, 1.95, 4.1, [ "// was: double", ], size=8.5, title="ONE TYPEDEF") callout(s, 8.5, 3.05, 4.1, - "Kernel, 9×9, measured with Nsight Systems\n**39.9 µs → 23.7 µs (−40.7%)**", - h=0.92, size=11) -caption(s, 8.5, 4.20, 4.1, - "Exclusive per-kernel time, 20 000 frames, 1 stream, both builds at the same " - "git rev. Transfers are unchanged, as they must be.") -callout(s, 8.5, 5.30, 4.1, + "Kernel, 9×9 **[s1 · cap 1700 · exclusive]**\n**39.86 µs → 23.70 µs (−40.5%)**\n" + "Shipped **[s4]**: **30.0 → 25.1 µs** end to end.", + h=1.32, size=10.5) +caption(s, 8.5, 4.48, 4.1, + "Both panels, both builds, same git rev, 20 000 frames. The deck's rule, and " + "the reason there are two: quote s1 for how long an engine takes, s4 for which " + "one sets the floor. A1 has the full grid.") +callout(s, 8.5, 5.24, 4.1, "Naive FP32 is **wrong** (next two slides), and even when correct it only pays " - "**because Act II came first**.", h=1.05, size=10, color=AMBER) + "**because Act II came first**.", h=0.80, size=10, color=AMBER) +callout(s, 8.5, 6.20, 4.1, + "At 3×3 the same typedef is worth **−70.6 %** (14.72 → 4.32 µs [s1]) yet buys " + "only **4.6 %** end to end: there the kernel was never the tallest bar.", + h=0.88, size=10) # =========================================================== 19 · OPT6 trap s = new_slide() @@ -1208,9 +1246,11 @@ chrome(s, 20, "Act III · opt7 · the correctness trap", "The naive FP32 pedestal loses the variance to cancellation") bullets(s, M, 1.95, 7.9, [ "The running variance was computed as **var = E[X²] − mean²**. With a pedestal " - "mean of ~4 655 ADU, both terms are ≈ 2.17 × 10⁷ while the answer is ≈ 2 000.", - "In FP32 the spacing between representable numbers at 2.17 × 10⁷ is **2 048**, " - "larger than the variance itself. This is **catastrophic cancellation**.", + "mean of ~4 655 ADU, both operands are ≈ 2.17 × 10⁷ while the answer is ≈ 2 000.", + "In FP32 each operand sits on a grid of **2 ADU²**, so the answer inherits " + "**±3 — an absolute error that does not shrink**. For a pixel whose true " + "variance is 9 that is a third of it; below rms ≈ 2 it goes negative and the " + "**rms clamps to 0**.", ], size=10.5) figure(s, "fig_cancellation", M, 3.25, 7.5) rail(s, [ @@ -1233,12 +1273,11 @@ chrome(s, 21, "Act III · opt7 · the variance rewrite", bullets(s, M, 1.95, COL, [ "Freeze a per-pixel baseline **X₀ = round(mean)** once, at the end of pedestal " "training, and never move it again.", - "Accumulate the **centred** value Y = X − X₀ instead of X. Now both sums are " - "O(rms)-sized: the huge common term is gone before the subtraction.", - "The reported pedestal mean is still the full value, **X₀ + sum/n**, so nothing " - "downstream changes.", + "Accumulate the **centred** value Y = X − X₀ instead of X, and report the mean " + "as **X₀ + sum/n** — so nothing downstream changes.", ]) -code(s, M, 4.0, COL, [ +figure(s, "fig_variance_rewrite", M, 2.95, COL) +code(s, M, 4.47, COL, [ "// before: both terms ~2.17e7, answer ~2000", "var = sum2/n - mean*mean;", "", @@ -1246,7 +1285,7 @@ code(s, M, 4.0, COL, [ "DEVICE_PED_TYPE resid = mean - «d_pd_off»[i]; // ~O(1)", "DEVICE_PED_TYPE var_px = sum2[i]/n - resid*resid; // no cancellation", ], title="clusterfinder_kernel.cuh") -callout(s, M, 6.05, COL, +callout(s, M, 5.94, COL, "Result: the 100% FP32 build now matches the FP64 build to " "**3 × 10⁻⁷**, 70 clusters out of 233 million.") rail(s, [ @@ -1265,20 +1304,29 @@ rail(s, [ # =========================================================== 21 · OPT6 when s = new_slide() chrome(s, 22, "Act III · why this act comes last", - "Before Act II the kernel win is not even measurable") -figure(s, "fig_bottleneck", M, 1.95, 11.9) -callout(s, M, 5.58, 11.9, - "**Through collect_view() the identical −40% kernel is worth −16.2%, resolvable " - "to 0.0 points. Through every allocating path the interval spans 19–41 points " - "and straddles zero.** The result path does not just shrink the win; it " - "destroys the ability to see it.", h=0.86, size=11) -caption(s, M, 6.56, 11.9, - "Placed anywhere before Act II this step measures as noise and would reasonably " - "have been abandoned; placed after, it is the second-largest step in the ladder " - "at 9×9. The ordering is not presentational; it decides whether the " - "optimization is visible at all. 9×9, 20 000 frames, 4 streams, cap 1 700, " - "5 reps per arm, both arms at the same git rev. The band is what those reps " - "allow: allocating paths oscillate between allocator states and swamp the effect.") + "The saving never grew — the frame around it shrank") +figure(s, "fig_f32_absolute", M, 1.95, 11.9) +callout(s, M, 4.86, 5.85, + "**The same typedef saves −4.63 µs at opt4 and −4.87 µs at opt6** — the two " + "readings whose fault counts match. What changes is the denominator: the frame " + "falls 79.8 → 30.0 µs, so an identical saving reads **−5.8 %, then −16.2 %**. " + "Act II did not make the kernel win bigger; it made it **separable**.", + h=1.28, size=10.5) +callout(s, 6.75, 4.86, 5.85, + "**And the act ends by handing the floor away.** At **s4**, the config that " + "ships, the f64 arm is kernel-bound: **32.66** against a **25.25 µs** D2H. The " + "typedef puts the kernel at **23.94** — **below a D2H that never moved**. " + "The constraint is now the result path, not the arithmetic. [engines: 19/34]", + h=1.28, size=10.5, color=AMBER) +caption(s, M, 6.30, 11.9, + "9×9 · 20 000 frames · 4 streams · cap 1 700 · warm · both arms at the same git " + "rev. opt3 is excluded: its two arms sat in different allocator states, so that " + "comparison does not report the typedef at all. † opt5's arms differ the same " + "way; its reading agrees with the other two but is not independently " + "attributable. Both are worked through in annex A4. The f32 bar at opt6 is " + "opt7, the shipped build; the f32 bars at opt4 and opt5 are the same typedef " + "applied at earlier steps — configurations that exist only to make this " + "comparison controlled.") # ------------------------------------------------------- divider · results section("Results · what came out of it", @@ -1320,10 +1368,11 @@ callout(s, 6.75, 5.80, 5.85, "the rule, twice.", h=0.8, color=AMBER) caption(s, M, 6.62, 11.9, "9×9 · cap 1 700 (lossless; 1 500 truncated 0.0095 % of clusters) · 20 000 frames · " - "opt1/opt2 are 3×3-only · CPU baseline ClusterFinderMT at 32 threads. Peak is the " + "opt1/opt2 are 3×3-only · CPU baseline ClusterFinderMT at 32 threads. The floor is the " "lower of the nsys estimate and the best sustained rate; both arms are sustained-bound, " "corroborated to 8.8 % (f64) and 0.4 % (f32). Why the arms differ: 544.5 kB of D2H " - "costs 25.2 µs on both, hidden under the 32.7 µs f64 kernel, not under the 23.9 µs f32 one.") + "costs 25.2 µs on both [s4], hidden under the 32.7 µs f64 kernel, not under " + "the 23.9 µs f32 one.") # =========================================================== 24 · WHERE TIME GOES s = new_slide() @@ -1332,17 +1381,19 @@ chrome(s, 25, "Where the time actually went", figure(s, "fig_overhead", 1.95, 1.95, 9.4) callout(s, M, 5.30, 5.85, "**Acts I and II never touch the arithmetic.** The GPU floor is a flat " - "16.2 µs / 30.0 µs; what collapses is everything stacked on top of it.", + "16.2 µs at 3×3 / 30.0 µs at 9×9; what collapses is everything stacked on it.", h=0.86, size=10.5) callout(s, 6.75, 5.30, 5.85, "**Act III is the only step that lowers the floor itself**, and it could not " "have been seen until the stack above it was gone.", h=0.86, size=10.5, color=AMBER) caption(s, M, 6.45, 11.9, - "Blue/white/amber = the GPU floor for that act's build (max of H2D, kernel, " - "D2H; PCIe is full duplex, so it is never the sum). Grey = everything the " - "host adds on top. At 9×9 the host contributes +50 µs at opt3 and nothing at " - "opt6.") + "Blue/white/amber = the GPU floor for that act's build — the LOWER of the s4 " + "engine max (max of H2D, kernel, D2H; PCIe is full duplex, so never the sum) " + "and the best sustained rate. At 3×3 the engine max sets it, 16.17 µs; at 9×9 " + "f64 the sustained rate does, 30.01 against a 32.66 µs max. Grey = everything " + "the host adds on top. At 9×9 the host contributes +50 µs at opt3 and nothing " + "at opt6.") # =========================================================== 24 · MEASUREMENT AUDIT s = new_slide() @@ -1654,7 +1705,8 @@ for lab, dx in hdr: run(para(tf, True), lab.upper(), 9, MUTED, bold=True, spc=1.3) y += 0.44 params = [ - ("n_streams", "How many frames are in flight at once.", + ("n_streams", "How many frames may be in flight at once — an upper bound, not a " + "count: the copy engines cap the real number below it.", "4 at both sizes. 8 buys no kernel concurrency at 9×9 (+1% instance time) and " "inflates the CUDA-event timer 3.5×."), ("max_clusters_per_frame", "Fixed size of the per-frame D2H transfer.", @@ -1694,13 +1746,15 @@ cards = [ "twin that isolates the port the decisions are identical; the shipping f32 " "pedestal adds 6 duplicates in 23 million."), ("DONE", ACCENT, "FP32 pedestal, safely", - "−40.7% kernel, and correct, because the variance is now accumulated on a frozen " + "−40.5% kernel, and correct, because the variance is now accumulated on a frozen " "per-pixel offset instead of a raw second moment."), ("NEXT", AMBER, "3×3: transfer granularity", - "The 16.31 µs achieved sits 3.16 µs above the uncontended 13.15 µs H2D rate: " + "The 16.31 µs sustained sits 3.16 µs above the uncontended 13.15 µs H2D rate " + "[f32: s4 vs s1]: " "2 000 separate 320 kB descriptors, plus 26% of H2D↔D2H contention."), ("NEXT", AMBER, "9×9: the result path, not the kernel", - "At a lossless cap D2H already binds: 25.24 µs against a 23.94 µs kernel. More " + "At a lossless cap D2H already binds [f32 · s4]: 25.24 µs against a 23.94 µs " + "kernel. More " "kernel work buys nothing until the D2H slot stops being cap-sized."), ] for i, (tag, col, title, body) in enumerate(cards): @@ -1723,9 +1777,11 @@ callout(s, M, 6.58, 11.2, section("Annex · the evidence behind the claims", "Kept back so the arc stays readable", "The measurement artefacts, the rejected route, and the opt5 code.", - [("A1–A3", "The three benchmark artefacts"), - ("A4–A5", "CUDA Graphs, and why they lost"), - ("A6", "opt5 · the overlap code")], + [("A1", "Every engine number, reconciled"), + ("A2", "CUDA Graphs, and why they lost"), + ("A3", "opt5 · the overlap code"), + ("A4", "The fault model, tested"), + ("A5", "The three benchmark artefacts")], rng=(1, N_ANNEX), col=AMBER, annex=True, carry=("Everything so far", "34 slides", "the arc is finished; what follows answers questions")) @@ -1734,115 +1790,60 @@ section("Annex · the evidence behind the claims", # ANNEX — the measurement detail behind slides 26–27, and the rejected route A # =========================================================================== -# ------------------------------------------------------------ A1 · FAULTS +# ---- A1 · THE CONVENTION ------------------------------------------------- s = new_slide() -annex_chrome(s, 1, "expands slide 26 · artifact 1", - "First-touch page faults: two sources, one counter") +annex_chrome(s, 1, "measurement convention · expands slide 12", + "Uncontended, or as the pipeline runs it") bullets(s, M, 1.90, 12.0, [ - "A page exists in the process's address space but has no physical frame yet. " - "On first touch the kernel finds one, **zeroes it** (mandatory), and maps it. " - "No disk I/O: ru_majflt stays 0 all campaign. At 4 kB/page, **1 GB = 262 144 faults**.", + "Every engine time in this deck is tagged **[build · s1|s4]**. **s1** is one " + "stream with nothing else running: what an engine does **on its own**, which is " + "the right number for a capability claim and for the headroom that remains. " + "**s4** is the shipped four-stream pipeline: each engine's **busy time per " + "frame**, the union of its intervals — the only number that can set a floor.", ], size=10.5) -table(s, M, 2.62, 12.0, - ["", "(a) result heap", "(b) pinned D2H slots"], - [["allocator", "malloc → mmap, one ClusterVector per frame", - "cudaMallocHost, in submit_batch"], - ["cost / page", "**0.7 µs**", "**1.0 µs**: same fault + pin + DMA map"], - ["recurs?", "**yes**, every alloc/free cycle above the mmap threshold", - "**no**: once per buffer, for its lifetime"], - ["removed by", "**collect_view()**: it allocates nothing", - "nothing; reserve_output_slots() only moves it out of the timer"]], - colw=[0.14, 0.43, 0.43], size=9) -callout(s, M, 5.62, 5.85, - "**The two are exactly additive.** Reserving subtracts precisely the pre-pin " - "count from run 0 and changes nothing else.", h=0.72, size=10) -code(s, 6.75, 5.55, 5.85, [ - "3x3: 572 292 - 455 129 = 117 163 vs 117 192 pre-pin", - "9x9: 2 759 037 - 2 278 567 = 480 470 vs 480 474", - "closed form: 2 slots x 2000 x 120 004 B / 4 kB = 117 191", -], size=7.5, title="NOT A CORRELATION, AN IDENTITY") -callout(s, M, 6.42, 11.9, - "**At 9×9 the heap never plateaus.** ~9.3 GB per pass is above glibc's mmap " - "threshold, so it is munmap'd and re-faulted every pass: ~292 k faults ≈ " - "**10 µs/frame, permanently**. No number of re-runs removes it, which is an " - "independent argument for opt6.", h=0.72, size=10, color=AMBER) +table(s, M, 2.78, 5.8, + ["3×3 · µs/frame", "s1 f64", "s1 f32", "s4 f64", "s4 f32"], + [["H2D", "13.14", "13.15", "16.17", "16.63"], + ["kernel", "14.72", "4.32", "15.17", "5.53"], + ["D2H", "5.31", "5.27", "7.69", "7.57"], + ["engine max [s4]", "—", "—", "16.17", "16.63"], + ["FLOOR = lower of max, sustained", "—", "—", "**16.17**", "**16.31**"]], + colw=[0.28, 0.18, 0.18, 0.18, 0.18], size=9, rowh=0.44) +table(s, 7.0, 2.78, 5.8, + ["9×9 · cap 1700", "s1 f64", "s1 f32", "s4 f64", "s4 f32"], + [["H2D", "13.20", "13.22", "20.77", "20.54"], + ["kernel", "39.86", "23.70", "32.66", "23.94"], + ["D2H", "21.97", "21.95", "25.25", "25.24"], + ["engine max [s4]", "—", "—", "32.66", "25.24"], + ["FLOOR = lower of max, sustained", "—", "—", "**30.01**", "**25.14**"]], + colw=[0.28, 0.18, 0.18, 0.18, 0.18], size=9, rowh=0.44) +callout(s, M, 5.66, 12.0, + "**Two traps this table closes.** (1) s4 is engine *occupancy*, not duration — " + "the 9×9 kernel row **falls** 39.86 → 32.66 while every transfer rises, and " + "nothing that measures a duration falls under load. (2) **the FLOOR is not the " + "engine max**: the max is profiled and runs 2–8 % high, so the floor is whichever is " + "lower, it or the best unprofiled sustained rate — which is why opt6 reports " + "**30.01 µs against a 32.66 max**.", h=0.92, size=10.5) +caption(s, M, 6.66, 12.0, + "THE D2H SHIFT IS AN s4 PHENOMENON: at s1 the kernel binds in BOTH arms " + "(39.86 and 23.70 against a 21.97 / 21.95 D2H); only under four-stream " + "contention does D2H climb to 25.24 and overtake the 23.94 f32 kernel. Any " + "claim about which engine binds must be read from the s4 columns. " + "NEITHER DIRECTION IS FASTER THAN THE OTHER. At s1, H2D moves 320 000 B in " + "13.15 µs = 24.3 GB/s; D2H moves 120 004 B in 5.27 µs = 22.8 GB/s; at 9×9 D2H " + "moves 557 604 B in 21.95 µs = 25.4 GB/s. All sit at 72–81 % of PCIe 4.0 ×16's " + "31.5 GB/s, the smallest transfer paying the most fixed cost per byte — so a " + "taller bar always means MORE BYTES, never a slower wire. That is why cap 1 700 " + "hands the 9×9 floor to D2H. Source: probes.csv in perf/results/" + "2026-08-18_{f64,f32}/ (3×3) and 2026-08-20_{f64,f32}_capAB/ (9×9). The floor is the " + "LOWER of the s4 estimate and the best unprofiled sustained rate — 16.31 vs " + "16.63 at 3×3 f32, 25.14 vs 25.24 at 9×9 f32 — so a probe roofline is never a " + "hard denominator.") -# ------------------------------------------------------------ A2 · EVENTS +# ---- A2 · ROUTE A · IDEA ------------------------------------------------- s = new_slide() -annex_chrome(s, 2, "expands slide 26 · artifact 2", - "CUDA events measure the stream, not the kernel") -bullets(s, M, 1.90, COL, [ - "avg_kernel_time_ms() brackets the launch with **cudaEventRecord on the " - "kernel's own stream**. What it returns is elapsed time on **that stream's " - "timeline**, which includes time spent **queued behind other streams**.", - "So it is honest at 1 stream and inflates under saturation: up to **3.5×** at " - "8 streams. The tell is that the derived *PCIe + overhead* = wall/N − kernel_ms " - "**goes negative**: kernels overlap, so wall/frame < kernel/frame.", - "It is also not free: **~3.6 µs/frame**, i.e. 10–15 % of end-to-end throughput " - "at 3×3, to produce a number that is unusable exactly when it matters.", -], size=10.5) -code(s, M, 4.75, COL, [ - "if (m_time_kernels) // OFF by default, all 3 finders", - " cudaEventRecord(start[slot][i], sc.stream);", - "find_clusters_in_single_frame<<>>(...);", - "if (m_time_kernels)", - " cudaEventRecord(stop[slot][i], sc.stream); // <- queue-wait lands here", -], size=8, title="ClusterFinderCUDA.hpp") -callout(s, M, 6.10, COL, - "The flag exists for **comparability**, not preference: with events on for one " - "finder and off for another, the step between them absorbs the tax.", - h=0.72, size=10) -rail(s, [ - ("label", "9×9 kernel · f64 · nsys"), - ("gap", 0.10), - ("stat", "1 stream: duration", "39.86 µs", ACCENT), - ("stat", "4 streams: occupancy", "32.66 µs", AMBER), - ("gap", 0.05), - ("row", "overlap factor", "1.32×", TEXT2), - ("gap", 0.20), - ("note", "Same kernel. The s4 column is the union of kernel intervals per frame, " - "not how long one kernel takes. Quote s1 for duration; s4 feeds the peak, " - "subject to the sustained-rate rule on slide 12."), -]) - -# ------------------------------------------------------------ A3 · NSYS -s = new_slide() -annex_chrome(s, 3, "expands slide 26 · artifact 3", - "Where nsys is sound, and where it is not") -bullets(s, M, 1.90, 12.0, [ - "Tracing a **host** call means running a callback on entry and on exit, " - "**inside the interval being measured**. GPU work is different: the hardware " - "stamps its own start/end and the host reads those records **afterwards**, so " - "nothing is injected into the execution path.", -], size=10.5) -table(s, M, 2.60, 12.0, - ["measurement", "sqlite table", "9×9 f64, 1 stream", "verdict"], - [["cudaLaunchKernel: the host call", "CUPTI_..._RUNTIME", "1.85 µs", - "**inflated ~4×**"], - ["the kernel executing", "CUPTI_..._KERNEL", "39.93 µs", "sound to ~2 %"], - ["cudaMemcpyAsync: the host call", "CUPTI_..._RUNTIME", "1.65 µs", - "**inflated ~4×**"], - ["the H2D / D2H transfer", "CUPTI_..._MEMCPY", "13.25 / 19.44 µs", - "sound to ~2 %"]], - colw=[0.36, 0.24, 0.22, 0.18], size=9, rowh=0.52) -callout(s, M, 5.30, 11.9, - "**Same cudaMemcpyAsync, two numbers: 1.65 µs to ask for the copy, 13.25 µs " - "for the copy to happen.** If the profiler must be present at the moment to " - "measure it, it distorts it; if the hardware records it anyway and the " - "profiler reads it later, it does not.", h=0.86, size=10.5) -caption(s, M, 6.40, 11.9, - "Every headline number in this deck (kernel times, transfer times, duty " - "cycles, overlap and every engine floor) comes from the _KERNEL and _MEMCPY " - "tables (see gpu_span.py). The one figure read from _RUNTIME is the graph " - "launch budget in A5, and it is quoted to one significant figure " - "for exactly this reason. Proof the GPU side is sound: opt7 sustains 25.14 µs " - "unprofiled against a 25.24 µs estimate measured under the profiler, 0.4 % " - "apart, which a 4× distortion could not survive.") - -# ----------------------------------------------------------- A4 · ROUTE A -s = new_slide() -annex_chrome(s, 4, "route A · the idea and the verdict", - "CUDA Graphs, a sound idea that the next act overtook") +annex_chrome(s, 2, "rejected route · CUDA graphs · expands slide 11", + "CUDA Graphs, a sound idea that the next act overtook", part=1, nparts=2) bullets(s, M, 1.95, 12.0, [ "Every cudaMemcpyAsync / kernel launch costs the **CPU** a few microseconds of " "driver work, per frame and per operation. After opt4 that looked like the budget.", @@ -1871,10 +1872,10 @@ caption(s, M, 6.62, 12.0, "Launch overhead stops binding one step later; the technique aimed at it can no " "longer pay.") -# ------------------------------------------------------------ A5 · BUDGET +# ---- A3 · ROUTE A · BUDGET ----------------------------------------------- s = new_slide() -annex_chrome(s, 5, "route A · the arithmetic", - "What a CUDA Graph actually saves, in microseconds") +annex_chrome(s, 2, "rejected route · CUDA graphs", + "What a CUDA Graph actually saves, in microseconds", part=2, nparts=2) bullets(s, M, 1.90, COL, [ "The stream path issues **four runtime calls per frame**, a memset to clear the " "cluster counter, H2D, the launch, D2H. A graph replaces all four with **one** " @@ -1889,7 +1890,7 @@ table(s, M, 2.72, COL, colw=[0.40, 0.18, 0.22, 0.20], size=9, rowh=0.44) code(s, M, 5.42, COL, [ "7.67 us/frame x 3/4 = 5.75 us/frame eliminated, AS MEASURED (under nsys)", - "5.75 / 4 (see A3) ~ 1.4 us/frame eliminated, unprofiled estimate", + "5.75 / 4 (see A5) ~ 1.4 us/frame eliminated, unprofiled estimate", ], size=8, title="THE ARITHMETIC") callout(s, M, 6.28, COL, "**~1.4 µs against a 16.17 µs floor = 8.7 %**, real while the host is the " @@ -1910,9 +1911,9 @@ rail(s, [ ]) -# ---------------------------------------------------- A6 · THE OPT5 CODE +# ---- A4 · THE OPT5 CODE -------------------------------------------------- s = new_slide() -annex_chrome(s, 6, "expands slide 16 · opt5", +annex_chrome(s, 3, "opt5 · the overlap code · expands slide 16", "The overlap, in six lines, and why you never write them") code(s, M, 1.86, 7.15, [ "tok = cf.«submit_batch»(data[a0:b0], first_frame=a0)", @@ -1955,5 +1956,147 @@ caption(s, M, 6.80, 12.0, "where the host term is roughly twice the GPU term. That gap is the diagnosis " "that motivates opt6 (report §8.2).") +# ---- A5 · THE FAULT MODEL ------------------------------------------------ +s = new_slide() +annex_chrome(s, 4, "the fault model · expands slide 22", + "The fault model, tested against every step") +bullets(s, M, 1.90, 12.0, [ + "Slide 26 fits **0.68 µs per first-touch fault**. If that rate is right it should " + "predict the f64→f32 gap at every step from the fault counts alone — and where it " + "does, the comparison is measuring **the allocator**, not the typedef.", +], size=10.5) +table(s, M, 2.78, 12.0, + ["step", "f64 warm (faults)", "f32 warm (faults)", "Δ wall", + "Δ faults", "predicted", "verdict"], + [["opt3", "82.44 µs (128 k)", "95.66 µs (521 k)", "**+13.22**", "+393 k", + "**+13.37**", "**the allocator**"], + ["opt4", "79.83 µs (128 k)", "75.20 µs (127 k)", "−4.63", "−0.5 k", + "−0.02", "clean"], + ["opt5", "66.39 µs (152 k)", "61.85 µs (10 k)", "−4.54", "−141 k", + "**−4.81**", "**not separable**"], + ["opt6", "30.01 µs (0)", "25.14 µs (0)", "−4.87", "0", "0.00", "clean"]], + colw=[0.08, 0.18, 0.18, 0.10, 0.10, 0.13, 0.23], size=8.5, rowh=0.50) +callout(s, M, 5.42, 12.0, + "**opt3 is the whole argument in one row.** A 393 k fault gap predicts +13.37 µs; " + "+13.22 was observed — agreement to **1 %**. The −40 % kernel is in there " + "somewhere, invisible under 13 µs of the OS zeroing pages. Only opt4 and opt6, " + "where the fault term is ~0, report the typedef at all.", h=0.90, size=10.5) +caption(s, M, 6.52, 12.0, + "ladder_9x9.csv in results/2026-08-20_{f64,f32}_cap1700/, warm = best of reps " + "1–4, faults are that rep's own getrusage minor-fault count. Predicted = Δfaults " + "× 0.68 µs ÷ 20 000 frames. This table replaces an earlier figure that quoted " + "opt3's +16 % as a measurement of the result path; it is a measurement of two " + "allocator states, and the model above is how that was established.") + +# ---- A6 · FAULTS --------------------------------------------------------- +s = new_slide() +annex_chrome(s, 5, "benchmark artefacts · expands slide 26", + "First-touch page faults: two sources, one counter", part=1, nparts=3) +bullets(s, M, 1.90, 12.0, [ + "A page exists in the process's address space but has no physical frame yet. " + "On first touch the kernel finds one, **zeroes it** (mandatory), and maps it. " + "No disk I/O: ru_majflt stays 0 all campaign. At 4 kB/page, **1 GB = 262 144 faults**.", +], size=10.5) +table(s, M, 2.62, 12.0, + ["", "(a) result heap", "(b) pinned D2H slots"], + [["allocator", "malloc → mmap, one ClusterVector per frame", + "cudaMallocHost, in submit_batch"], + ["cost / page", "**0.7 µs**", "**1.0 µs**: same fault + pin + DMA map"], + ["recurs?", "**yes**, every alloc/free cycle above the mmap threshold", + "**no**: once per buffer, for its lifetime"], + ["removed by", "**collect_view()**: it allocates nothing", + "nothing; reserve_output_slots() only moves it out of the timer"]], + colw=[0.14, 0.43, 0.43], size=9) +callout(s, M, 5.62, 5.85, + "**The two are exactly additive.** Reserving subtracts precisely the pre-pin " + "count from run 0 and changes nothing else.", h=0.72, size=10) +code(s, 6.75, 5.55, 5.85, [ + "3x3: 572 292 - 455 129 = 117 163 vs 117 192 pre-pin", + "9x9: 2 759 037 - 2 278 567 = 480 470 vs 480 474", + "closed form: 2 slots x 2000 x 120 004 B / 4 kB = 117 191", +], size=7.5, title="NOT A CORRELATION, AN IDENTITY") +callout(s, M, 6.42, 11.9, + "**At 9×9 the heap never plateaus.** ~9.3 GB per pass is above glibc's mmap " + "threshold, so it is munmap'd and re-faulted every pass: ~292 k faults ≈ " + "**10 µs/frame, permanently**. No number of re-runs removes it, which is an " + "independent argument for opt6.", h=0.72, size=10, color=AMBER) + +# ---- A7 · EVENTS --------------------------------------------------------- +s = new_slide() +annex_chrome(s, 5, "benchmark artefacts", + "CUDA events measure the stream, not the kernel", part=2, nparts=3) +bullets(s, M, 1.90, COL, [ + "avg_kernel_time_ms() brackets the launch with **cudaEventRecord on the " + "kernel's own stream**. What it returns is elapsed time on **that stream's " + "timeline**, which includes time spent **queued behind other streams**.", + "So it is honest at 1 stream and inflates under saturation: up to **3.5×** at " + "8 streams. The tell is that the derived *PCIe + overhead* = wall/N − kernel_ms " + "**goes negative**: kernels overlap, so wall/frame < kernel/frame.", + "It is also not free: **~3.6 µs/frame**, i.e. 10–15 % of end-to-end throughput " + "at 3×3, to produce a number that is unusable exactly when it matters.", +], size=10.5) +code(s, M, 4.75, COL, [ + "if (m_time_kernels) // OFF by default, all 3 finders", + " cudaEventRecord(start[slot][i], sc.stream);", + "find_clusters_in_single_frame<<>>(...);", + "if (m_time_kernels)", + " cudaEventRecord(stop[slot][i], sc.stream); // <- queue-wait lands here", +], size=8, title="ClusterFinderCUDA.hpp") +callout(s, M, 6.10, COL, + "The flag exists for **comparability**, not preference: with events on for one " + "finder and off for another, the step between them absorbs the tax.", + h=0.72, size=10) +rail(s, [ + ("label", "9×9 kernel · f64 · nsys"), + ("gap", 0.10), + ("stat", "s1 · duration", "39.86 µs", ACCENT), + ("stat", "s4 · occupancy", "32.66 µs", AMBER), + ("gap", 0.05), + ("row", "overlap factor", "1.32×", TEXT2), + ("gap", 0.20), + ("note", "Same kernel. The s4 column is the union of kernel intervals per frame, " + "not how long one kernel takes. Quote s1 for duration; s4 feeds the floor, " + "subject to the sustained-rate rule on slide 12. Full grid: A1."), +]) + +# ---- A8 · NSYS ----------------------------------------------------------- +s = new_slide() +annex_chrome(s, 5, "benchmark artefacts", + "Where nsys is sound, and where it is not", part=3, nparts=3) +bullets(s, M, 1.90, 12.0, [ + "Tracing a **host** call means running a callback on entry and on exit, " + "**inside the interval being measured**. GPU work is different: the hardware " + "stamps its own start/end and the host reads those records **afterwards**, so " + "nothing is injected into the execution path.", +], size=10.5) +table(s, M, 2.60, 12.0, + ["measurement", "sqlite table", "9×9 [f64 · s1 · cap 1500]", "verdict"], + [["cudaLaunchKernel: the host call", "CUPTI_..._RUNTIME", "1.85 µs", + "**inflated ~4×**"], + ["the kernel executing", "CUPTI_..._KERNEL", "39.93 µs", "sound to ~2 %"], + ["cudaMemcpyAsync: the host call", "CUPTI_..._RUNTIME", "1.65 µs", + "**inflated ~4×**"], + ["the H2D / D2H transfer", "CUPTI_..._MEMCPY", "13.25 / 19.44 µs", + "sound to ~2 %"]], + colw=[0.36, 0.24, 0.22, 0.18], size=9, rowh=0.52) +callout(s, M, 5.30, 11.9, + "**Same cudaMemcpyAsync, two numbers: 1.65 µs to ask for the copy, 13.25 µs " + "for the copy to happen.** If the profiler must be present at the moment to " + "measure it, it distorts it; if the hardware records it anyway and the " + "profiler reads it later, it does not.", h=0.86, size=10.5) +caption(s, M, 6.40, 11.9, + "Every headline number in this deck (kernel times, transfer times, duty " + "cycles, overlap and every engine floor) comes from the _KERNEL and _MEMCPY " + "tables (see gpu_span.py). ⚠ CAP 1500 — the ONLY cap-1500 numbers left in the " + "deck, kept because the RUNTIME column exists only in the trace and the host " + "call cost does not depend on the cap. So the " + "19.44 µs D2H here is NOT the shipped 9×9 bar — at cap 1700 that engine reads " + "21.95 [s1] and 25.24 [s4], which is what binds the f32 build. The one figure " + "read from _RUNTIME is the graph " + "launch budget in A2, and it is quoted to one significant figure " + "for exactly this reason. Proof the GPU side is sound: opt7 sustains 25.14 µs " + "unprofiled against a 25.24 µs estimate measured under the profiler, 0.4 % " + "apart, which a 4× distortion could not survive.") + prs.save(OUT) print(f"saved {OUT} ({len(prs.slides._sldIdLst)} slides)") diff --git a/docs/deck/make_figs.py b/docs/deck/make_figs.py index b1bf9376..6703d0ee 100644 --- a/docs/deck/make_figs.py +++ b/docs/deck/make_figs.py @@ -36,14 +36,15 @@ Ladder, warm, us/frame 3x3 9x9 (cap 1700) opt6 zero-copy 17.10 30.01 opt7 = opt6 on f32 16.31 25.14 -Peak throughput = 1 / max(H2D, kernel, D2H), each term being that engine's BUSY +The GPU FLOOR (the deck says "floor" everywhere; older drafts said "peak") is +1 / max(H2D, kernel, D2H), each term being that engine's BUSY TIME PER FRAME (the union of its intervals) at the ladder's 4 streams -- and taken as the LOWER of two estimates: the profiled engine occupancy, and the best rate the unprofiled pipeline sustained. A sustained rate is an existence proof; the probe is an estimate made in a loop nsys slows to ~69 us/frame, where kernels overlap less and the union per frame reads high. - probe (nsys) best sustained PEAK binds + probe (nsys) best sustained FLOOR binds 3x3 f64 16.17 us 17.10 us 16.17 us -> 61 859 H2D 3x3 f32 16.63 us 16.31 us 16.31 us -> 61 312 H2D 9x9 f64 32.66 us 30.01 us 30.01 us -> 33 323 KERNEL @@ -412,49 +413,101 @@ def fig_overhead(): # ------------------------------------------------------ 4. streams timeline +# 3x3 [f64 · s1] proportions, 1 unit = 1 us: H2D 13.14, kernel 14.72, D2H 5.31. +# These slides are 3x3-only, and at 3x3 H2D is the tallest bar -- which the +# schedule below then reproduces on its own rather than being asserted. +H_, K_, D_ = 13, 15, 5 +LANE_ = 0.68 + + +def _frame_bars(ax, y, t0): + ax.broken_barh([(t0, H_)], (y, LANE_), facecolors=AMBER, zorder=3) + ax.broken_barh([(t0 + H_, K_)], (y, LANE_), facecolors=ACCENT, zorder=3) + ax.broken_barh([(t0 + H_ + K_, D_)], (y, LANE_), facecolors=PALE, zorder=3) + + +def _schedule(n_frames, n_streams, H=None, K=None, D=None, round_size=None): + """Greedy list schedule that honours the real engine constraints. + + H2D and D2H are ONE FIFO resource each -- the GPU has a single copy engine + per direction, which is why H2D_overlap and D2H_overlap are 1.00 in every + row of probes.csv. Kernels may overlap (measured 1.02-1.32x). A frame's + three stages stay ordered, and a stream cannot start its next frame until + its previous D2H has freed the per-stream buffers. + + round_size=n reproduces opt2's cudaDeviceSynchronize() after every round of + n frames: every resource resets to the round's end, so the drain is a + consequence of the barrier rather than a drawn-in gap. + + Returns (frames, drains) with frames = [(stream, h0, k0, d0)] and drains = + [(t_h2d_goes_idle, t_round_end)] -- exactly the time the barrier costs. + """ + H, K, D = H or H_, K or K_, D or D_ + h_free = d_free = 0.0 + ready = [0.0] * n_streams + frames, drains = [], [] + last_h_end = 0.0 + for i in range(n_frames): + if round_size and i and i % round_size == 0: + t = max(max(ready), h_free, d_free) + drains.append((last_h_end, t)) + h_free = d_free = t + ready = [t] * n_streams + s = i % n_streams + h0 = max(ready[s], h_free); h_free = last_h_end = h0 + H + k0 = h_free; k1 = k0 + K + d0 = max(k1, d_free); d_free = ready[s] = d0 + D + frames.append((s, h0, k0, d0)) + if round_size: + drains.append((last_h_end, max(max(ready), h_free, d_free))) + return frames, drains + + +def _draw_schedule(ax, frames, top_lane=3.0, H=None, K=None, D=None): + H, K, D = H or H_, K or K_, D or D_ + for s, h0, k0, d0 in frames: + y = top_lane - s * 1.0 + ax.broken_barh([(h0, H)], (y, LANE_), facecolors=AMBER, zorder=3) + ax.broken_barh([(k0, K)], (y, LANE_), facecolors=ACCENT, zorder=3) + ax.broken_barh([(d0, D)], (y, LANE_), facecolors=PALE, zorder=3) + + def fig_streams(): fig, axes = plt.subplots(3, 1, figsize=(7.7, 3.9)) - H, K, D = 12, 22, 12 - FR = H + K + D - LANE = 0.68 - - def frame(ax, lane_y, t0): - ax.broken_barh([(t0, H)], (lane_y, LANE), facecolors=AMBER, zorder=3) - ax.broken_barh([(t0 + H, K)], (lane_y, LANE), facecolors=ACCENT, zorder=3) - ax.broken_barh([(t0 + H + K, D)], (lane_y, LANE), facecolors=PALE, zorder=3) + FR = H_ + K_ + D_ # --- opt1: one stream, strictly serial ax = axes[0] for i in range(3): - frame(ax, 1.0, i * FR) + _frame_bars(ax, 1.0, i * FR) ax.set_ylim(0.4, 2.3) - ax.text(3 * FR + 6, 1.34, "GPU idle between every stage", color=MUTED, fontsize=7.5, + ax.text(3 * FR + 6, 1.34, "one engine at a time", color=MUTED, fontsize=7.5, va="center") - # --- opt2: 4 streams, barrier after each round + # --- opt2: 4 streams, barrier after each round. The drain is not drawn in; + # it falls out of the schedule, because after the round's last H2D the copy + # engine has nothing left to feed until the barrier releases. ax = axes[1] - ROUND = FR + 3 * 8 - for r in range(2): - for st in range(4): - frame(ax, 3 - st * 1.0, r * (ROUND + 26) + st * 8) - ax.axvspan(ROUND, ROUND + 26, color=AMBER, alpha=0.13, zorder=1) - ax.text(ROUND + 13, 4.15, "barrier — GPU drains", color=AMBER, fontsize=7.5, - ha="center", va="bottom") + frames, drains = _schedule(8, 4, round_size=4) + _draw_schedule(ax, frames) + for h_idle, t_end in drains: + ax.axvspan(h_idle, t_end, color=AMBER, alpha=0.13, zorder=1) + ax.text((drains[0][0] + drains[0][1]) / 2, 4.15, "barrier — H2D starves", + color=AMBER, fontsize=7.5, ha="center", va="bottom") ax.set_ylim(-0.4, 4.9) # --- opt3: no barriers, continuous ax = axes[2] - for i in range(11): - frame(ax, 3 - (i % 4) * 1.0, i * 11) + frames, _ = _schedule(11, 4) + _draw_schedule(ax, frames) ax.set_ylim(-1.5, 4.5) - ax.text(0, -0.25, "streams never wait on each other — the GPU is continuously busy", + ax.text(0, -0.25, "streams never wait on each other — the H2D engine never goes idle", color=ACCENT, fontsize=7.5, va="top") - titles = ["opt1 · 1 stream, synchronous", "opt2 · 4 streams, sync barrier per round", "opt3 · 4 streams, barriers removed"] for ax, t in zip(axes, titles): - ax.set_xlim(-2, 190) + ax.set_xlim(-2, 178) ax.set_yticks([]); ax.set_xticks([]) bare(ax, keep=()) ax.set_title(t, color=TEXT2, fontsize=9, loc="left", pad=4) @@ -464,7 +517,11 @@ def fig_streams(): fontsize=8, labelcolor=TEXT2, ncol=3, loc="lower right", bbox_to_anchor=(1.02, 0.98), handlelength=1.1) axes[2].set_xlabel("time →", color=MUTED, fontsize=8.5, loc="left") - fig.subplots_adjust(hspace=0.75) + fig.subplots_adjust(hspace=0.75, bottom=0.10) + fig.text(0.5, 0.005, + "Scheduled, not sketched: H2D and D2H are one FIFO engine each, as in hardware — measured overlap 1.00 on both — while " + "kernels may\noverlap. 3×3 proportions [f64 · s1], 13 : 15 : 5. The H2D lane fills first because at 3×3 H2D is the tallest bar.", + color=MUTED, fontsize=6.4, ha="center", va="bottom", linespacing=1.6) save(fig, "fig_streams") @@ -597,7 +654,7 @@ def fig_resultpath(): (a2, "9×9 · 467 kB / frame", 30.01, 40.0, "×2.21", "copy is larger than the GPU\n→ cannot hide at any overlap", AMBER), ]: - ax.bar([0], [floor], width=0.5, color=ACCENT, zorder=3, linewidth=0) + ax.bar([0], [floor], width=0.5, color=col, zorder=3, linewidth=0) ax.bar([1], [copy_us], width=0.5, color=col, zorder=3, linewidth=0) ax.axhline(floor, color=GREEN, lw=1.3, ls="--", zorder=4) ax.text(-0.55, floor + 1.2, "GPU floor", color=GREEN, fontsize=8.5, ha="left") @@ -622,39 +679,59 @@ def fig_resultpath(): # ------------------------------------------------ 8. f32 kernel (nsys truth) def fig_f32_kernel(): - fig, (ax, ax2) = plt.subplots(1, 2, figsize=(7.7, 2.6), - gridspec_kw={"width_ratios": [1, 1.35]}) - v = [39.93, 23.67] - ax.bar([0, 1], v, width=0.5, color=[AMBER, ACCENT], zorder=3) - for i, val in enumerate(v): - ax.text(i, val + 1.2, f"{val:.1f} µs", ha="center", color=PALE, - fontsize=11, fontweight="bold") - ax.annotate("", xy=(1, 25.5), xytext=(0, 41.5), - arrowprops=dict(arrowstyle="-|>", color=MUTED, lw=1.2, - connectionstyle="arc3,rad=-0.25")) - ax.text(0.5, 34, "−40.7 %", ha="center", color=PALE, fontsize=10, - fontweight="bold") - ax.set_xticks([0, 1]); ax.set_xticklabels(["f64 pedestal", "f32 pedestal"], - color=TEXT2, fontsize=8.5) - ax.set_ylim(0, 50); ax.set_yticks([]); bare(ax, keep=("bottom",)) - ax.set_title("kernel, exclusive (nsys, 9×9, 1 stream)", color=MUTED, - fontsize=8, pad=8) + """The same three engines, read at s1 and at s4 — cap 1700, both arms. + s1 is the DURATION claim: overlap is exactly 1.00 on every engine, so the + union per frame IS the mean duration, and the kernel is tallest in BOTH + arms. s4 is the BINDING claim: the copy engines still serialize (overlap + 1.00) but the kernel overlaps itself across streams, so its per-frame + number is occupancy. Only at s4 does the f32 kernel fall below D2H, which + is why the left panel cannot be used to argue the handover. + """ + fig, (ax1, ax2) = plt.subplots(1, 2, figsize=(8.0, 3.15)) labels = ["kernel", "D2H", "H2D"] - f64 = [39.93, 19.44, 13.24] - f32 = [23.67, 19.44, 13.24] - y = np.arange(3); h = 0.35 - ax2.barh(y + h / 2, f64, height=h, color=AMBER, zorder=3, label="f64 ped") - ax2.barh(y - h / 2, f32, height=h, color=ACCENT, zorder=3, label="f32 ped") - for yi, (a, b) in enumerate(zip(f64, f32)): - ax2.text(a + 1, yi + h / 2, f"{a:.1f}", va="center", color=TEXT2, fontsize=8) - ax2.text(b + 1, yi - h / 2, f"{b:.1f}", va="center", color=TEXT2, fontsize=8) - ax2.set_yticks(y); ax2.set_yticklabels(labels, color=TEXT2, fontsize=8.5) - ax2.invert_yaxis(); ax2.set_xlim(0, 52); ax2.set_xticks([]) - bare(ax2, keep=("left",)) - ax2.legend(frameon=False, fontsize=8, labelcolor=TEXT2, loc="lower right") - ax2.set_title("per-frame GPU operations (µs) — only the kernel moves", - color=MUTED, fontsize=8, pad=8) + y = np.arange(3); h = 0.34 + panels = [ + (ax1, "s1 · UNCONTENDED — duration", + [39.86, 21.97, 13.20], [23.70, 21.95, 13.22], "−40.5 %", + "kernel binds in BOTH arms"), + (ax2, "s4 · SHIPPED — what binds", + [32.66, 25.25, 20.77], [23.94, 25.24, 20.54], "−26.7 %", + "f32 kernel drops BELOW D2H"), + ] + for ax, title, f64, f32, delta, verdict in panels: + ax.barh(y + h / 2, f64, height=h, color=AMBER, zorder=3, label="f64 ped") + ax.barh(y - h / 2, f32, height=h, color=ACCENT, zorder=3, label="f32 ped") + # the engine that sets the floor in each arm is named, not ringed + i64, i32 = int(np.argmax(f64)), int(np.argmax(f32)) + for yi, (a, b) in enumerate(zip(f64, f32)): + for val, off, top in ((a, +h / 2, yi == i64), (b, -h / 2, yi == i32)): + # the kernel row carries the argument, so it stays legible even + # in the arm where it no longer binds + ax.text(val + 0.7, yi + off, f"{val:.2f}" + (" ◀ binds" if top else ""), + va="center", color=PALE if (top or yi == 0) else TEXT2, + fontsize=7.5, fontweight="bold" if top else "normal") + ax.text(0.985, 0.055, verdict, transform=ax.transAxes, ha="right", + color=PALE, fontsize=8, fontweight="bold") + ax.text(f64[0] * 0.42, -0.02 - h, delta, ha="center", va="center", + color=PALE, fontsize=9.5, fontweight="bold") + ax.set_yticks(y); ax.set_yticklabels(labels, color=TEXT2, fontsize=8.5) + ax.invert_yaxis(); ax.set_xlim(0, 46); ax.set_xticks([]) + ax.set_ylim(2.62, -0.72) + bare(ax, keep=("left",)) + ax.set_title(title, color=PALE, fontsize=9, pad=10, loc="left") + + # the wall the f32 kernel has to clear, drawn only where it matters + ax2.axvline(25.24, color=MUTED, lw=1.0, ls="--", zorder=4) + ax1.legend(frameon=False, fontsize=8, labelcolor=TEXT2, loc="lower right", + bbox_to_anchor=(1.0, 0.13)) + fig.subplots_adjust(bottom=0.20, top=0.86, left=0.09, right=0.99, wspace=0.22) + fig.text(0.545, 0.015, + "s4 is busy time per frame — the UNION of an engine's intervals, not a duration. The copy engines " + "still serialize (overlap 1.00), so H2D and D2H\nstay durations and rise under contention; the kernel " + "overlaps itself 1.32× across streams, so its 43.2 µs mean duration reads as 32.66 µs of " + "occupancy. → A1", + ha="center", va="bottom", color=MUTED, fontsize=6.8, linespacing=1.6) save(fig, "fig_f32_kernel") @@ -662,33 +739,57 @@ def fig_f32_kernel(): def fig_cancellation(): fig, (ax, ax2) = plt.subplots(1, 2, figsize=(7.7, 2.7), gridspec_kw={"width_ratios": [1.25, 1]}) - names = ["E[X²]\n2.17e7", "mean²\n2.17e7", "variance\n2025"] - vals = [2.17e7, 2.17e7, 2025] - ax.bar([0, 1], vals[:2], width=0.5, color=[PALE, PALE], zorder=3) + names = ["E[X²]\n2.17e7", "mean²\n2.17e7", "var, bulk\n2025", "var, quiet\n9"] + ax.bar([0, 1], [2.17e7, 2.17e7], width=0.5, color=[PALE, PALE], zorder=3) ax.bar([2], [2025], width=0.5, color=AMBER, zorder=3) - ax.set_yscale("log"); ax.set_ylim(1e2, 2e8) - ax.set_xticks([0, 1, 2]); ax.set_xticklabels(names, color=TEXT2, fontsize=8) - ax.set_yticks([1e3, 1e5, 1e7]) - ax.axhline(2048, color=ACCENT, lw=1.3, ls="--", zorder=4) - ax.text(2.42, 3000, "f32 rounding step\nat 2.17e7 = 2048", color=ACCENT, - fontsize=7.5, ha="right", va="bottom") + ax.bar([3], [9], width=0.5, color=AMBER, zorder=3) + ax.set_yscale("log"); ax.set_ylim(1, 2e8) + ax.set_xticks([0, 1, 2, 3]); ax.set_xticklabels(names, color=TEXT2, fontsize=7.5) + ax.set_yticks([1e0, 1e2, 1e4, 1e6, 1e8]) + ax.axhline(3, color=ACCENT, lw=1.3, ls="--", zorder=4) + ax.text(3.46, 2.2e5, "the error: ±3 ADU², the same\nfor every pixel — so it is the\n" + "quiet ones it swallows", color=ACCENT, fontsize=7.5, ha="right", + va="center", linespacing=1.5) + ax.plot([3.0, 3.0], [3.6, 4e4], color=ACCENT, lw=0.7, ls=":", zorder=4) bare(ax) - ax.set_title("var = E[X²] − mean² (f32, mean ≈ 4655 ADU)", + ax.set_title("var = E[X²] − mean² · each operand on a grid of 2 ADU²", color=MUTED, fontsize=8, pad=8) - rms = np.linspace(0, 12, 200) - ax2.fill_between(rms, 0, np.where(rms < 6.5, 1, 0), color=AMBER, alpha=0.16, - step="pre") - ax2.plot(rms, rms**2, color=PALE, lw=1.8, label="true variance") - ax2.axhline(42, color=ACCENT, lw=1.4, ls="--", label="f32 error floor") + # The axes are scaled to the region the error actually reaches: the floor is + # +-3-4 ADU^2 (docs/pedestal_precision_f32_cancellation.md SS5), so variance is + # lost below rms ~2 and merely corrupted up to rms ~5. An earlier version drew + # the floor at 42 with the damage running to rms 6.5 -- that is 6.5^2, i.e. the + # line and the shading were each other's source, and both were ~14x too large. + rms = np.linspace(0, 6, 300) + ax2.axvspan(0, 2.0, color=AMBER, alpha=0.20, zorder=1) + ax2.axvspan(2.0, 5.0, color=AMBER, alpha=0.07, zorder=1) + ax2.fill_between([0, 6], 3, 4, color=ACCENT, alpha=0.22, zorder=2, lw=0) + ax2.plot(rms, rms**2, color=PALE, lw=1.8, zorder=4) + ax2.axhline(3, color=ACCENT, lw=1.4, ls="--", zorder=5) + # the three rows of the doc's per-pixel table, labelled left of the curve + # (it is convex, so the whole upper-left of the panel is free) + pts = ((2.0, "rms 2 → var 4 ± 3 → 0", 0.30, 11.2), + (3.0, "rms 3 → 17 % rms err", 0.30, 17.4), + (5.0, "rms 5 → 6 % rms err", 2.30, 24.6)) + for r, lab, tx, ty in pts: + ax2.plot([r], [r * r], marker="o", ms=4, color=PALE, zorder=6) + ax2.annotate(lab, xy=(r, r * r), xytext=(tx, ty), color=TEXT2, + fontsize=6.8, va="center", zorder=6, + arrowprops=dict(arrowstyle="-", color=MUTED, lw=0.6, + shrinkA=2, shrinkB=3)) ax2.set_xlabel("pixel rms (ADU)", color=TEXT2, fontsize=8.5) - ax2.set_ylabel("variance", color=TEXT2, fontsize=8.5) - ax2.set_ylim(0, 150); ax2.set_xlim(0, 12) - ax2.set_yticks([]); ax2.tick_params(labelsize=8) - bare(ax2) - ax2.text(1.0, 108, "quiet pixels:\nerror > variance\n→ rms clamped to 0\n→ fires every frame", - color=AMBER, fontsize=7.5, va="top") - ax2.legend(frameon=False, fontsize=7.5, labelcolor=TEXT2, loc="lower right") + ax2.set_ylabel("variance (ADU²)", color=TEXT2, fontsize=8.5) + ax2.set_ylim(0, 40); ax2.set_xlim(0, 6) + ax2.set_yticks([0, 10, 20, 30, 40]); ax2.set_xticks([0, 2, 4, 6]) + ax2.tick_params(labelsize=7.5, colors=MUTED) + bare(ax2, keep=("left", "bottom")) + ax2.text(0.12, 38.8, "rms < 2\n→ clamped to 0\n→ fires every frame", + color=AMBER, fontsize=7.0, va="top", linespacing=1.5, zorder=6) + ax2.text(2.20, 38.8, "rms 2–5: threshold\ncorrupted, not clamped", + color=MUTED, fontsize=6.8, va="top", linespacing=1.5, zorder=6) + ax2.text(3.05, 31.5, "true variance = rms²", color=PALE, fontsize=7.2, zorder=6) + ax2.text(5.90, 5.3, "f32 error floor ±3–4 ADU²", color=ACCENT, fontsize=7.0, + ha="right", zorder=6) save(fig, "fig_cancellation") @@ -1048,3 +1149,240 @@ def fig_regpressure(): fig_regpressure() + + +# ------------------------------------------ 4b. the timeline, revealed in steps +# fig_streams shows all three stages at once and belongs on the opt3 slide, where +# removing the barrier is the point. These two are the earlier beats of the same +# picture, so opt1 and opt2 can each show the state of play at their own step. + +def _engine_legend(ax, y=0.92): + handles = [Rectangle((0, 0), 1, 1, color=c) for c in (AMBER, ACCENT, PALE)] + ax.legend(handles, ["H2D copy", "kernel", "D2H copy"], frameon=False, + fontsize=8, labelcolor=TEXT2, ncol=3, loc="lower right", + bbox_to_anchor=(1.02, y), handlelength=1.1) + + +def fig_opt1_timeline(): + """opt1 — one stream, synchronous: one engine at a time, host idle between. + + Proportions are opt1's own: 13.1 + 14.7 + 5.3 = 33.2 us of serialized engine + time inside a 63.3 us frame, so the host gap is drawn ~48 % of the period. The + contiguous-bars version of this picture (fig_streams, panel 1) implies the GPU + is busy end to end, which is exactly what opt1 is not. + """ + WORK = H_ + K_ + D_ # 33 units = 33.2 us of engine work + HOST = 30 # 63.3 - 33.2 = 30.1 us the host holds + PER = WORK + HOST + fig, ax = plt.subplots(figsize=(7.7, 1.05)) + for i in range(3): + t0 = i * PER + _frame_bars(ax, 1.0, t0) + ax.axvspan(t0 + WORK, t0 + PER, color=MUTED, alpha=0.09, zorder=1) + ax.annotate("host blocks — every engine idle", + xy=(WORK + HOST / 2, 0.98), xytext=(WORK + HOST / 2, 0.56), + color=AMBER, fontsize=8, ha="center", va="top", + arrowprops=dict(arrowstyle="-", color=AMBER, lw=0.8)) + ax.text(3 * PER + 4, 1.34, "one engine\nat a time", color=TEXT2, fontsize=8, + va="center", linespacing=1.5) + ax.set_xlim(-4, 3 * PER + 40) + ax.set_ylim(0.02, 2.45) + ax.set_yticks([]); ax.set_xticks([]) + bare(ax, keep=()) + _engine_legend(ax) + fig.subplots_adjust(left=0.01, right=0.99, top=0.86, bottom=0.04) + save(fig, "fig_opt1_timeline") + + +def fig_opt2_timeline(): + """opt2 — four streams, ONE round: what streaming buys, and nothing else. + + Scheduled by _schedule(), so the four H2D bars queue on the single copy + engine instead of being drawn on top of each other. The stagger is therefore + not a drawing choice: it is exactly H2D's duration, which is why the lanes + step by 13 units. The barrier is opt3's subject and is left to that slide. + """ + fig, ax = plt.subplots(figsize=(7.7, 1.50)) + frames, _ = _schedule(4, 4) + _draw_schedule(ax, frames) + for st in range(4): + ax.text(-4, (3 - st) + LANE_ / 2, f"stream {st}", color=MUTED, + fontsize=7.5, ha="right", va="center") + # The window where the most frames are simultaneously in flight -- measured + # off the schedule, not asserted. At 3x3 proportions it is THREE, not four: + # the frame span (33) is only 2.5x the H2D stagger (13), so stream 0 has + # already retired by the time stream 3 gets the copy engine. The old drawing + # claimed all four, which the single copy engine makes impossible here. + spans = [(h0, d0 + D_) for _, h0, _, d0 in frames] + edges = sorted({t for sp in spans for t in sp}) + counts = [(a, b, sum(1 for s0, s1 in spans if s0 <= a and s1 >= b)) + for a, b in zip(edges, edges[1:])] + best = max(c for _, _, c in counts) + lo = min(a for a, _, c in counts if c == best) + hi = max(b for _, b, c in counts if c == best) + ax.axvspan(lo, hi, color=ACCENT, alpha=0.10, zorder=1) + ax.text((lo + hi) / 2, 4.05, f"{best} frames in flight", color=ACCENT, + fontsize=7.5, ha="center", va="bottom") + ax.set_xlim(-26, 100) + ax.set_ylim(-0.35, 4.6) + ax.set_yticks([]); ax.set_xticks([]) + bare(ax, keep=()) + ax.text(74, 0.34, "a copy in one stream runs\nwhile another computes", + color=TEXT2, fontsize=8, va="center", linespacing=1.5) + _engine_legend(ax, y=0.94) + ax.set_xlabel("time →", color=MUTED, fontsize=8.5, loc="left") + fig.subplots_adjust(left=0.01, right=0.99, top=0.90, bottom=0.19) + fig.text(0.01, 0.015, + "Scheduled, not sketched: the H2D bars queue because there is one copy engine per direction (measured overlap 1.00). 3×3 proportions " + "[f64 · s1], 13 : 15 : 5 — so a frame's span is 2.5× the H2D stagger and three are in flight, not four.", + color=MUTED, fontsize=6.4, va="bottom") + save(fig, "fig_opt2_timeline") + + +fig_opt1_timeline() +fig_opt2_timeline() + + +# ------------------------------- 10b. the same typedef, the same absolute gain +def fig_f32_absolute(): + """The same typedef, the same ~4.7 us -- and a percentage that triples. + + 9x9 end-to-end, f64 vs f32, warm, cap 1700, from ladder_9x9.csv in + results/2026-08-20_{f64,f32}_cap1700/. + + opt3 is EXCLUDED, not dropped for space: its two arms differ by 393k page + faults, which the deck's own 0.68 us/fault model turns into +13.4 us against + an observed +13.2. That reading measured the allocator, not the kernel, and + quoting it as "+16 %" would repeat the error this slide exists to expose. + + opt5 is shown but greyed: its arms differ by 141k faults (-4.81 us predicted + against -4.54 observed), so the fault term alone accounts for the whole + effect. opt4 (faults matched to 493) and opt6 (zero faults in both arms) are + the two clean readings -- and they agree, -4.63 and -4.87 us. + + Two panels because one cannot carry it: a 4.6 us delta on an 80 us axis is + invisible, which is itself the point. Left = the frame shrinking; right = + the saving that does not. + """ + steps = ["opt4", "opt5", "opt6"] + sub = ["+ pinned input", "host↔GPU overlap", "zero-copy"] + f64v = [79.83, 66.39, 30.01] + f32v = [75.20, 61.85, 25.14] + dv = [a - b for a, b in zip(f64v, f32v)] + pct = [-5.8, -6.8, -16.2] + dagger = [False, True, False] + + fig, (axA, axB) = plt.subplots( + 1, 2, figsize=(11.4, 3.15), gridspec_kw={"width_ratios": [1.5, 1]}) + x = np.arange(3) + + # ---- A: the frame, shrinking + w = 0.34 + axA.bar(x - w / 2 - 0.015, f64v, width=w, color=ACCENT, zorder=3, linewidth=0) + axA.bar(x + w / 2 + 0.015, f32v, width=w, color=AMBER, zorder=3, linewidth=0) + for i, (a, b) in enumerate(zip(f64v, f32v)): + axA.text(i - w / 2 - 0.015, a + 1.6, f"{a:.1f}", ha="center", + color=TEXT2, fontsize=8.5) + axA.text(i + w / 2 + 0.015, b + 1.6, f"{b:.1f}", ha="center", + color=TEXT2, fontsize=8.5) + axA.annotate("opt6 + f32 = opt7\nthe shipped build", + xy=(2.19, 19.0), xytext=(2.36, 56), + arrowprops=dict(arrowstyle="->", color=AMBER, lw=1.2), + color=AMBER, fontsize=8.5, fontweight="bold", linespacing=1.5, + ha="center") + axA.set_xlim(-0.62, 3.02) + axA.set_ylim(0, 94) + axA.set_yticks([0, 25, 50, 75]) + axA.set_yticklabels(["0", "25", "50", "75"], fontsize=8) + axA.set_ylabel("end-to-end µs / frame", color=MUTED, fontsize=8.5) + axA.set_xticks(x) + axA.set_xticklabels([f"{s}\n{t}" for s, t in zip(steps, sub)], + fontsize=8.5, linespacing=1.6, color=TEXT2) + axA.tick_params(axis="x", length=0, pad=7) + bare(axA, keep=("left", "bottom")) + axA.spines["bottom"].set_color(RULE) + handles = [Rectangle((0, 0), 1, 1, color=c) for c in (ACCENT, AMBER)] + axA.legend(handles, ["f64 pedestal", "f32 pedestal"], frameon=False, + fontsize=8.5, labelcolor=TEXT2, ncol=2, loc="upper right", + bbox_to_anchor=(1.02, 1.10), handlelength=1.1) + axA.set_title("the frame shrinks by 2.7×", color=MUTED, fontsize=9, + pad=14, loc="left") + + # ---- B: the saving, which does not + axB.bar(x, dv, width=0.52, color=AMBER, zorder=3, linewidth=0) + axB.axhline(np.mean(dv), color=PALE, lw=0.9, ls="--", zorder=4) + for i, (d, p_) in enumerate(zip(dv, pct)): + axB.text(i, d + 0.24, f"−{d:.2f} µs", ha="center", color=PALE, + fontsize=10.5, fontweight="bold") + axB.set_ylim(0, 7.6) + axB.set_yticks([0, 2, 4]) + axB.set_yticklabels(["0", "2", "4"], fontsize=8) + axB.set_ylabel("µs saved by the f32 pedestal", color=MUTED, fontsize=8.5) + axB.set_xticks(x) + axB.set_xticklabels([f"{s}{' †' if d else ''}\n{p:+.1f} %" + for s, d, p in zip(steps, dagger, pct)], + fontsize=9, linespacing=1.7, color=TEXT2) + axB.tick_params(axis="x", length=0, pad=7) + bare(axB, keep=("left", "bottom")) + axB.spines["bottom"].set_color(RULE) + axB.set_title(f"the saving does not · dashed = {np.mean(dv):.2f} µs mean", + color=MUTED, fontsize=9, pad=14, loc="left") + + fig.subplots_adjust(bottom=0.30, top=0.84, left=0.055, right=0.985, + wspace=0.26) + save(fig, "fig_f32_absolute") + + +fig_f32_absolute() + + +# ------------------------------------------- 10c. what the rewrite changes +def fig_variance_rewrite(): + """The rewrite in one axis: the SIZE of the numbers you subtract. + + Numbers from docs/pedestal_precision_f32_cancellation.md §4-5 and §11. + Before, var = E[X²] − mean² subtracts two ~2.17e7 operands to recover ~2025: + each operand sits on an f32 grid of 2 ADU², so the answer inherits ±3 — an + ABSOLUTE error that does not shrink for quiet pixels, which is what kills + them. After, with Y = X − X0 accumulated instead, the operands are 2025 and + ~0.25; the grid under them is 1.2e-4 and the cancellation is simply gone. + """ + ANS = 2025.0 + rows = [(1.0, 2.17e7, "2.17 × 10⁷", "±3 ADU² — fatal below rms ≈ 2", MUTED), + (0.0, ANS, "2.02 × 10³", "±0.0001 ADU² — 30 000× smaller", AMBER)] + + fig, ax = plt.subplots(figsize=(7.7, 1.36)) + ax.set_xscale("log") + for y, operand, mag, note, col in rows: + if operand > ANS: + ax.plot([ANS, operand], [y, y], color=col, lw=11, alpha=0.5, + solid_capstyle="butt", zorder=3) + ax.plot([operand], [y], "o", color=col, ms=9, zorder=5) + ax.text(operand * 2.2, y + 0.17, f"operands {mag}", color=TEXT2, + fontsize=8, va="center") + ax.text(operand * 2.2, y - 0.19, note, color=col, fontsize=8, + va="center", fontweight="bold") + ax.axvline(ANS, color=PALE, lw=1.2, ls="--", zorder=4) + ax.text(ANS * 1.25, -0.44, "the answer: variance ≈ 2025", color=PALE, + fontsize=7.5, ha="left", va="center") + ax.annotate("", xy=(2.17e7, 1.44), xytext=(ANS, 1.44), + arrowprops=dict(arrowstyle="<->", color=MUTED, lw=1.0)) + ax.text(2.1e5, 1.52, "4 decades of common term to cancel", color=MUTED, + fontsize=7.5, ha="center") + ax.set_yticks([1, 0]) + ax.set_yticklabels(["before\naccumulate X", "after\naccumulate Y = X − X₀"], + fontsize=8.5, linespacing=1.5) + for lab, c in zip(ax.get_yticklabels(), (TEXT2, PALE)): + lab.set_color(c) + ax.tick_params(axis="y", length=0) + ax.set_xlim(3e2, 4e9) + ax.set_ylim(-0.62, 1.78) + ax.set_xticks([1e3, 1e5, 1e7, 1e9]) + ax.tick_params(axis="x", labelsize=7.5) + bare(ax, keep=("bottom",)) + ax.spines["bottom"].set_color(RULE) + fig.subplots_adjust(left=0.19, right=0.99, top=0.97, bottom=0.20) + save(fig, "fig_variance_rewrite") + + +fig_variance_rewrite() diff --git a/docs/figures/fig_arc.png b/docs/figures/fig_arc.png index f0418cf5..7bbe2341 100644 Binary files a/docs/figures/fig_arc.png and b/docs/figures/fig_arc.png differ diff --git a/docs/figures/fig_arc_9x9.png b/docs/figures/fig_arc_9x9.png index 2bf44ec2..830dda02 100644 Binary files a/docs/figures/fig_arc_9x9.png and b/docs/figures/fig_arc_9x9.png differ diff --git a/docs/figures/fig_bottleneck.png b/docs/figures/fig_bottleneck.png index 0470265f..2b355d93 100644 Binary files a/docs/figures/fig_bottleneck.png and b/docs/figures/fig_bottleneck.png differ diff --git a/docs/figures/fig_cancellation.png b/docs/figures/fig_cancellation.png index 56aad03e..a3667984 100644 Binary files a/docs/figures/fig_cancellation.png and b/docs/figures/fig_cancellation.png differ diff --git a/docs/figures/fig_correctness.png b/docs/figures/fig_correctness.png index 0e86a1cc..a7a4d99c 100644 Binary files a/docs/figures/fig_correctness.png and b/docs/figures/fig_correctness.png differ diff --git a/docs/figures/fig_f32_absolute.png b/docs/figures/fig_f32_absolute.png new file mode 100644 index 00000000..3cc9d553 Binary files /dev/null and b/docs/figures/fig_f32_absolute.png differ diff --git a/docs/figures/fig_f32_kernel.png b/docs/figures/fig_f32_kernel.png index 65a4e3c3..142cdff1 100644 Binary files a/docs/figures/fig_f32_kernel.png and b/docs/figures/fig_f32_kernel.png differ diff --git a/docs/figures/fig_first_run.png b/docs/figures/fig_first_run.png index 0ad3017e..2576595c 100644 Binary files a/docs/figures/fig_first_run.png and b/docs/figures/fig_first_run.png differ diff --git a/docs/figures/fig_graphs.png b/docs/figures/fig_graphs.png index bacad131..4144eebd 100644 Binary files a/docs/figures/fig_graphs.png and b/docs/figures/fig_graphs.png differ diff --git a/docs/figures/fig_mismatch147.png b/docs/figures/fig_mismatch147.png index ff7a4a4c..28eef876 100644 Binary files a/docs/figures/fig_mismatch147.png and b/docs/figures/fig_mismatch147.png differ diff --git a/docs/figures/fig_opt1_timeline.png b/docs/figures/fig_opt1_timeline.png new file mode 100644 index 00000000..e995efa2 Binary files /dev/null and b/docs/figures/fig_opt1_timeline.png differ diff --git a/docs/figures/fig_opt2_timeline.png b/docs/figures/fig_opt2_timeline.png new file mode 100644 index 00000000..09b5c30f Binary files /dev/null and b/docs/figures/fig_opt2_timeline.png differ diff --git a/docs/figures/fig_overhead.png b/docs/figures/fig_overhead.png index 9525c11f..7fec7ec3 100644 Binary files a/docs/figures/fig_overhead.png and b/docs/figures/fig_overhead.png differ diff --git a/docs/figures/fig_overlap.png b/docs/figures/fig_overlap.png index d2918d75..39b58938 100644 Binary files a/docs/figures/fig_overlap.png and b/docs/figures/fig_overlap.png differ diff --git a/docs/figures/fig_pedtiming.png b/docs/figures/fig_pedtiming.png index ebb178c3..91ed8e76 100644 Binary files a/docs/figures/fig_pedtiming.png and b/docs/figures/fig_pedtiming.png differ diff --git a/docs/figures/fig_pinning.png b/docs/figures/fig_pinning.png index 7d7a7760..4fc6b01a 100644 Binary files a/docs/figures/fig_pinning.png and b/docs/figures/fig_pinning.png differ diff --git a/docs/figures/fig_regpressure.png b/docs/figures/fig_regpressure.png index 417715fa..4a803e9e 100644 Binary files a/docs/figures/fig_regpressure.png and b/docs/figures/fig_regpressure.png differ diff --git a/docs/figures/fig_resultpath.png b/docs/figures/fig_resultpath.png index a612f663..fe62fd54 100644 Binary files a/docs/figures/fig_resultpath.png and b/docs/figures/fig_resultpath.png differ diff --git a/docs/figures/fig_streams.png b/docs/figures/fig_streams.png index c00b1b77..c6eb1e32 100644 Binary files a/docs/figures/fig_streams.png and b/docs/figures/fig_streams.png differ diff --git a/docs/figures/fig_variance_rewrite.png b/docs/figures/fig_variance_rewrite.png new file mode 100644 index 00000000..9866dc62 Binary files /dev/null and b/docs/figures/fig_variance_rewrite.png differ diff --git a/docs/figures/img_gpu_die.png b/docs/figures/img_gpu_die.png new file mode 100644 index 00000000..6b652480 Binary files /dev/null and b/docs/figures/img_gpu_die.png differ diff --git a/docs/figures/img_gpu_sm.png b/docs/figures/img_gpu_sm.png new file mode 100644 index 00000000..5eb2e629 Binary files /dev/null and b/docs/figures/img_gpu_sm.png differ