detectors/aare
25
0
Fork 0
mirror of https://github.com/slsdetectorgroup/aare.git synced 2026-09-02 23:20:43 +02:00
Code Issues Actions Packages Releases Activity

docs: correct floor/peak definition, s1-vs-s4 evidence, and engine timelines
Build on RHEL8 / build (push) Successful in 3m19s
Details
Build on RHEL9 / build (push) Successful in 3m41s
Details
Run tests using data on local RHEL8 / build (push) Successful in 4m9s
Details

Browse Source 
Resolve a terminology collision where "floor" meant both max(H2D, kernel,
D2H) and the best sustained rate. They coincide at 3x3 (16.17) and diverge
at 9x9 f64 (32.66 profiled max vs 30.01 sustained), so the deck read as
self-contradictory only at 9x9. Unify on "floor" across the deck, split A1
into separate `engine max [s4]` and `FLOOR = lower of max, sustained` rows,
and document in the report that the two numbers come from two different
runs made by two different tools -- 30.01 is unprofiled ladder wall clock,
not an nsys number.

Rebuild fig_f32_kernel as paired s1/s4 panels: the old one was hardcoded to
cap-1500 s1 values and showed the kernel binding in both arms, refuting the
D2H-handover claim on the same slide. Move that conclusion to 22/34, the end
of the Act III arc.

Redraw the opt2/opt3 timelines from a real schedule honouring one copy
engine per direction (measured overlap 1.00), with 3x3 proportions instead
of unsourced ones. Retires the "all four in flight" claim -- three are.

Redraw fig_cancellation's error floor from 42 to the documented +-3-4 ADU2;
42 was 6.5 squared, i.e. the line and the shading were each other's source.

Normalise 9x9 numbers to cap 1700 and fix eight cap-1500 residues in the
report, plus the per-instance kernel figure (43.2 us at s4, not 39.9).
This commit is contained in:
kferjaoui
2026-08-24 20:57:24 +02:00
parent 4c0a093e9f
commit d14b2c0088
28 changed files with 1047 additions and 340 deletions
Download Patch File Download Diff File Expand all files Collapse all files
docs/ClusterFinderCUDA_benchmark_results.md
+97 -25
View File
@@ -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]`
docs/cf_cuda_fused.pptx
BIN
View File
Binary file not shown.
docs/cf_cuda_kernel.pptx
BIN
View File
Binary file not shown.
docs/deck/CHANGELOG_2026-08-24.md
+154
View File
@@ -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.
docs/deck/build_fused_deck.py
+378 -235
View File
@@ -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<ClusterType, FRAME_TYPE>",
" <<<grid, block, shmem>>>(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<<<grid, block, shmem, «sc.stream»>>>(...);",
"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<<<grid, block, shmem, «sc.stream»>>>(...);",
"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)")
docs/deck/make_figs.py
+418 -80
View File
@@ -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()
docs/figures/fig_arc.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 116 KiB

After

Width:  |  Height:  |  Size: 127 KiB

docs/figures/fig_arc_9x9.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 113 KiB

After

Width:  |  Height:  |  Size: 124 KiB

docs/figures/fig_bottleneck.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 71 KiB

After

Width:  |  Height:  |  Size: 79 KiB

docs/figures/fig_cancellation.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 73 KiB

After

Width:  |  Height:  |  Size: 111 KiB

docs/figures/fig_correctness.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 36 KiB

After

Width:  |  Height:  |  Size: 41 KiB

docs/figures/fig_f32_absolute.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 100 KiB

docs/figures/fig_f32_kernel.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 51 KiB

After

Width:  |  Height:  |  Size: 106 KiB

docs/figures/fig_first_run.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 124 KiB

After

Width:  |  Height:  |  Size: 138 KiB

docs/figures/fig_graphs.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 86 KiB

After

Width:  |  Height:  |  Size: 94 KiB

docs/figures/fig_mismatch147.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 80 KiB

After

Width:  |  Height:  |  Size: 85 KiB

docs/figures/fig_opt1_timeline.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 21 KiB

docs/figures/fig_opt2_timeline.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 58 KiB

docs/figures/fig_overhead.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 75 KiB

After

Width:  |  Height:  |  Size: 80 KiB

docs/figures/fig_overlap.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 62 KiB

After

Width:  |  Height:  |  Size: 67 KiB

docs/figures/fig_pedtiming.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 106 KiB

After

Width:  |  Height:  |  Size: 118 KiB

docs/figures/fig_pinning.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 77 KiB

After

Width:  |  Height:  |  Size: 84 KiB

docs/figures/fig_regpressure.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 66 KiB

After

Width:  |  Height:  |  Size: 72 KiB

docs/figures/fig_resultpath.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 79 KiB

After

Width:  |  Height:  |  Size: 86 KiB

docs/figures/fig_streams.png
BIN
View File
Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 54 KiB

After

Width:  |  Height:  |  Size: 88 KiB

docs/figures/fig_variance_rewrite.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 56 KiB

docs/figures/img_gpu_die.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 2.3 MiB

docs/figures/img_gpu_sm.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 864 KiB