diff --git a/docs/cf_cuda_fused.pptx b/docs/cf_cuda_fused.pptx index 800444fc..080c09de 100644 Binary files a/docs/cf_cuda_fused.pptx and b/docs/cf_cuda_fused.pptx differ diff --git a/docs/deck/CHANGELOG_2026-08-25.md b/docs/deck/CHANGELOG_2026-08-25.md new file mode 100644 index 00000000..82c65d22 --- /dev/null +++ b/docs/deck/CHANGELOG_2026-08-25.md @@ -0,0 +1,304 @@ +# 2026-08-25 — projection legibility, and the 30-minute restructure + +Session log for `docs/cf_cuda_fused.pptx` (via `docs/deck/build_fused_deck.py` + +`make_figs.py`). Deck rebuilt at **51 slides**: title + hero + **33 numbered** +(3–33) + 7 dividers + **11 annex** (A1–A6). + +--- + +## 1. One projection floor, enforced in two places + +The back row is 6–7 m from the screen. On a 13.33 × 7.5 in slide, **9 pt is +about 1/60 of the slide height** — the conventional lower bound for supporting +detail. That is now the floor for *everything*, on the slide and inside every +PNG, and it is enforced by code rather than by hand. + +**Text set in PowerPoint** — `run()` clamps to `MIN_PT = 9.0`. Enforced in the +helper, not at the ~400 call sites, because a floor applied by hand is a floor +one new caption silently drops through. Consequences: table headers 8 → 9, +statstrip labels 8 → 9, rail labels 8.5 → 9, code titles 7.5 → 9, page numbers +8.5 → 9, the 8.5 pt captions → 9. + +`code()` had a **constant** line height of 0.148 in tuned for 8.5 pt, so raising +the type pushed the last line out of the panel. It now scales: `lh = 0.0174 × +size`, which reproduces the old value at 8.5. + +**Text inside figures** — the real problem, and it was invisible where the font +size is written. A matplotlib fontsize is in points of the *figure's* inches; +the PNG is then placed at some other width, so what the audience reads is + +``` +effective_pt = raw_pt × (placement_width_in / figure_width_in) +``` + +`make_figs.py`'s `save()` now takes the placement width, walks every `Text` +object in the figure after drawing it, and prints/collects anything below the +floor. `legibility_report()` at the end of the run is the gate. The placement +widths are **parsed out of `build_fused_deck.py`** rather than hardcoded, so the +check can never drift from the layout it is checking. + +Measured before the fix: + +| figure | PNG | placed | scale | smallest text | +|---|--:|--:|--:|--:| +| `fig_mismatch147` | 8.97 in | 9.40 | 1.05 | **4.8 pt** | +| `fig_opt2_timeline` | 10.64 in | 7.90 | 0.74 | **4.8 pt** | +| `fig_regpressure` | 10.71 in | 7.90 | 0.74 | **6.3 pt** | +| `fig_f32_kernel` | 8.00 in | 7.50 | 0.94 | **6.4 pt** | +| `fig_streams` | 6.36 in | 6.55 | 1.03 | **6.6 pt** | + +13 of 22 placed figures failed. All 22 now pass. + +Three causes, three fixes: + +- **Self-inflating captions.** `savefig(bbox_inches="tight")` grew the saved + width to fit a long in-figure caption, which shrank the scale factor, which + shrank the caption further — the longer the text, the smaller it rendered. + Four such strings (`fig_opt2_timeline`, `fig_streams`, `fig_f32_kernel`, + `fig_graphs`) were provenance or verdict text, not data labels. They moved to + slide captions and speaker notes, where a point is a point. +- **Declared width far above placement.** `fig_regpressure` 11.2 → 7.9 in, + `fig_varfloor` 5.6 → 4.35 in, `fig_graphs` re-laid out at 7.6 in. +- **Unfittable text.** `fig_mismatch147` printed an ADU value in each of its + masked-in cells at 4.6 pt. Dropped here, then **restored in §5** once the + figure was given a wider placement — see there. + +Then 75 individual sizes were raised — only those below their figure's floor, so +the internal hierarchy is preserved. + +**Layout fallout, all fixed:** slides 9, 10, 20, 25, 33 and A5 were re-laid out +where the larger type pushed content into a figure or past the footer. Verified +with `pdftotext -bbox` across all 51 pages: no text below the footer line except +the PSI template's own title page. + +--- + +## 2. Restructure for a 30-minute slot + +### New: the measurement convention (12/33 here, **moved to 18/33 in §5**) + +Promoted out of A1. Everything from slide 13 on is quoted as `[s1]` or `[s4]` +and compared against a "floor", and none of those three words was defined +anywhere the audience would see it. Three definition cards (s1 / s4 / floor) and +one figure; the numbers stay in A1. + +`fig_measure` is new: four stream lanes from the **real** `_schedule()`, so the +H2D bars queue because there is one copy engine per direction (measured overlap +1.000 in every row of `probes.csv`) while kernels sit on top of each other in +time. The two union strips beneath are computed off that schedule. No +microseconds on that panel — it reproduces the *shape* of the 9×9 pipeline, not +its exact overlap factor, and measured numbers next to a schematic invite being +read across. The right panel is measured: 20.77 / 32.66 / 25.25 at s4, the +dashed profiled max, and the 30.01 sustained rate that is the actual floor. + +### Reordered: 9 and 10 swapped + +Registers (the cause) now precede occupancy (the effect). The old order taught a +metric and then said it was not the point. Slide 10 opens with occupancy in one +sentence — *when a warp stalls the SM switches to another resident warp; +occupancy is how many alternatives it has* — so the 33 % lands as a consequence +of the register budget, not as a failing grade. + +### Merged: 20 + 21 → one slide + +The correctness trap and the fix. The ULP arithmetic of the error floor is the +least transferable part of the story; what stays is *a small answer computed as +the difference of two large numbers is not computable in f32*, what it did to +the physics, and the two-line change. + +`fig_cancellation`'s right panel is now the **deformed spectrum**. The f64 curve +is measured (`validation_tiers.json`, 23.2 M clusters). The f32 curve is +**reconstructed, and labelled as such on the slide**: its area is the measured ++28.06 % from §1 of the write-up, its shape is where §9 derives the excess (a +gate collapsed from ~225 ADU to ~0 admits the whole positive side, so the +population smears upward from zero). Nobody kept the broken build to re-run. + +The old right panel — variance vs rms, the ±3–4 ADU² floor — is now +`fig_varfloor`, in A5. + +### Cut to the annex + +| was | now | +|---|---| +| 18 · two rejected routes (B′, B″) | **A2·3** | +| 21 · the variance rewrite | **A5** (+ `fig_varfloor`) | +| 26 · three ways a benchmark lies | **A6·1**, as it stood | + +Slide 25 keeps only the artefact that moves a number the audience is shown — +first-touch page faults — plus one amber callout naming the other two (nsys +inflates host-side API calls ~4×, CUDA events measure a stream) and pointing at +A6. CUDA-event detail is not discussed in the arc at all. + +Slide 19 (opt7) lost the s1-vs-s4 caption, which the measurement slide now owns. + +### Detector context + +MÖNCH03 specifications on slide 3: hybrid silicon, 400 × 400 at 25 µm pitch, +10 × 10 mm² active area, **1.3 kHz standard and 3–6 kHz with optimised readout**. +Slide 33 closes the loop: 58 495 FPS is **45× the standard rate and ~10× the +optimised ceiling**, with the two qualifications in the notes (frames are assumed +already in host RAM; at 9×9 the margin shrinks if the cap has to grow). + +--- + +## 3. Renumbering + +`N_SLIDES` 34 → **33**, `N_ANNEX` 5 → **6**. All `chrome()` indices reassigned +sequentially; six divider slide-lists and ranges rewritten; every cross-reference +retargeted (`expands slide N` on A2·3, A3, A4, A5, A6; `[engines: 19/33]`; +"slides 27–30"; the A2 part counts 1/2, 2/2 → 1/3, 2/3). + +--- + +## 4. Verified + +- `pdftotext -bbox`, all 51 pages: no text past the footer (PSI title page aside). +- No unrendered `**` markup anywhere. +- `legibility_report()`: every string in every figure ≥ 9.0 pt as placed. +- No stale cap-1500 values (`82.17`, `80.44`, `94.78`, "18 % SLOWER"); the one + remaining "1 500" is the sentence explaining *why* the cap is 1 700. +- 20 em dashes on-slide, all pivots, appositions, or table "—" for empty. +- Visual pass over all 51 rendered pages. + +## Still open + +- `fig_bottleneck`, `fig_correctness` and `fig_variance_rewrite` are generated + but placed on no slide. Harmless; `fig_variance_rewrite` joined the list when + slide 21 moved to A5, where the code block says the same thing in less space. +- Timing: 40 slides in the arc. At 30 minutes that is ~45 s each, which is tight. + Next cut candidates, in order: 14/15 (opt2 and opt3 could be one slide), 24 + (where the time went — the arc slides already carry it), 26 (first run — the + point survives as a sentence on 25). + +--- + +## 5. Second review pass (same day) + +**Code-panel clearance, deck-wide.** The type floor grew every `code()` panel by +0.06–0.09 in, and eleven of them had less clearance than that to whatever was +drawn under. Found with a parser that computes each panel's height from its own +line count and size and compares against every later element with an overlapping +x-range, so this is not a slide-by-slide eyeball. All eleven now clear by +≥ 0.12 in; slides 8, 15 and 30 needed re-laying out around it. + +**12/33 moved into Act III, now 18/33.** It talked about four streams (which +arrive at opt2) and 9×9 (which is not discussed until Act III), so it was +spending the audience's attention before either was earned. It now sits +immediately after the Act III divider, ahead of opt7 — whose two-panel figure is +the first place s1 and s4 are put side by side. Retitled "How the engine times +are measured". + +Slide 12 (opt1) is where "% of floor" first appears, so it keeps a one-sentence +definition — *floor = max(H2D, kernel, D2H), the fastest a frame can go if the +host cost nothing* — and defers how each term is measured to 18. That callout +had been carrying the whole s1/s4/union argument; it is now four lines shorter. + +**One colour code for the two builds in Act III: f32 = amber, f64 = blue.** +`fig_f32_absolute` and `fig_cancellation` already were; `fig_f32_kernel` had them +the other way round, so the same two builds swapped colour between slides 19 and +20. (The engine palette — H2D amber, kernel blue, D2H white — is a separate axis +and unchanged; both figures carry their own legend.) + +**Titles.** 19 back to "FP32 halves pedestal traffic: −41 % kernel time". +20 → "Accumulate what is small, not what is large", with catastrophic +cancellation moved into the eyebrow; A5 renamed "The rewrite in full, and which +pixels the error reached" so the two no longer collide. The annex divider is +just "Annexes". + +**29/33: the per-cell ADU values are back.** They were dropped because at a 9.4 in +placement a 25 × 25 grid gives each cell 13.5 pt of width, which cannot hold a +three-character number legibly. The figure is now placed at **10.6 in** — 15.3 pt +per cell — which fits the values at **9.0 pt as rendered**, so they are back +*and* they clear the projection floor. The centre marker moved behind the text +(it used to paint over the value in its own cell) and the slide's caption moved +to the notes to pay for the extra height. + +**32/33: the knob that loses data is ringed.** New `frame_rect()` helper (an +outline, not a fill, so the row is not recoloured) and a new `RED` token used for +pointing only, never as a data colour. The bottom callout switched from amber to +the same red, so the ring and the warning read as one thing. + +**Also:** `fig_f32_kernel`'s "39.86 ◀ binds" ran past its own x-limit onto the +right panel's tick labels (xlim 46 → 60, wspace 0.22 → 0.30); `fig_pinning`'s DMA +arrow label was wider than the gap it sat in. + +Re-verified after every change: no text past the footer on any of the 51 pages, +no unrendered `**`, every figure string ≥ 9 pt as placed, and a visual pass over +all 51 rendered pages. + +--- + +## 6. Third review pass (same day) + +**19/33 · `fig_f32_kernel` label clearance.** A bar pair spans `y ± h`, so the +step percentage (`−40.5 %`, `−26.7 %`) at `−0.02 − h` sat exactly on the top edge +of the f32 kernel bar, and the verdict line at 5.5 % of axes height sat on the +H2D bars. Both now carry a full `h` of clearance beyond the bar pair, and the y +limits (`3.05, −0.98`) hold that margin explicitly rather than relying on default +padding — which the larger type had eaten. The legend moved to the one region no +label reaches, between the D2H and H2D rows, and the dashed 25.24 wall dropped +below the labels in z-order so it stops crossing the `23.94`. + +**18/33 · the floor's units, spelled out.** The card read "1 / the busiest +engine" while the panel under it was labelled in µs and the annotation said +"FLOOR 30.01" — reciprocal in one place, microseconds in the other, with nothing +saying they are the same quantity. They are: the busiest engine's busy time per +frame is µs/frame, and one divided by it is FPS. **30.01 µs/frame is 33 323 +FPS.** + +- Card retitled "Set by the busiest engine", body now names both units. +- The figure annotation reads `FLOOR 30.01 µs/frame = 33 323 FPS`. +- The slide's closing callout says the deck quotes the floor both ways. +- The speaker notes gain a paragraph on when each unit is used: µs/frame when + comparing engines to each other, FPS when comparing against the CPU baseline + or the detector's frame rate. `% of floor` is the same ratio either way. + +One consequence worth recording: the notes previously said "the floor is +1 / max(H2D, kernel, D2H) at s4 … 1 / 32.66 µs = 30 600 FPS", which mixed the two +forms in a single sentence and rounded loosely. It now reads "the max is 32.66 µs += 30 618 FPS", with the reciprocal taken once and named. + +Cards grew 1.52 → 1.70 in to hold the longer floor definition; `fig_measure` +re-placed at 10.3 in so the stack still clears the footer. + +--- + +## 7. The CPU die, mirroring the GPU die (5/33) + +Slide 6 had a die photograph with the compute units boxed; slide 5 had only a +block diagram, so the two machines were compared as schematics and never as +objects. Added `img_cpu_die.png` — Intel "Comet Lake" 10-core Core i9, same +CS149 source and the same crop recipe (`scratchpad/crops.py`: content box, light +card, rounded alpha edge). + +**Kept whole rather than cropped to the ten cores.** The L3 slab on the right and +the I/O block on the left are close to half the die area, which is the callout's +own argument — *"the ALUs are the small part"* — standing next to it in silicon. + +**On placement.** A strict mirror of slide 6 is not available: the AD102 die is +square (1.01:1) so it fits the narrow left slot, while the Comet Lake die is +2.35:1 and at any width where its core labels could be read it is 3.7 in tall. +Stacking it under the Skylake diagram at a shared width does not fit either — +both are landscape, and two of them plus the callout overruns the column by +~1.1 in. What did fit, without shrinking the Skylake diagram below the size at +which its `Fetch/Decode` and `ALU` labels survive projection: + +| | | +|---|---| +| Skylake core diagram | unchanged at 7.3 in, `M, 1.86` | +| callout | narrowed 7.9 → 4.55 in, `M, 5.42`, h 1.06 | +| **Comet Lake die** | **`5.60, 5.42`, 3.00 in wide** | + +The die's own text is illegible at 3.0 in, which is fine and matches slide 6: +neither die is read, both are counted. The caption carries the claim ("10 cores +boxed; nearly half the die is cache and I/O") exactly as slide 6's does ("144 +blocks, 128 enabled on this card. One SM boxed."). + +What moved to pay for it: the four-line Skylake-vs-Zen 4 caveat caption is now +speaker notes, along with a new paragraph on why ten cores do not fill the die. + +**Also on 6/33** (pre-existing, same class): the closing caption sat at 6.90 and +its second line rode on the progress track. Trimmed to one line at 6.82, with the +FP64-scarcity remark moved into the notes — where it now says explicitly that the +V100 diagram *understates* it, since a GeForce part runs FP64 at 1/64, and that +this is half of why opt7 pays. diff --git a/docs/deck/build_fused_deck.py b/docs/deck/build_fused_deck.py index 930d4790..7018e7b4 100644 --- a/docs/deck/build_fused_deck.py +++ b/docs/deck/build_fused_deck.py @@ -52,6 +52,7 @@ PALE = RGBColor(0xE7, 0xED, 0xF4) TEXT2 = RGBColor(0xA5, 0xB2, 0xC4) MUTED = RGBColor(0x6B, 0x7A, 0x90) CARD = RGBColor(0xF4, 0xF6, 0xF9) # light card for white figures +RED = RGBColor(0xE2, 0x54, 0x54) # pointing only, never a data colour UI, MONO = "Segoe UI", "Consolas" W, H = 13.333, 7.5 @@ -66,7 +67,7 @@ R = "{http://schemas.openxmlformats.org/officeDocument/2006/relationships}" prs = Presentation(str(BASE)) prs.slide_width, prs.slide_height = In(W), In(H) BLANK = prs.slide_layouts[0] # 'Blank Slide' — zero shapes -N_SLIDES = 34 +N_SLIDES = 33 # ------------------------------------------------------------------- helpers @@ -139,10 +140,24 @@ def para(tf, first=False, space_after=0, space_before=0, line=None, align=None): return p +# The projection floor, in points, for text set directly in PowerPoint. +# +# The deck is shown in a room where the back row is 6-7 m from the screen. On a +# 13.33 x 7.5 in slide, 9 pt is about 1/60 of the slide height, which is the +# conventional lower bound for supporting detail. It is deliberately the SAME +# number that make_figs.py checks every string inside every PNG against, so the +# deck has one floor rather than two, and nothing is legible on the slide but +# not inside the picture next to it. +# +# It is enforced here, in run(), rather than at the ~400 call sites, because a +# floor applied by hand is a floor that one new caption silently drops through. +MIN_PT = 9.0 + + def run(p, text, size=11, color=TEXT2, font=UI, bold=False, italic=False, spc=None): r = p.add_run(); r.text = text f = r.font - f.name, f.size, f.bold, f.italic = font, Pt(size), bold, italic + f.name, f.size, f.bold, f.italic = font, Pt(max(size, MIN_PT)), bold, italic f.color.rgb = color if spc is not None: r.font._rPr.set("spc", str(int(spc * 100))) @@ -166,7 +181,7 @@ 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 = 5 # annex GROUPS, not slides +N_ANNEX = 6 # annex GROUPS, not slides def annex_chrome(s, grp, eyebrow, title, part=None, nparts=None, title_size=27): @@ -230,15 +245,19 @@ def bullets(s, x, y, w, items, size=11, gap=7): return tf -def code(s, x, y, w, lines, size=8.5, title=None): - lh = 0.148 - h = 0.24 + len(lines) * lh + (0.22 if title else 0) +def code(s, x, y, w, lines, size=9, title=None): + # Line height has to follow the font size: it was a constant 0.148 in tuned + # for 8.5 pt, so raising the type to the projection floor pushed the last + # line out of the panel. 0.0174 in/pt reproduces the old value at 8.5. + size = max(size, MIN_PT) + lh = 0.0174 * size + h = 0.24 + len(lines) * lh + (0.24 if title else 0) rect(s, x, y, w, h, CODEBG, MSO_SHAPE.ROUNDED_RECTANGLE) ty = y + 0.12 if title: tf = tb(s, x + 0.18, ty, w - 0.36, 0.2) - run(para(tf, True), title, 7.5, MUTED, bold=True, spc=1.2) - ty += 0.22 + run(para(tf, True), title, 9, MUTED, bold=True, spc=1.2) + ty += 0.24 tf = tb(s, x + 0.18, ty, w - 0.36, h - 0.24) for i, ln in enumerate(lines): p = para(tf, i == 0, line=1.12) @@ -321,6 +340,17 @@ def card_figure(s, name, x, y, w, pad=0.10): return h + 2 * pad +def frame_rect(s, x, y, w, h, color=RED, wpt=1.75): + """An outline, not a fill: used to point at one row without recolouring it.""" + sh = s.shapes.add_shape(MSO_SHAPE.ROUNDED_RECTANGLE, In(x), In(y), In(w), In(h)) + sh.fill.background() + sh.line.color.rgb = color + sh.line.width = Pt(wpt) + sh.shadow.inherit = False + sh.adjustments[0] = 0.06 + return sh + + def notes(s, text): """Speaker notes. Detail that belongs in the talk, not on the slide. @@ -412,8 +442,8 @@ def section(kicker, title, thesis, items, rng, col=ACCENT, carry=None, Deliberately sparse — it exists to buy 10–15 s of stage setting, so it has to be readable at a glance and finished before the audience starts reading ahead. It carries no slide number and takes no tick of its own on the - progress track: slides 3–34 keep the numbers they have, so the annex's - cross-references ("expands slide 26") stay true. What it lights up instead + progress track: slides 3–33 keep the numbers they have, so the annex's + cross-references ("expands slide 25") stay true. What it lights up instead is the *range* the section covers, which is the thing the audience wants. """ s = new_slide() @@ -529,8 +559,9 @@ section("Context · what the code does", [("3–4", "The algorithm"), ("5–6", "The two machines"), ("7–8", "The CUDA kernel"), - ("9–10", "The hardware limit")], - rng=(3, 10)) + ("9–10", "What limits it"), + ("11", "The roadmap")], + rng=(3, 11)) # =========================================================== 3 · THE PHYSICS s = new_slide() @@ -550,14 +581,32 @@ caption(s, RAIL_X, 2.32 + h + 0.10, RAIL_W, "Each peak is one beam energy; its width is the resolution being measured.", size=8.5) rail(s, [ - ("row", "Pixels per frame", "400 × 400 = 160 000", TEXT2), + ("label", "MÖNCH03 · the detector this feeds"), + ("row", "Array · pitch · active area", "400 × 400 · 25 µm · 10 × 10 mm²", TEXT2), + ("row", "Frames per second", "1.3 k standard, 3–6 k optimised", AMBER), ("row", "Peak pixels = photons / frame", "~2 330 · 1.5 %", ACCENT), - ("row", "Pixels above 5σ (2.4 per photon)", "~5 700 · 3.6 %", ACCENT), -], y0=5.10, divider=False) +], y0=4.92, divider=False) caption(s, M, 6.68, 7.9, "Real MOENCH data, Cu fluorescence, MAX IV beamtime. One cluster is emitted " "per local maximum, so 2 330 counts photons, not lit pixels; the recorded " - "3×3 windows cover 12.7 % of the frame.") + "3×3 windows cover 12.7 % of the frame. The detector's frame rate is the " + "number every throughput figure in this deck should be read against.") +notes(s, """The detector, and why its frame rate is the yardstick. + +MOENCH03 is a hybrid silicon pixel detector: charge integration with analog +readout, 25 x 25 um^2 pitch, 400 x 400 pixels over a 10 x 10 mm^2 active area. +Its standard frame rate is 1.3 kHz; with optimised readout boards the design +reaches 3-6 kHz depending on configuration. + +That is the number to hold on to. 400 x 400 = 160 000 pixels is 312.5 kB per +frame at 16 bit, so 1.3 kHz is ~0.4 GB/s off the detector and 6 kHz is ~1.9 +GB/s. The 24-thread CPU finder does 6 762 FPS at 3x3, so it already keeps up +with the standard mode and roughly matches the optimised ceiling -- with nothing +left over for anything else on the machine, and nothing left at 9x9, where it +manages 1 503 FPS. Slide 33 closes this loop. + +Pixels above 5 sigma are ~5 700, about 3.6 % of the frame and 2.4 per photon; +that is the number that sets how much of the frame the 3x3 windows cover.""") # =========================================================== 4 · PER FRAME s = new_slide() @@ -567,8 +616,8 @@ 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 pixel that sees no photon** — about 80 % of " - "them — 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", @@ -607,17 +656,21 @@ rail(s, [ s = new_slide() chrome(s, 5, "The machine we are starting from", "CPU: latency-oriented, built to finish one thread fast") -figure(s, "img_cpu_core", M, 1.98, 7.3) -callout(s, M, 5.58, COL, +figure(s, "img_cpu_core", M, 1.86, 7.3) +# The die photo mirrors slide 6's: same grammar, compute units boxed, so the two +# machines are compared as objects and not only as block diagrams. It is kept +# WHOLE rather than cropped to the ten cores, because the L3 slab on the right +# and the I/O block on the left are half the die area -- which is the callout's +# point standing next to it in silicon. +callout(s, M, 5.42, 4.55, "Count the boxes: **6 fetch/decode**, out-of-order instruction selection, two " "levels of private cache, all of it to keep **two** instruction streams fed. " - "The ALUs are the small part.", h=0.86, size=10) -caption(s, M, 6.54, COL, - "One core, schematically. Intel Skylake is shown. The Zen 4 core in this machine " - "differs in detail (4 FP pipes rather than 3; AVX-512 double-pumped on 256-bit " - "datapaths) but not in kind: ~6-wide front end, 4 scalar ALUs, 2 SMT contexts, " - "private L1 + L2. Load/store units not drawn. After Stanford CS149, Fall 2025.", - size=8) + "The ALUs are the small part.", h=1.06, size=10) +figure(s, "img_cpu_die", 5.60, 5.42, 3.00) +caption(s, 5.60, 6.76, 3.00, + "Comet Lake · 10 cores boxed; nearly half the die is cache and I/O.", size=8) +caption(s, M, 6.60, 4.55, + "Skylake core and Comet Lake die after Stanford CS149, Fall 2025.", size=8) rail(s, [ ("label", "pc-moench-04 · AMD Ryzen 9 7950X"), ("gap", 0.10), @@ -634,6 +687,17 @@ rail(s, [ notes(s, """The CPU slide. The point is not that CPUs are bad -- it is what the silicon is SPENT on. +On the diagram: Intel Skylake is shown, one core, schematically. The Zen 4 core +in this machine differs in detail -- 4 FP pipes rather than 3, AVX-512 +double-pumped on 256-bit datapaths -- but not in kind: ~6-wide front end, 4 +scalar ALUs, 2 SMT contexts, private L1 + L2. Load/store units are not drawn. + +On the die photo: Comet Lake, a 10-core Core i9, at the same scale of argument +as the AD102 die on the next slide. Ten cores, and they do not fill the die -- +the orange slab on the right is L3 and the block on the left is I/O and the +memory controller. If someone asks why a 16-core Ryzen is not simply 16x a +1-core Ryzen, that picture is the answer. + Six fetch/decode units, an out-of-order instruction selector, branch prediction (not even drawn), L1 + L2 private cache: all of that machinery exists to find independent work INSIDE one instruction stream, and to hide memory latency @@ -665,10 +729,9 @@ callout(s, M, 6.02, COL, "contexts** and a wall of lanes. Nothing reorders instructions: when a warp " "stalls on memory, the selector just **runs a different one**.", h=0.86, size=10) -caption(s, M, 6.90, COL, - "One SM: a V100 is shown; this card's is the same idea (128 FP32 lanes, 48 warp " - "slots, 100 kB shared memory). Note how few of the units are FP64, and remember " - "it at opt7. Diagrams after Stanford CS149, Fall 2025.", size=8) +caption(s, M, 6.82, COL, + "One SM: a V100 is shown; this card's is the same idea — 128 FP32 lanes, " + "48 warp slots, 100 kB shared memory. After Stanford CS149, Fall 2025.", size=8) rail(s, [ ("label", "NVIDIA GeForce RTX 4090"), ("gap", 0.12), @@ -686,13 +749,19 @@ notes(s, """The GPU slide, and the number to land. 128 SMs x 1 536 threads = 196 608 thread slots that can be RESIDENT at the same time. Our frame is 160 000 pixels. The entire frame fits in the machine at once, one thread per pixel, with room left over -- which is why slide 7's "one thread -per pixel" is not a figure of speech, and why slide 9 can talk about occupancy +per pixel" is not a figure of speech, and why slide 10 can talk about occupancy as a real quantity (3x3 achieves 100 %: 1 536 threads resident per SM). Latency hiding, in one sentence: the CPU hides memory latency with a cache and out-of-order execution; the GPU hides it by having 48 other warps ready to run. That is why there is no reorder buffer on this diagram and no branch predictor. +Count the pink cells in the legend before moving on: FP64 units are a small +minority of the SM, 8 MUL/ADD per clock against 16 for fp32 and 16 for int. On +a GeForce part the ratio is far worse than this V100 diagram suggests -- 1/64. +That is half of why opt7 pays, and it is worth planting here so opt7 lands as a +consequence rather than a surprise. + The trap to pre-empt: 16 384 lanes vs 32 streams is a factor of 512, and we measure x9.1. Say so before someone else does. We are not lane-limited; we are BANDWIDTH-limited -- 1 TB/s of device memory, and 25 GB/s of PCIe to reach it. @@ -721,7 +790,7 @@ code(s, M, 4.35, 7.5, [ " d_frame, d_pd_mean, d_pd_sum, d_pd_sum2, d_pd_off, n_pd_samples,", " nSigma, nrows, ncols, d_clusters, d_cluster_count, max_clusters);", ], size=8, title="LAUNCH CONFIGURATION · ClusterFinderCUDA.hpp") -caption(s, M, 5.80, 7.5, +caption(s, M, 5.94, 7.5, "400×400 pixels → a 25×25 grid of 16×16 blocks = 625 blocks per frame, " "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.") @@ -737,9 +806,9 @@ code(s, 8.5, 2.85, 4.1, [ " pedestal.«push_fast»(iy, ix, ...);", "}", ], size=8, title="THREE OUTCOMES, NOT TWO") -callout(s, 8.5, 4.80, 4.1, +callout(s, 8.5, 5.04, 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 " + "shadow is **neither recorded nor fed back**, and the CUDA kernel " "reproduces it exactly.", h=1.15, size=10) callout(s, 8.5, 6.10, 4.1, @@ -768,11 +837,11 @@ code(s, 8.4, 3.05, 4.3, [ "// pedestal subtraction fused into the load", "sh[tid] = d_frame[gid] - «d_pd_mean»[gid];", ], size=8, title="clusterfinder_kernel.cuh") -callout(s, 8.4, 4.60, 4.3, +callout(s, 8.4, 4.80, 4.3, "Only **odd** cluster sizes are supported (3×3, 5×5, 7×7, 9×9), so that " "the centre pixel is unique and local-maximum suppression is well defined.", h=1.15, size=10) -caption(s, 8.4, 5.95, 4.3, +caption(s, 8.4, 6.06, 4.3, "The tile is stored in COMPUTE_TYPE (float), not in the pedestal type: " "1.3 KB for 3×3, 2.3 KB for 9×9, against 100 KB of shared memory per SM. " "Even the old double-precision tile only reached 4.5 KB.") @@ -781,80 +850,100 @@ caption(s, M, 6.70, 7.5, "13 % at 32×32, which is the first half of the block-size argument. " "The second half is registers, next slide.") -# =========================================================== 7 · OCCUPANCY +# ================================================= 7 · REGISTERS (the input) +# Registers come BEFORE occupancy, not after. Occupancy is an output of the +# register budget, so presenting the percentage first and the cause second asked +# the audience to learn a metric and then be told it was not the point. s = new_slide() -chrome(s, 9, "Hardware · occupancy", "3×3 reaches 100 % occupancy, 9×9 only 33 %") -statstrip(s, M, 1.82, 11.9, [ - ("block size", "16 × 16"), - ("threads / block", "256"), - ("max threads / SM · 48 warps", "1 536"), - ("registers / SM", "65 536"), - ("shared mem / block", "1.3 – 2.3 KB"), -]) -figure(s, "fig_occupancy", 0.9, 2.78, 11.5) -callout(s, M, 6.12, 11.9, - "**16×16 is the balance point**, enough threads to amortise the halo, few " - "enough that 6 blocks still fit per SM. At 3×3 that is 48 of 48 warps " - "resident: **fully occupied**. Shared memory never binds; **registers do**.", - h=0.66, size=10.5) -caption(s, M, 6.84, 11.9, - "Shipped f32 build, measured not estimated: cuobjdump -res-usage for registers " - "and spills, cudaOccupancyMaxActiveBlocksPerMultiprocessor for blocks/SM. " - "RTX 4090 (sm_89). Reproduce: python/tests/perf/kernel_resources.py") - -# =========================================================== 8 · REGISTERS -s = new_slide() -chrome(s, 10, "Hardware · the real limiter", "Occupancy is a means, not the goal") -bullets(s, M, 1.92, COL, [ - "The limiter is **register pressure**: every thread keeps a private " - "clusterData[CSX × CSY] staging array, so demand grows with the **square** of " - "the cluster size. 3×3 costs **38 registers**, 9×9 costs **128**.", - "**Zero spills either way.** ptxas would rather give up occupancy than spill to " - "local memory, and on this kernel that is the right trade.", - "Low occupancy is not automatically a problem: at 9×9 one kernel nearly fills " - "the machine on its own.", -], size=10) -figure(s, "fig_regpressure", M, 3.24, 7.9) -callout(s, M, 5.00, COL, - "**Occupancy is an output.** At 9×9 the register file is full at two blocks, " - "so two thirds of the thread slots are stranded: nothing else can be resident " - "no matter how much work is queued.", h=0.80, size=10) -code(s, M, 5.92, COL, [ +chrome(s, 9, "Hardware · what runs out first", + "38 registers per thread at 3×3, 128 at 9×9") +bullets(s, M, 1.90, COL, [ + "An SM has a fixed budget: **65 536 registers** and **1 536 thread slots**. " + "A 16×16 block claims 256 slots and 256 × (registers per thread), and whichever " + "budget runs out first decides how many blocks fit on the SM at once.", + "Every thread keeps a private **clusterData[CSX × CSY]** staging array, so " + "register demand grows with the **square** of the cluster size. Neither build " + "spills: ptxas would rather hold fewer blocks than go to local memory.", +], size=10.5) +figure(s, "fig_regpressure", M, 3.34, 7.9) +callout(s, M, 5.30, COL, + "**At 9×9 the register file is exactly full at two blocks** (2 × 128 × 256 " + "= 65 536), so two thirds of the thread slots are stranded. At 3×3 the " + "slots run out first and the registers still have room.", h=0.80, size=10.5) +code(s, M, 6.22, COL, [ "«cuobjdump -res-usage» build/aare/_aare_cuda*.so | c++filt", " 3x3: REG:«38» STACK:0 LOCAL:0 # STACK/LOCAL 0 = no spills", " 9x9: REG:«128» STACK:0 LOCAL:0", - "python python/tests/perf/«kernel_resources.py» # blocks/SM, runtime check", -], size=7.5, title="HOW THESE NUMBERS WERE OBTAINED · READ FROM THE BUILT .SO, NO REBUILD") +], size=8.5, title="MEASURED, NOT ESTIMATED · READ FROM THE BUILT .SO") rail(s, [ - ("label", "Register budget · sm_89 · f32"), + ("label", "Per SM · sm_89 · f32 build"), ("gap", 0.15), - ("stat", "3×3 · occupancy", "100 %", ACCENT), - ("stat", "9×9 · occupancy", "33 %", AMBER), + ("stat", "3×3 · blocks resident", "6", ACCENT), + ("stat", "9×9 · blocks resident", "2", AMBER), ("gap", 0.05), ("row", "Spills, either case", "0 bytes", TEXT2), - ("row", "3×3 on the f64 build", "47 regs → 83 %", AMBER), + ("row", "3×3 on the f64 build", "47 regs → 5 blocks", AMBER), ("gap", 0.20), - ("note", "Which is why the 9×9 kernel is the one worth optimising, see opt7."), + ("note", "How many blocks fit is the whole hardware story. The next slide is " + "what that buys."), ]) -notes(s, """The arithmetic behind the bars, and the one caveat. +notes(s, """The arithmetic, and the one build-dependent caveat. One SM holds 65 536 registers and 1 536 thread slots. A 16x16 block is 256 threads, so a block costs regs_per_thread x 256 registers. At 3x3 that is 9 728, -six blocks fit, and the THREAD SLOTS run out first at 100 %. At 9x9 it is -32 768, so two blocks exactly fill the REGISTER FILE and strand two thirds of -the slots. Same kernel, same block size, opposite binding resource. - -The caveat on "low occupancy is fine": it is build-dependent. On the shipped f32 -build the 9x9 kernel is long enough that four streams overlap by only 1.02x, so -the extra streams buy nothing in kernel concurrency. On the f64 build the kernel -is 3.5x longer and DOES overlap, 1.32x, and there four streams genuinely lower -the floor. Quote the overlap factor, never "occupancy is fine" on its own. +six blocks fit, and the THREAD SLOTS run out first. At 9x9 it is 32 768, so two +blocks exactly fill the REGISTER FILE and strand two thirds of the slots. Same +kernel, same block size, opposite binding resource. Build dependence of the register count itself: the f64 pedestal costs 3x3 nine -extra registers (47), which loses a block per SM and drops it to 83 %. 9x9 is -unmoved at 128 -- there the limiter is the clusterData[9][9] staging array, not -the pedestal, which is why opt7 helps 3x3's occupancy and not 9x9's.""") +extra registers (47), which loses a block per SM. 9x9 is unmoved at 128 -- there +the limiter is the clusterData[9][9] staging array, not the pedestal, which is +why opt7 helps 3x3's occupancy and not 9x9's. +Reproduce blocks/SM at runtime with python/tests/perf/kernel_resources.py, which +calls cudaOccupancyMaxActiveBlocksPerMultiprocessor on the built kernel.""") + +# ================================================ 8 · OCCUPANCY (the output) +s = new_slide() +chrome(s, 10, "Hardware · occupancy", + "Occupancy is a latency-hiding budget") +bullets(s, M, 1.86, 11.9, [ + "When a warp stalls on memory the SM switches to another warp that is already " + "resident. **Occupancy = resident warps / the maximum the SM can hold**: how " + "many alternatives it has to switch to. More resident warps, more stalls hidden.", +], size=11) +statstrip(s, M, 2.44, 11.9, [ + ("block size", "16 × 16"), + ("threads / block", "256"), + ("thread slots / SM", "1 536"), + ("max warps / SM", "48"), + ("occupancy 3×3 · 9×9", "100 % · 33 %"), +], h=0.74) +figure(s, "fig_occupancy", 1.37, 3.34, 10.6) +callout(s, M, 6.44, 11.9, + "**16×16 is the balance point**: enough threads to amortise the halo, few " + "enough that 6 blocks still fit. But 33 % is not a failure to fix — it is " + "what the register budget allows, and at 9×9 one kernel nearly fills the " + "machine on its own.", h=0.62, size=10.5) +notes(s, """Why "33 % occupancy" is not the alarm it looks like. + +Occupancy only buys latency hiding, and latency hiding only matters if there is +latency left unhidden. At 9x9 the kernel is long and each thread does a lot of +work per byte loaded, so 16 resident warps are enough to keep the SM fed. + +It is also build-dependent, and that is the honest caveat. On the shipped f32 +build the 9x9 kernel overlaps across four streams by only 1.02x -- extra streams +buy nothing in kernel concurrency. On the f64 build the kernel is 3.5x longer +and DOES overlap, 1.32x, and there four streams genuinely lower the floor. So +quote the overlap factor, never "occupancy is fine" on its own. + +Shared memory never binds here: 1.3 KB per block at 3x3, 2.3 KB at 9x9, against +100 kB per SM. + +Provenance for both hardware slides: shipped f32 build, measured not estimated. +cuobjdump -res-usage for registers and spills; +cudaOccupancyMaxActiveBlocksPerMultiprocessor for blocks per SM; RTX 4090 +(sm_89). Reproduce with python/tests/perf/kernel_resources.py.""".rstrip()) # =========================================================== 10 · THE LADDER s = new_slide() chrome(s, 11, "Roadmap", @@ -868,7 +957,7 @@ rows = [ ("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", "the kernel is the tallest bar — at 9×9", "at 9×9 [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), ] @@ -894,10 +983,10 @@ for i, (tag, name, sub, gain, col) in enumerate(rows): run(para(tf, True, align=PP_ALIGN.RIGHT), gain, 12.5, col, bold=True) y += 0.52 caption(s, M, 6.62, 11.9, - "Speedups are 3×3 vs the best CPU configuration, 24 threads. Two routes were measured and " - "rejected, two result transports inside Act II, shown next to the step they " - "lost to, and CUDA Graphs, which is in the annex. The rule that predicts " - "the wins predicts the failures too.") + "Speedups are 3×3 vs the best CPU configuration, 24 threads. Three routes were " + "measured and rejected: CUDA Graphs, and two faster ways of copying the results. " + "All three are in annex A2, because the rule that predicts the wins predicts " + "the failures too.") # --------------------------------------------------------- divider · ACT I section("Act I of III · feed the GPU", @@ -931,16 +1020,14 @@ code(s, M, 3.18, COL, [ " <<>>(d_frame, d_pd_mean, ...);", "cudaMemcpy(h_out, d_out, out_bytes, cudaMemcpyDeviceToHost);", ], title="ClusterFinderCUDAOpt2.hpp · find_clusters()") -callout(s, M, 4.50, COL, +callout(s, M, 4.62, 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 **[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) + "overlap, so the **floor** — the fastest a frame can go if the host cost " + "nothing — is **max(H2D, kernel, D2H)**, never the sum. At 3×3 that is " + "max(**16.17**, 15.17, 7.69) = **16.2 µs → 61 859 FPS**. Exactly how each of " + "those three is measured is slide 18; it does not change this one.", + h=1.00, size=10.5) +figure(s, "fig_opt1_timeline", M, 5.80, COL) rail(s, [ ("label", "opt1 · 3×3 · 100 k frames · f64"), ("gap", 0.10), @@ -1042,28 +1129,27 @@ code(s, 8.5, 3.15, 4.1, [ "", "«cudaHostUnregister»(ptr);", ], size=8, title="ClusterFinderCUDA.hpp") -callout(s, 8.5, 4.72, 4.1, +callout(s, 8.5, 4.90, 4.1, "**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, +caption(s, 8.5, 5.94, 4.1, "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 " + "**The rule, first sighting: ×1.32 at 3×3 but only ×1.03 at 9×9.** Pinning " "attacks H2D, the tallest bar at 3×3, the shortest at 9×9.", h=0.72, size=10, color=AMBER) # -------------------------------------------------------- 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. Still 3×3 — " + "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")], - rng=(16, 18), col=PALE, + (17, "opt6 · zero-copy")], + rng=(16, 17), col=PALE, carry=("Arriving at", "38 486 FPS", "opt4 · 26.0 µs per frame · 62 % of the GPU floor")) @@ -1152,162 +1238,221 @@ rail(s, [ "f64 the run BEATS the 32.66 µs max, so the 30.0 sustained sets it [A1]."), ]) -# =========================================================== 20 · ACT II rejected -s = new_slide() -chrome(s, 18, "Act II · two rejected routes", - "The copy is allocation-bound, not bandwidth-bound") -rows = [ - ("B′", "One allocation per chunk", "collect_packed()", - "Removes the per-frame malloc but keeps the copy, and the copy is ~80% of the " - "cost. Worse, the replacement allocation is 1.17 GB, far above glibc's mmap " - "threshold, so it is mmap'd and munmap'd every chunk: 606 566 faults, ~21 µs/frame.", - "69.3 µs · deleted from the API"), - ("B″", "Parallel materialisation", "8-thread copy pool", - "Each worker gets its own glibc arena, which destroys the cross-run heap reuse " - "that makes the single-threaded path cheap. Faults went 9 700 → 2 270 000; " - "MALLOC_ARENA_MAX=1 collapsed them back to 138 k, which is the proof.", - "+6% at best, −33% when results are freed promptly"), -] -y = 2.00 -for tag, name, how, body, verdict in rows: - rect(s, M, y, 11.9, 2.05, PANEL) - rect(s, M, y, 0.035, 2.05, AMBER) - tf = tb(s, M + 0.30, y + 0.22, 1.0, 0.4) - run(para(tf, True), tag, 20, AMBER, bold=True, font=MONO) - tf = tb(s, M + 1.30, y + 0.20, 6.0, 0.3) - run(para(tf, True), name, 14, PALE, bold=True) - tf = tb(s, M + 1.30, y + 0.52, 6.0, 0.3) - run(para(tf, True), how, 9.5, MUTED, font=MONO) - tf = tb(s, M + 1.30, y + 0.90, 10.2, 1.0) - run(para(tf, True, line=1.3), body, 10, TEXT2) - tf = tb(s, M + 1.30, y + 1.68, 10.2, 0.3) - run(para(tf, True), verdict, 10.5, AMBER, bold=True) - y += 2.25 -callout(s, M, 6.38, 11.9, - "**Copying faster does not help when the cost is the OS populating pages.** " - "The only winning move is not to allocate — which is exactly what opt6 does. " - "materialize_slot() is deliberately single-threaded and carries a comment " - "saying so, to stop the experiment being repeated.", h=0.80, size=10.5) - # ------------------------------------------------------- divider · ACT III section("Act III of III · the kernel", "Only now is the kernel the tallest bar", - "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"), - (22, "why it comes last")], - rng=(19, 22), col=AMBER, + "The story moves to 9×9, where the kernel is finally the tallest bar. First, " + "how the engine times in this act are measured.", + [(18, "how the engine times are measured"), + (19, "opt7 · FP32 pedestal"), + (20, "catastrophic cancellation"), + (21, "why it comes last")], + rng=(18, 21), col=AMBER, carry=("Arriving at", "58 495 FPS", "opt6 · 3×3, and 33 323 FPS at 9×9, where this act pays")) +# ============================================ 10b · THE MEASUREMENT CONVENTION +# Promoted out of the annex. Everything from here on is quoted as "[build · s1]" +# or "[build · s4]" and compared against a "floor", and none of those three words +# had been defined anywhere the audience would see them. The annex keeps the full +# grid of numbers; this slide keeps only the three definitions and the one +# picture that makes the middle one make sense. +s = new_slide() +chrome(s, 18, "How the engine times are measured", + "Two configurations, one floor") +cards = [ + ("s1", ACCENT, "One stream, nothing else running", + "How long an operation actually takes. The right number for a capability " + "claim, and for the headroom that is left."), + ("s4", PALE, "The shipped pipeline, four streams", + "How BUSY each engine is per frame: the union of its intervals, which is " + "not the sum of the durations. The only number that can set a floor."), + ("floor", AMBER, "Set by the busiest engine", + "The fastest a frame could go if the host cost nothing. In µs per frame, or " + "its reciprocal in FPS. The LOWER of the profiled estimate and the best " + "rate sustained."), +] +x = M +for tag, col, title, body in cards: + rect(s, x, 1.82, 3.83, 1.70, PANEL) + rect(s, x, 1.82, 3.83, 0.035, col) + tf = tb(s, x + 0.26, 2.00, 3.3, 0.30) + run(para(tf, True), tag, 14, col, bold=True, font=MONO) + tf = tb(s, x + 0.26, 2.34, 3.35, 0.30) + run(para(tf, True, line=1.1), title, 11, PALE, bold=True) + tf = tb(s, x + 0.26, 2.68, 3.35, 0.80) + run(para(tf, True, line=1.22), body, 9.5, TEXT2) + x += 4.03 +figure(s, "fig_measure", 1.52, 3.62, 10.3) +callout(s, M, 6.48, 11.9, + "**Nothing that measures a duration falls under load** — yet the 9×9 kernel " + "row falls 39.9 → 32.7 µs from s1 to s4. That is the tell: s4 is occupancy, " + "not duration. **The floor is quoted both ways in this deck** — 30.01 µs per " + "frame is 33 323 FPS. Full engine grid: **annex A1**.", h=0.62, size=10.5) +notes(s, """Say the floor out loud, in this order. + +1. There are three engines and they are independent: the H2D copy engine, the + SMs, and the D2H copy engine. PCIe is full duplex, so a frame's cost is + never the sum of the three -- it is the slowest of them. + +2. There is exactly ONE copy engine per direction. So two H2D copies can never + run at the same time; they queue. That is not a modelling assumption, it is + measured: H2D_overlap and D2H_overlap are 1.000 in every row of probes.csv. + Kernels are the only row that ever overlaps. + +3. That is why s4 is a UNION, not a sum. Four kernels, each 43.2 us long, but + the SMs are busy only 32.66 us per frame: overlap factor 1.32x. + +4. On units, because the card says "1 / the busiest engine" and the axis says + microseconds: they are the same number. The busiest engine's busy time per + frame is in us/frame; one divided by it is frames per second. 30.01 us/frame + IS 33 323 FPS. The deck quotes whichever reads better in context -- us/frame + when comparing engines, FPS when comparing against the CPU or the detector -- + and "% of floor" is the same ratio either way. + +5. The floor is max(H2D, kernel, D2H) at s4. At 9x9 f64 that max is 32.66 us + = 30 618 FPS. But the unprofiled pipeline actually sustained 30.01 us/frame, + which is FASTER than the probe says is possible -- because under nsys + submission is sparser, kernels overlap less, and a less-overlapped interval + set has a LARGER union. A sustained rate is an existence proof; a probe is an + estimate. So the floor is the lower of the two, and the 1.32x is a lower + bound on the real overlap. + +If time is short, say only: "s1 is how long it takes, s4 is how busy it is, and +the floor is the busiest engine." The rest is in A1.""") + + # =========================================================== 18 · OPT7 why 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, [ - "**~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.", - "Third, quietly: the narrower accumulators free **9 registers at 3×3**, 47 → 38, " - "which buys back a block per SM and takes occupancy **83 % → 100 %**.", + "**~80 % of pixels** take the **pedestal-update** branch: it reads off, sum and " + "sum², and writes back sum, sum² and mean. All four pedestal arrays are " + "DEVICE_PED_TYPE, so one typedef halves all six accesses: **48 bytes per " + "updating pixel in FP64, 24 in FP32**.", + "The kernel is **bandwidth-bound**, so halving that traffic nearly halves the " + "time. On a GeForce part there is a second effect: **FP64 arithmetic runs at " + "1/64 of FP32**, and the pedestal update was paying that tax on every pixel.", + "Quietly, a third: the narrower accumulators free **9 registers at 3×3**, 47 → " + "38, which buys back a block per SM.", ], size=10.5) -figure(s, "fig_f32_kernel", M, 4.05, 7.5) +figure(s, "fig_f32_kernel", M, 3.90, 7.5) code(s, 8.5, 1.95, 4.1, [ "// clusterfinder_kernel.cuh", "using COMPUTE_TYPE = float;", "using DEVICE_PED_TYPE = «float»;", "// was: double", -], size=8.5, title="ONE TYPEDEF") -callout(s, 8.5, 3.05, 4.1, - "Kernel, 9×9 **[s1 · cap 1700 · exclusive]**\n**39.86 µs → 23.70 µs (−40.5%)**\n" +], size=9, title="ONE TYPEDEF") +callout(s, 8.5, 3.24, 4.1, + "Kernel, 9×9 **[s1 · cap 1700]**\n**39.86 → 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=0.80, size=10, color=AMBER) -callout(s, 8.5, 6.20, 4.1, +callout(s, 8.5, 4.68, 4.1, + "Naive FP32 is **wrong** (next slide), and even when correct it only pays " + "**because Act II came first**.", h=0.86, size=10.5, color=AMBER) +callout(s, 8.5, 5.76, 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) + h=1.00, size=10.5) +caption(s, 8.5, 6.94, 4.1, + "Both builds, same git rev, 20 000 frames. Full engine grid: A1.", size=9) +notes(s, """Reading the two panels, and the number to quote. -# =========================================================== 19 · OPT6 trap -s = new_slide() -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 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, [ - ("label", "What it looked like"), - ("gap", 0.15), - ("stat", "Extra clusters", "+28.06%", AMBER), - ("gap", 0.05), - ("note", "Quiet pixels got rms → 0, so their threshold became 0 and they fired " - "on every single frame, producing a large unphysical high-energy tail " - "in the spectrum."), - ("gap", 0.35), - ("row", "Affected pixels", "~1–2% of the sensor", TEXT2), - ("row", "Written up in", "docs/pedestal_precision_…", MUTED), -]) +Left is s1: one stream, nothing else running, so those are true durations. The +kernel binds in BOTH arms there -- 39.86 and 23.70 against a D2H of ~21.95 -- so +s1 alone would say the kernel is the thing to optimise and stop there. -# =========================================================== 20 · OPT6 fix +Right is s4, the shipped four-stream pipeline, and it says something different: +the f64 kernel's busy time falls to 32.66 (self-overlap, 1.32x) while every +transfer RISES under contention, and once opt7 puts the kernel at 23.94 the D2H +bar at 25.24 is above it. That is the handover slide 21 is about. + +So: -40.5 % is the s1 kernel claim and is the honest headline for what the +typedef does to the arithmetic. -26.7 % is the same change measured at s4, where +self-overlap had already hidden part of the win. Both are true, they measure +different things, and the previous slide is the rule for which to quote. + +The 3x3 occupancy effect is real but not worth stage time: 47 registers -> 38 +takes 3x3 from five blocks per SM to six, 83 % -> 100 % occupancy. It changes +nothing end to end, because at 3x3 H2D is the floor.""") + + +# ================================================ 19 · THE TRAP AND THE FIX +# Was two slides. The split spent one whole slide on the ULP arithmetic of the +# error floor, which is the least transferable part of the story: what the +# audience needs is that a tiny answer computed as the difference of two huge +# numbers is not computable in f32, what that did to the physics, and the +# two-line change that removes it. The error-floor curve and the full rewrite +# are annex A5. s = new_slide() -chrome(s, 21, "Act III · opt7 · the variance rewrite", +chrome(s, 20, "Act III · opt7 · catastrophic cancellation", "Accumulate what is small, not what is large") -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, and report the mean " - "as **X₀ + sum/n** — so nothing downstream changes.", -]) -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;", - "", - "// after: centred on a frozen per-pixel offset X0", - "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, 5.94, COL, - "Result: the 100% FP32 build now matches the FP64 build to " - "**3 × 10⁻⁷**, 70 clusters out of 233 million.") +bullets(s, M, 1.90, COL, [ + (TEXT2, "**The trap.** The variance was computed as **var = E[X²] − mean²**. " + "With a pedestal at ~4 655 ADU both operands are ≈ 2.17 × 10⁷ while the answer " + "is ≈ 2 000. FP32 carries ~7 digits, so the answer inherits an **absolute** " + "error of **±3 ADU², which does not shrink as the answer does**."), + (TEXT2, "**What it cost.** For a quiet pixel whose true variance is 9, ±3 is a " + "third of it; below rms ≈ 2 the variance goes negative, the rms **clamps to " + "zero**, and its 5σ gate becomes a 0σ gate. **~1–2 % of the sensor** then fires " + "every frame: **+28.06 % clusters** and a population of clusters below the " + "physical threshold."), + (TEXT2, "**The fix.** Freeze **X₀ = round(mean)** once at the end of pedestal " + "training and accumulate the **centred** value **Y = X − X₀**. Both operands " + "become O(rms) and the cancellation is gone."), +], size=10.5) +figure(s, "fig_cancellation", M, 3.92, COL) +caption(s, M, 6.42, COL, + "Left: the two operands and the answer, log scale, against the ±3 ADU² error. " + "Right: the f64 curve is measured (23.2 M clusters); the f32 curve is " + "reconstructed — measured area, modelled shape, method in the notes. " + "The two-line patch itself is in annex A5.", size=9) rail(s, [ - ("label", "Why it works"), - ("gap", 0.15), - ("note", "Precision is relative. Floats resolve small numbers finely and large " - "numbers coarsely, so never let a small answer be the difference of " - "two large numbers."), - ("gap", 0.5), - ("row", "f32 vs f64 counts", "3 × 10⁻⁷", ACCENT), - ("row", "vs CPU", "0.0039%", ACCENT), - ("gap", 0.3), + ("label", "naive f32 · what it did"), + ("gap", 0.12), + ("stat", "Extra clusters", "+28.06 %", AMBER), + ("row", "Pixels affected", "~1–2 % of the sensor", AMBER), + ("gap", 0.28), + ("label", "after the rewrite"), + ("gap", 0.12), + ("stat", "f32 vs f64 counts", "3 × 10⁻⁷", ACCENT), + ("row", "vs the CPU baseline", "0.0039 %", ACCENT), + ("gap", 0.20), ("note", "X₀ must never be updated: the accumulators are defined relative to it."), ]) +notes(s, """The one sentence to leave the room with. + +"Precision is relative." A float resolves small numbers finely and large numbers +coarsely, so a small answer must never be computed as the difference of two +large numbers. That is the transferable lesson; everything else on this slide is +this detector's instance of it. + +Why the tail matters more than the count. +28 % is a number you might argue +about. A spectrum with a population sitting below the 5 sigma cut is not +arguable: those clusters cannot physically be there, and any gain or resolution +fit done on that spectrum is wrong. + +Honesty about the right-hand panel: nobody kept the broken build around to +re-run, so the f32 curve is reconstructed. What is measured is the f64 curve and +the +28.06 % excess (CPU 116 010 113 vs CUDA 148 559 598, SS1 of the write-up). +The placement follows SS9: a correct 5 sigma gate cuts at ~225 ADU, a collapsed +gate admits the whole positive side of the pixel's distribution, so the excess +is smeared upward from zero. The write-up's TL;DR calls it a high-energy tail; +SS9 derives it from ~0 upward. Both describe the same corrupted pixels. + +Welford's online variance is the other correct answer and is mentioned in the +write-up. The frozen-offset form was chosen because it is two lines and does not +change the update's arithmetic cost.""") + # =========================================================== 21 · OPT6 when s = new_slide() -chrome(s, 22, "Act III · why this act comes last", +chrome(s, 21, "Act III · why this act comes last", "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 " + "**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**.", @@ -1315,8 +1460,8 @@ callout(s, M, 4.86, 5.85, 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]", + "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/33]", 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 " @@ -1325,31 +1470,32 @@ caption(s, M, 6.30, 11.9, "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 " + "applied at earlier steps, configurations that exist only to make this " "comparison controlled.") # ------------------------------------------------------- divider · results section("Results · what came out of it", "The whole ladder, and how to use it", "Both cluster sizes end to end, and the audit behind the numbers.", - [("23–24", "Results, both cluster sizes"), - ("25", "Where the time went"), - ("26–27", "What the numbers survived")], - rng=(23, 27), col=PALE, + [("22–23", "Results, both cluster sizes"), + ("24", "Where the time went"), + ("25–26", "What the numbers survived")], + rng=(22, 26), col=PALE, carry=("Arriving at", "58 495 FPS", "opt6 · everything after this is the ladder seen whole")) # =========================================================== 22 · RESULTS s = new_slide() -chrome(s, 23, "Results · 3×3", "×9.1 at 3×3, sitting on the H2D floor") +chrome(s, 22, "Results · 3×3", "×9.1 at 3×3, sitting on the H2D floor") figure(s, "fig_arc", 1.95, 1.88, 9.4) callout(s, M, 5.80, 5.85, "**×9.1 over 24 CPU threads**, 14.8 s → 1.63 s for 100 000 frames, " "and **at the H2D floor**.", h=0.8) callout(s, 6.75, 5.80, 5.85, - "Every step is **monotonic**, and correctness is held constant **throughout**, " - "0.004 % against the CPU baseline, and **exact** against the CPU twin " - "that isolates the port (slides 28–31).", h=0.8, color=AMBER) + "Every step is **monotonic**, and correctness is held constant **throughout**: " + "0.004 % against the CPU baseline; against the CPU twin that isolates the port, " + "**exact on the f64 pedestal** and **6 clusters in 23 M** on the shipped f32 " + "(slides 27–30).", h=0.8, color=AMBER) caption(s, M, 6.62, 11.9, "3×3 clusters · nσ = 5 · 100 000 frames · batch 2 000 · 4 streams · 5 reps · " "warm = best of reps 1–4 (collect() does not converge, it oscillates between " @@ -1358,7 +1504,7 @@ caption(s, M, 6.62, 11.9, # =========================================================== 23 · RESULTS 9x9 s = new_slide() -chrome(s, 24, "Results · 9×9", "×26.5 at 9×9, and opt7 hands the floor to D2H") +chrome(s, 23, "Results · 9×9", "×26.5 at 9×9, and opt7 hands the floor to D2H") figure(s, "fig_arc_9x9", 1.95, 1.88, 9.4) callout(s, M, 5.80, 5.85, "**×26.5 over 32 CPU threads**; the kernel is the tallest bar for the whole " @@ -1376,7 +1522,7 @@ caption(s, M, 6.62, 11.9, # =========================================================== 24 · WHERE TIME GOES s = new_slide() -chrome(s, 25, "Where the time actually went", +chrome(s, 24, "Where the time actually went", "The host bar dies first, then the floor itself drops") figure(s, "fig_overhead", 1.95, 1.95, 9.4) callout(s, M, 5.30, 5.85, @@ -1388,55 +1534,66 @@ callout(s, 6.75, 5.30, 5.85, "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 — the LOWER of the s4 " + "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 +# ======================================================= 24 · MEASUREMENT AUDIT +# Was a three-card slide that gave equal weight to page faults, CUDA-event +# timing and profiler overhead. Only the first of the three changes a number the +# audience is about to be shown, and only the first is something they will hit +# themselves. The other two are instrument caveats: named here, worked through +# in A6. The full three-card version is A6·1. s = new_slide() -chrome(s, 26, "Behind the numbers · what they had to survive", - "Three ways a GPU benchmark lies") -items = [ - ("First-touch page faults", AMBER, - "Each run materialises ~10 GB of clusters. The first pass faults in ~2.6 M " - "pages at 0.7 µs each, up to 4 s of pure OS work inside the timer.", - "Fix: re-run until getrusage() minor faults plateau (< 200 k)."), - ("CUDA-event kernel timing", AMBER, - "avg_kernel_time_ms() measures elapsed time on a stream, including waiting " - "for other streams. Under 8-stream load it over-reads by up to 3.5×.", - "Fix: Nsight Systems per-instance times; 1 stream for exclusive numbers."), - ("The profiler itself", AMBER, - "Under nsys, wall time per frame inflates ~4× from API tracing.", - "Fix: GPU op times from nsys, wall times from unprofiled runs."), -] -x = M -for title, col, body, fix in items: - rect(s, x, 2.0, 3.83, 3.15, PANEL) - rect(s, x, 2.0, 3.83, 0.035, col) - tf = tb(s, x + 0.26, 2.28, 3.3, 0.6) - run(para(tf, True, line=1.15), title, 13, PALE, bold=True) - tf = tb(s, x + 0.26, 3.02, 3.3, 1.5) - run(para(tf, True, line=1.3), body, 10, TEXT2) - tf = tb(s, x + 0.26, 4.42, 3.3, 0.65) - run(para(tf, True, line=1.25), fix, 9.5, ACCENT) - x += 4.03 -code(s, M, 5.4, 11.9, [ +chrome(s, 25, "Behind the numbers · the artefact that dominates", + "A GPU benchmark mostly measures the operating system") +bullets(s, M, 1.90, COL, [ + "Every run that keeps its results materialises **~10 GB of clusters**. The " + "first pass **faults in ~2.6 M pages**, and each fault costs the kernel a page " + "it must find and **zero** before the write can proceed.", + "At **0.7 µs a fault**, that is up to **4 seconds of pure OS work inside the " + "timer** — on a run whose GPU work is under 2 seconds. Nothing about the GPU " + "changed; kernel time is constant across every one of these runs. Nor is it a " + "warm-up you can wave away: it is what a **user's first run** looks like, " + "which is the next slide.", +], size=10.5) +code(s, M, 3.76, COL, [ "# every timed cell in the benchmark notebook is bracketed with:", "mf0 = resource.getrusage(resource.RUSAGE_SELF).ru_minflt", "... t = time.perf_counter() - t0 ...", "print(f'minor faults: {mf1-mf0:,}') # quote the run where this plateaus", -], title="THE FAULT PROTOCOL · python/tests/ClusterFinderCUDA_perf.ipynb") -callout(s, M, 6.52, 11.2, - "Validated: **wall = steady-state + faults × 0.68 µs** reproduced a 6.110 s " - "run to within **1 ms**. Kernel time stayed constant throughout; the GPU was never the variable.", - h=0.70, size=9.5) +], size=9, title="THE PROTOCOL · python/tests/ClusterFinderCUDA_perf.ipynb") +callout(s, M, 5.10, COL, + "**The protocol: re-run until getrusage() minor faults plateau** (< 200 k), " + "and quote that run. Validated: **wall = steady-state + faults × 0.68 µs** " + "reproduced a 6.110 s run to within **1 ms**.", h=0.86, size=10.5) +callout(s, M, 6.06, COL, + "**Two instrument caveats, named and then set aside.** Under nsys, host-side " + "CUDA API calls read ~4× high and wall time inflates with them, so every wall " + "time in this deck comes from an **unprofiled** run and only GPU-side " + "timestamps come from nsys. CUDA events measure a stream, not a kernel. " + "Both worked through in **A6**.", h=0.86, size=10.5, color=AMBER) +rail(s, [ + ("label", "one 100 000-frame pass · 9×9"), + ("gap", 0.12), + ("stat", "Pages faulted in", "~2.6 M", AMBER), + ("stat", "Cost inside the timer", "up to 4 s", AMBER), + ("gap", 0.08), + ("row", "Cost per fault, fitted", "0.68 µs", TEXT2), + ("row", "Major faults, all campaign", "0", ACCENT), + ("gap", 0.24), + ("note", "No disk is involved. A minor fault is the kernel finding a physical " + "page and zeroing it, which is mandatory and cannot be avoided by " + "copying faster. Only by not allocating, which is opt6."), +]) + # =========================================================== 25 · FIRST RUN s = new_slide() -chrome(s, 27, "Behind the numbers · what a user actually gets", +chrome(s, 26, "Behind the numbers · what a user actually gets", "A first run loses a third of its throughput to page faults") figure(s, "fig_first_run", 1.37, 1.70, 10.6) callout(s, M, 5.98, 5.85, @@ -1458,17 +1615,17 @@ caption(s, M, 6.76, 11.9, section("Validation · does it find the same photons", "Every number so far assumed the answers are identical", "Whether the CUDA finder returns the same clusters as the CPU.", - [("28–29", "The fair comparison"), - ("30–31", "The residual, dissected"), - ("32–33", "For users"), - ("34", "What is next")], - rng=(28, 34), col=PALE, + [("27–28", "The fair comparison"), + ("29–30", "The residual, dissected"), + ("31–32", "For users"), + ("33", "What is next")], + rng=(27, 33), col=PALE, carry=("Established", "×9.1 and ×26.5", "on the hardware floor at both cluster sizes, if the physics holds")) # ============================================ 26 · PEDESTAL UPDATE TIMING s = new_slide() -chrome(s, 28, "Validation · why a CPU twin was needed", +chrome(s, 27, "Validation · why a CPU twin was needed", "CPU and CUDA update the pedestal at different moments") figure(s, "fig_pedtiming", M - 0.15, 1.90, 12.2) callout(s, M, 5.30, 5.85, @@ -1514,7 +1671,7 @@ Frozen ships in the library as a diagnostic, not as the recommended finder.""") # =========================================================== 26 · CORRECTNESS s = new_slide() -chrome(s, 29, "Validation · isolating one variable at a time", +chrome(s, 28, "Validation · isolating one variable at a time", "CUDA and its CPU twin agree exactly: 0 in 23 million") bullets(s, M, 1.86, 12.0, [ "**ClusterFinderFrozen** makes byte-for-byte the same decisions as ClusterFinder " @@ -1548,25 +1705,26 @@ caption(s, M, 6.58, 11.9, "centre-set difference at tol = 0 · python/tests/validation_tiers.py. [f64 ped] " "and [f32 ped] are the same source built with DEVICE_PED_TYPE double / float; " "COMPUTE_TYPE is float in both, so the stencil arithmetic is identical across " - "the two rows.") + "the two rows. Either row is reproduced by rebuilding with that typedef and " + "re-running python/tests/ClusterFinderFrozen_vs_CUDA.ipynb, which counts and " + "localises every disagreement rather than only totalling them.") # ================================================ 28 · THE MISMATCH, SEEN s = new_slide() -chrome(s, 30, "Validation · the disagreement, seen", +chrome(s, 29, "Validation · the disagreement, seen", "The whole disagreement is one duplicate centre") -figure(s, "fig_mismatch147", 1.97, 1.78, 9.4) -callout(s, M, 5.82, 5.85, - "Only the **3×3 footprints of each finder's own centres** are drawn, " - "everything else is blank, so the panels differ **exactly** where the finders " - "do. Values are pedestal-subtracted ADU.", h=0.94, size=10.5) -callout(s, 6.75, 5.82, 5.85, +figure(s, "fig_mismatch147", 1.37, 1.72, 10.6) +callout(s, M, 6.30, 5.85, + "Only the **3×3 footprints of each finder's own centres** are drawn, so the " + "panels differ **exactly** where the finders do. Cell values are " + "pedestal-subtracted ADU.", h=0.78, size=10.5) +callout(s, 6.75, 6.30, 5.85, "cuda's patch is **one row taller**: a second centre directly below the one " - "both found, so the two 3×3 windows overlap. **The charge is already " - "counted**: a duplicate, not a new photon.", h=0.94, size=10.5, color=AMBER) -caption(s, M, 6.92, 11.9, - "Frame 147, the strongest of the six residuals · shipping f32 build · every " - "other centre in the patch agrees, including the ordinary photon at bottom right.") -notes(s, """The masked view, and how to read it. + "both found. **The charge is already counted** — a duplicate, not a new " + "photon.", h=0.78, size=10.5, color=AMBER) +notes(s, """Frame 147, the strongest of the six residuals, shipping f32 build. +Every other centre in the patch agrees, including the ordinary photon at bottom +right. Each panel masks every pixel that is not inside the 3x3 footprint of one of THAT finder's own cluster centres. So a pixel is visible on the left only if frozen @@ -1588,11 +1746,10 @@ reproducible from ClusterFinderFrozen_vs_CUDA.ipynb directly.""") # ========================================================== 27 · RESIDUALS s = new_slide() -chrome(s, 31, "Validation · the six residuals, dissected", +chrome(s, 30, "Validation · the six residuals, dissected", "float32 cannot tell these two pixels apart") -code(s, M, 1.92, 6.35, [ +code(s, M, 1.88, 6.35, [ "frame 147 centre (x=202, y=8) 3×3 window", - "", "raw window (ADU) pedestal-subtracted, 1 decimal", "[[4646 5282 4703] [[ 45.3 «638.4» -12.1] frozen and cuda", " [4857 5318 4950] [ 43.7 «638.4» -7.5] print the SAME", @@ -1607,18 +1764,18 @@ code(s, M, 1.92, 6.35, [ " frozen 638.382773664 >= 638.383019956 -> reject", " cuda 638.382812500 >= 638.382812500 -> «ACCEPT»", ], size=8, title="THE ONLY TEST THAT FLIPS, AND WHY") -callout(s, M, 4.62, 6.35, +callout(s, M, 4.72, 6.35, "Separation under the f64 pedestal: **0.000246 ADU**. One float32 ULP at " "4 679.6 ADU is **0.000488**. The two pixels are **half a ULP apart**; in " "float32 they are *the same number*, and the gate accepts on a tie.", - h=1.04, size=10) + h=1.00, size=10) figure(s, "fig_spectra_valid", 7.30, 1.92, 5.40) -callout(s, M, 5.82, 11.9, +callout(s, M, 5.84, 11.9, "**With the f64 pedestal there are none at all.** In the shipping f32 build " "each of the six sits **one pixel from a cluster both finders found**, a " "duplicate neighbour, never a spurious photon and never a missed one.", h=0.76, size=10.5, color=AMBER) -caption(s, M, 6.74, 11.9, +caption(s, M, 6.70, 11.9, "The same frame 147 as the previous slide, recomputed under each finder's " "decision-time pedestal. Both finders run the same gate (CPU value == max, CUDA " "!(val < max)), so the tie separates them, not the expression.") @@ -1656,7 +1813,7 @@ Consequences worth stating out loud if asked: # =========================================================== 28 · API 1 s = new_slide() -chrome(s, 32, "For users · Python API", "The fast path in eight lines") +chrome(s, 31, "For users · Python API", "The fast path in eight lines") code(s, M, 1.95, 7.6, [ "from aare import File, ClusterFinderCUDA", "", @@ -1696,7 +1853,8 @@ callout(s, 8.6, 6.35, 4.1, # =========================================================== 27 · API 2 s = new_slide() -chrome(s, 33, "For users · choosing the knobs", "Four knobs, and the one that silently truncates") +chrome(s, 32, "For users · choosing the knobs", + "Five knobs, and the one that silently truncates") hdr = [("Parameter", 1.05), ("What it does", 3.6), ("Guidance", 5.2)] y = 2.0 rect(s, M, y, 11.9, 0.4, PANEL) @@ -1705,7 +1863,7 @@ 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 may be in flight at once — an upper bound, not a " + ("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×."), @@ -1721,6 +1879,7 @@ params = [ ("register_input_buffer", "Page-locks the host array for DMA.", "Always, if the data is already in RAM. Check the pinning budget first."), ] +flagged = None for i, (p_, what, guide) in enumerate(params): if i % 2 == 0: rect(s, M, y, 11.9, 0.82, PANEL) @@ -1730,24 +1889,32 @@ for i, (p_, what, guide) in enumerate(params): run(para(tf, True, line=1.2), what, 9.5, PALE) tf = tb(s, M + 6.15, y + 0.14, 5.4, 0.6) run(para(tf, True, line=1.2), guide, 9.5, TEXT2) + if p_ == "max_clusters_per_frame": + flagged = y y += 0.80 +# The one row that loses data if you get it wrong, ringed so it is found without +# reading the table. Drawn last so the outline sits over the zebra fill. +frame_rect(s, M - 0.04, flagged - 0.05, 11.98, 0.92) +# The ring and the warning are the same point, so they carry the same colour. callout(s, M, 6.55, 11.2, "The single most common mistake: leaving **max_clusters_per_frame** too low. " "It does not error; it truncates, and every frame quietly returns the same count.", - h=0.66, size=10, color=AMBER) + h=0.66, size=10, color=RED) # =========================================================== 28 · NEXT s = new_slide() -chrome(s, 34, "Where this leaves us", +chrome(s, 33, "Where this leaves us", "The bottleneck has walked from the host, to the GPU, to the wire") cards = [ ("DONE", ACCENT, "×9.1 at 3×3, ×26.5 at 9×9", - "16.3 and 25.1 µs/frame end to end, both on their hardware floor. Against the CPU " - "twin that isolates the port the decisions are identical; the shipping f32 " - "pedestal adds 6 duplicates in 23 million."), + "16.3 and 25.1 µs/frame end to end, both sitting on their hardware floor. " + "At 3×3 that is 58 495 FPS against MÖNCH03's 1.3 kHz standard frame rate: " + "45× the detector, and ~10× its optimised 3–6 kHz ceiling."), ("DONE", ACCENT, "FP32 pedestal, safely", - "−40.5% kernel, and correct, because the variance is now accumulated on a frozen " - "per-pixel offset instead of a raw second moment."), + "−40.5 % kernel, and correct, because the variance is accumulated on a frozen " + "per-pixel offset instead of a raw second moment. Against the CPU twin that " + "isolates the port the decisions are identical; the shipped f32 pedestal adds " + "6 duplicates in 23 million."), ("NEXT", AMBER, "3×3: transfer granularity", "The 16.31 µs sustained sits 3.16 µs above the uncontended 13.15 µs H2D rate " "[f32: s4 vs s1]: " @@ -1771,35 +1938,52 @@ for i, (tag, col, title, body) in enumerate(cards): callout(s, M, 6.58, 11.2, "Full numbers, methodology and reproduction steps: **docs/ClusterFinderCUDA_benchmark_results.md** · " "notebook **python/tests/ClusterFinderCUDA_perf.ipynb** · measurement campaign **python/tests/perf/**", - h=0.66, size=9.5) + h=0.60, size=10) +notes(s, """The headroom claim, stated carefully. + +MOENCH03 runs at 1.3 kHz as standard and 3-6 kHz with optimised readout boards. +The finder does 58 495 FPS at 3x3 and 39 775 at 9x9, so it is 45x and 31x the +standard rate, and about 10x and 7x the optimised ceiling. The point is not the +multiple: it is that cluster finding has stopped being the thing that decides +how fast you can take data, on one GPU, at either cluster size. + +Two honest qualifications. First, this is throughput on frames already in host +RAM -- getting them there from the detector is a separate problem and is not +measured here. Second, at 9x9 the D2H slot is cap-sized, so the margin shrinks +if max_clusters_per_frame has to grow for a busier beam. + +What the 24-thread CPU finder does, for contrast: 6 762 FPS at 3x3, which +matches the standard mode with nothing to spare, and 1 503 FPS at 9x9, which +does not.""") # --------------------------------------------------------- divider · ANNEX -section("Annex · the evidence behind the claims", - "Kept back so the arc stays readable", - "The measurement artefacts, the rejected route, and the opt5 code.", +section("", + "Annexes", + "", [("A1", "Every engine number, reconciled"), - ("A2", "CUDA Graphs, and why they lost"), + ("A2", "The three rejected routes"), ("A3", "opt5 · the overlap code"), ("A4", "The fault model, tested"), - ("A5", "The three benchmark artefacts")], + ("A5", "The variance rewrite in full"), + ("A6", "The three benchmark artefacts")], rng=(1, N_ANNEX), col=AMBER, annex=True, - carry=("Everything so far", "34 slides", + carry=("Everything so far", "33 slides", "the arc is finished; what follows answers questions")) # =========================================================================== -# ANNEX — the measurement detail behind slides 26–27, and the rejected route A +# ANNEX — the measurement detail behind slides 24–25, and the rejected routes # =========================================================================== # ---- A1 · THE CONVENTION ------------------------------------------------- s = new_slide() -annex_chrome(s, 1, "measurement convention · expands slide 12", +annex_chrome(s, 1, "measurement convention · expands slide 18", "Uncontended, or as the pipeline runs it") bullets(s, M, 1.90, 12.0, [ "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.", + "frame**, the union of its intervals, the only number that can set a floor.", ], size=10.5) table(s, M, 2.78, 5.8, ["3×3 · µs/frame", "s1 f64", "s1 f32", "s4 f64", "s4 f32"], @@ -1818,32 +2002,38 @@ table(s, 7.0, 2.78, 5.8, ["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 — " + "**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 " + "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 " +caption(s, M, 6.62, 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.") + "contention does D2H climb to 25.24 and overtake the 23.94 f32 kernel, so any " + "claim about which engine binds must be read from the s4 columns. Source: " + "probes.csv in perf/results/2026-08-18_{f64,f32}/ (3×3) and " + "2026-08-20_{f64,f32}_capAB/ (9×9). Bandwidth arithmetic in the notes.") +notes(s, """Neither direction is faster than the other, and the bar heights say so. + +At s1, H2D moves 320 000 B in 13.15 us = 24.3 GB/s; D2H moves 120 004 B in +5.27 us = 22.8 GB/s; at 9x9, D2H moves 557 604 B in 21.95 us = 25.4 GB/s. All +three sit at 72-81 % of PCIe 4.0 x16's 31.5 GB/s, the smallest transfer paying +the most fixed cost per byte. So a taller bar in this deck always means MORE +BYTES, never a slower wire -- which is exactly why raising the cap to 1 700 +hands the 9x9 floor to D2H: it is 544.5 kB per frame either way the arm is +built, and it does not care about the typedef. + +On the last row: the floor is the LOWER of the s4 estimate and the best +unprofiled sustained rate. 16.31 vs 16.63 at 3x3 f32; 25.14 vs 25.24 at 9x9 f32; +30.01 vs 32.66 at 9x9 f64, which is the widest gap and the reason the rule +exists. A probe roofline is an estimate, never a hard denominator.""") # ---- A2 · ROUTE A · IDEA ------------------------------------------------- s = new_slide() 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) + "CUDA Graphs, a sound idea that the next act overtook", part=1, nparts=3) 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.", @@ -1861,7 +2051,7 @@ code(s, 8.5, 3.15, 4.1, [ "// per batch: one call", "«cudaGraphLaunch»(sc.graphExec, sc.stream);", ], size=8, title="ClusterFinderCUDA_graph.hpp") -callout(s, 8.5, 4.88, 4.1, +callout(s, 8.5, 5.06, 4.1, "**REJECTED**\n3×3: 39 752 FPS, inside noise of opt4.\n9×9: **11 072 FPS, 12 % slower**.", h=1.10, size=10.5, color=AMBER) caption(s, M, 6.62, 12.0, @@ -1875,7 +2065,7 @@ caption(s, M, 6.62, 12.0, # ---- A3 · ROUTE A · BUDGET ----------------------------------------------- s = new_slide() annex_chrome(s, 2, "rejected route · CUDA graphs", - "What a CUDA Graph actually saves, in microseconds", part=2, nparts=2) + "What a CUDA Graph actually saves, in microseconds", part=2, nparts=3) 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** " @@ -1892,7 +2082,7 @@ 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 A5) ~ 1.4 us/frame eliminated, unprofiled estimate", ], size=8, title="THE ARITHMETIC") -callout(s, M, 6.28, COL, +callout(s, M, 6.38, COL, "**~1.4 µs against a 16.17 µs floor = 8.7 %**, real while the host is the " "critical path, and **worth nothing after opt5**, which hides host work under " "the GPU entirely.", h=0.80, size=10, color=AMBER) @@ -1911,9 +2101,49 @@ rail(s, [ ]) +# ---- A2·3 · ACT II REJECTED ROUTES --------------------------------------- +# Was a main-arc slide. It answers "did you try just making the copy faster?", +# which is a question, not a step in the argument, so it belongs here. +s = new_slide() +annex_chrome(s, 2, "rejected routes · the result copy · expands slide 18", + "The copy is allocation-bound, not bandwidth-bound", part=3, nparts=3) +rows = [ + ("B\u2032", "One allocation per chunk", "collect_packed()", + "Removes the per-frame malloc but keeps the copy, and the copy is ~80% of the " + "cost. Worse, the replacement allocation is 1.17 GB, far above glibc's mmap " + "threshold, so it is mmap'd and munmap'd every chunk: 606 566 faults, ~21 \u00b5s/frame.", + "69.3 \u00b5s \u00b7 deleted from the API"), + ("B\u2033", "Parallel materialisation", "8-thread copy pool", + "Each worker gets its own glibc arena, which destroys the cross-run heap reuse " + "that makes the single-threaded path cheap. Faults went 9 700 \u2192 2 270 000; " + "MALLOC_ARENA_MAX=1 collapsed them back to 138 k, which is the proof.", + "+6% at best, \u221233% when results are freed promptly"), +] +y = 2.00 +for tag, name, how, body, verdict in rows: + rect(s, M, y, 11.9, 2.05, PANEL) + rect(s, M, y, 0.035, 2.05, AMBER) + tf = tb(s, M + 0.30, y + 0.22, 1.0, 0.4) + run(para(tf, True), tag, 20, AMBER, bold=True, font=MONO) + tf = tb(s, M + 1.30, y + 0.20, 6.0, 0.3) + run(para(tf, True), name, 14, PALE, bold=True) + tf = tb(s, M + 1.30, y + 0.52, 6.0, 0.3) + run(para(tf, True), how, 10, MUTED, font=MONO) + tf = tb(s, M + 1.30, y + 0.90, 10.2, 1.0) + run(para(tf, True, line=1.3), body, 10, TEXT2) + tf = tb(s, M + 1.30, y + 1.68, 10.2, 0.3) + run(para(tf, True), verdict, 10.5, AMBER, bold=True) + y += 2.25 +callout(s, M, 6.38, 11.9, + "**Copying faster does not help when the cost is the OS populating pages.** " + "The only winning move is not to allocate — which is exactly what opt6 does. " + "materialize_slot() is deliberately single-threaded and carries a comment " + "saying so, to stop the experiment being repeated.", h=0.80, size=10.5) + + # ---- A4 · THE OPT5 CODE -------------------------------------------------- s = new_slide() -annex_chrome(s, 3, "opt5 · the overlap code · expands slide 16", +annex_chrome(s, 3, "opt5 · the overlap code · expands slide 17", "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)", @@ -1958,12 +2188,13 @@ caption(s, M, 6.80, 12.0, # ---- A5 · THE FAULT MODEL ------------------------------------------------ s = new_slide() -annex_chrome(s, 4, "the fault model · expands slide 22", +annex_chrome(s, 4, "the fault model · expands slide 21", "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.", + "Slide 26 fits **0.68 µs per first-touch fault** on the **3×3 f32** ladder, where " + "wall = steady-state + faults × 0.68 µs reproduced a 6.110 s run to **1 ms**. " + "Every row below is **9×9, and an f64-vs-f32 gap**: a different cluster size and a " + "different comparison, so the rate is applied **out of sample**, never refitted.", ], size=10.5) table(s, M, 2.78, 12.0, ["step", "f64 warm (faults)", "f32 warm (faults)", "Δ wall", @@ -1978,20 +2209,103 @@ table(s, M, 2.78, 12.0, 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 " + "+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.") + "× 0.68 µs ÷ 20 000 frames, with 0.68 carried in unchanged from the 3×3 fit " + "(slide 25): nothing on this slide is tuned to make the columns agree. Observed " + "vs predicted: +13.22 / +13.37, −4.63 / −0.02, −4.54 / −4.81, −4.87 / 0.00. " + "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.") + +# ---- A5 · THE VARIANCE REWRITE IN FULL ----------------------------------- +# Was main-arc slide 21, plus the error-floor panel that used to share +# fig_cancellation. Both are the quantitative backing for slide 20's third +# bullet, and neither is needed to follow the argument. +s = new_slide() +annex_chrome(s, 5, "the variance rewrite · expands slide 20", + "The rewrite in full, and which pixels the error reached") +bullets(s, M, 1.90, 7.4, [ + "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, and report the mean " + "as **X₀ + sum/n**, so nothing downstream changes.", + "Both operands of the subtraction are now **O(rms)-sized**: the huge common " + "term is gone before the subtraction rather than after it.", +], size=10.5) +code(s, M, 3.62, 7.4, [ + "// before: both terms ~2.17e7, answer ~2000", + "var = sum2/n - mean*mean;", + "", + "// after: centred on a frozen per-pixel offset X0", + "DEVICE_PED_TYPE resid = mean - «d_pd_off»[i]; // ~O(1)", + "DEVICE_PED_TYPE var_px = sum2[i]/n - resid*resid; // no cancellation", +], size=9, title="clusterfinder_kernel.cuh") +callout(s, M, 5.42, 7.4, + "Result: the 100 % FP32 build matches the FP64 build to **3 × 10⁻⁷**, " + "70 clusters out of 233 million.", h=0.72, size=10.5) +callout(s, M, 6.30, 7.4, + "**X₀ must never be updated.** The accumulators are defined relative to it, " + "so moving it invalidates every sum already collected. Welford's online " + "variance is the other correct answer; this one was chosen because it is " + "two lines and does not change the update's arithmetic cost.", + h=0.86, size=10.5, color=AMBER) +h = figure(s, "fig_varfloor", 8.30, 2.00, 4.32) +caption(s, 8.30, 2.00 + h + 0.18, 4.32, + "Which pixels the ±3 ADU² floor actually reaches. Below rms ≈ 2 the variance " + "is lost outright; from 2 to 5 the threshold is corrupted but not clamped. " + "Source: §5–§6 of docs/pedestal_precision_f32_cancellation.md. The naive f32 " + "build's +28.06 % excess is this shaded band, integrated over the sensor.", + size=9) + +# ---- A6·1 · THE THREE ARTEFACTS ------------------------------------------ +# Main-arc slide 25 keeps only the first of these three, because it is the only +# one that moves a number the audience is shown. This is that slide as it stood. +s = new_slide() +annex_chrome(s, 6, "benchmark artefacts · expands slide 25", + "Three ways a GPU benchmark lies", part=1, nparts=4) +items = [ + ("First-touch page faults", AMBER, + "Each run materialises ~10 GB of clusters. The first pass faults in ~2.6 M " + "pages at 0.7 µs each, up to 4 s of pure OS work inside the timer.", + "Fix: re-run until getrusage() minor faults plateau (< 200 k)."), + ("CUDA-event kernel timing", AMBER, + "avg_kernel_time_ms() measures elapsed time on a stream, including waiting " + "for other streams. Under 8-stream load it over-reads by up to 3.5×.", + "Fix: Nsight Systems per-instance times; 1 stream for exclusive numbers."), + ("The profiler itself", AMBER, + "Under nsys, wall time per frame inflates ~4× from API tracing.", + "Fix: GPU op times from nsys, wall times from unprofiled runs."), +] +x = M +for title, col, body, fix in items: + rect(s, x, 2.0, 3.83, 3.15, PANEL) + rect(s, x, 2.0, 3.83, 0.035, col) + tf = tb(s, x + 0.26, 2.28, 3.3, 0.6) + run(para(tf, True, line=1.15), title, 13, PALE, bold=True) + tf = tb(s, x + 0.26, 3.02, 3.3, 1.5) + run(para(tf, True, line=1.3), body, 10, TEXT2) + tf = tb(s, x + 0.26, 4.42, 3.3, 0.65) + run(para(tf, True, line=1.25), fix, 10, ACCENT) + x += 4.03 +code(s, M, 5.4, 11.9, [ + "# every timed cell in the benchmark notebook is bracketed with:", + "mf0 = resource.getrusage(resource.RUSAGE_SELF).ru_minflt", + "... t = time.perf_counter() - t0 ...", + "print(f'minor faults: {mf1-mf0:,}') # quote the run where this plateaus", +], size=9, title="THE FAULT PROTOCOL · python/tests/ClusterFinderCUDA_perf.ipynb") +callout(s, M, 6.68, 11.9, + "Validated: **wall = steady-state + faults × 0.68 µs** reproduced a 6.110 s " + "run to within **1 ms**. Kernel time stayed constant throughout; the GPU was " + "never the variable.", h=0.62, size=10) # ---- 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) +annex_chrome(s, 6, "benchmark artefacts · expands slide 25", + "First-touch page faults: two sources, one counter", part=2, nparts=4) 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. " @@ -2010,12 +2324,12 @@ table(s, M, 2.62, 12.0, 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, [ +code(s, 6.75, 5.42, 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, +callout(s, M, 6.50, 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 " @@ -2023,8 +2337,8 @@ callout(s, M, 6.42, 11.9, # ---- A7 · EVENTS --------------------------------------------------------- s = new_slide() -annex_chrome(s, 5, "benchmark artefacts", - "CUDA events measure the stream, not the kernel", part=2, nparts=3) +annex_chrome(s, 6, "benchmark artefacts", + "CUDA events measure the stream, not the kernel", part=3, nparts=4) 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 " @@ -2042,7 +2356,7 @@ code(s, M, 4.75, COL, [ "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, +callout(s, M, 6.18, 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) @@ -2056,13 +2370,13 @@ rail(s, [ ("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."), + "subject to the sustained-rate rule on slide 18. 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) +annex_chrome(s, 6, "benchmark artefacts", + "Where nsys is sound, and where it is not", part=4, nparts=4) 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 " @@ -2085,18 +2399,28 @@ callout(s, M, 5.30, 11.9, "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.") + "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). Proof that 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. ⚠ Cap 1500: the only such numbers left in " + "the deck. Why, and what the shipped bar is, in the notes.") +notes(s, """Why this one slide is still cap 1500, and what not to read off it. + +The RUNTIME column exists only inside the trace: it is the cost of the host call +itself, which does not depend on how large the output slot is. So there is +nothing to re-run -- a cap-1700 trace would report the same ~1.65 and ~1.85 us +for cudaMemcpyAsync and cudaLaunchKernel. + +What you must NOT read off this slide is the 19.44 us D2H. That is the cap-1500 +figure. At the shipped cap of 1 700 the same engine reads 21.95 us [s1] and +25.24 us [s4], and it is the s4 value that binds the f32 build. A1 has the +shipped grid. + +The one number in the whole deck taken from _RUNTIME is the CUDA Graph launch +budget in A2, and it is quoted to a single significant figure for exactly this +reason: a ~4x inflated host-call time can support "launch cost is about 2 us, +against a 24 us floor", and nothing finer than that.""") prs.save(OUT) print(f"saved {OUT} ({len(prs.slides._sldIdLst)} slides)") diff --git a/docs/deck/make_figs.py b/docs/deck/make_figs.py index 6703d0ee..afc302bb 100644 --- a/docs/deck/make_figs.py +++ b/docs/deck/make_figs.py @@ -56,9 +56,11 @@ opt7 cuts the kernel 40 % to 23.94 us -- below the D2H bar -- so the f32 floor i D2H, not the kernel. Optimizing the kernel in bottleneck order ended by handing the constraint to the result path, which is what success looks like. """ +import re import matplotlib matplotlib.use("Agg") import matplotlib.pyplot as plt +import matplotlib.text import numpy as np from matplotlib.patches import FancyArrowPatch, Rectangle from pathlib import Path @@ -86,11 +88,103 @@ plt.rcParams.update({ }) -def save(fig, name): - fig.savefig(OUT / f"{name}.png", dpi=220, transparent=False, +DPI = 220 + +# ------------------------------------------------------------ legibility gate +# A matplotlib fontsize is in points of the FIGURE's own inches. The PNG is then +# placed on the slide at some other width, so what the audience actually reads is +# +# effective_pt = raw_pt x (placement_width_in / figure_width_in) +# +# and the second factor is invisible at the point where the font size is written. +# fig_opt2_timeline used to save 10.64 in wide and be placed at 7.9, turning a +# 6.4 pt provenance line into 4.8 pt on the screen -- two thirds of the size of +# the smallest text set directly in PowerPoint. Worse, bbox_inches="tight" grew +# the saved width to fit that very caption, so the longer the caption the smaller +# it rendered. +# +# So the placement width is now an argument, and every figure is measured against +# the deck's projection floor after it is drawn. The floor is set for a room +# where the back row is 6-7 m from the screen: on a 13.33 x 7.5 in slide, 9 pt is +# about 1/60 of the slide height, which is the conventional lower bound for +# supporting detail, and 10.5 pt (~1/50) is the bound for anything the audience is +# asked to read a number off. +MIN_EFF_PT = 9.0 +_VIOLATIONS = [] + + +def _placements(): + """Read the placement width of every figure straight out of the deck script. + + Hardcoding the widths here would drift the first time a slide is re-laid out, + and drift silently, because nothing renders the figure at both sizes. Parsing + the deck means the generator always checks against the width the deck will + actually use. Where a figure appears twice, the narrowest placement wins, + since that is the one that sets the smallest text. + """ + deck = (Path(__file__).resolve().parent / "build_fused_deck.py") + if not deck.exists(): + return {} + env = {"M": 0.7, "COL": 7.9, "RAIL_W": 3.5, "RAIL_X": 9.2, + "W": 13.333, "H": 7.5} + out = {} + for m in re.finditer(r'\b(?:card_)?figure\(s,\s*"([a-z0-9_]+)",([^)]*)\)', + deck.read_text()): + args = [a.strip() for a in m.group(2).split(",")] + if len(args) < 3: + continue + try: + w = float(eval(args[2], {}, env)) + except Exception: + continue + out[m.group(1)] = min(out.get(m.group(1), 99.0), w) + return out + + +PLACE_W = _placements() + + +def save(fig, name, place_w=None): + """Write the PNG, then check every string in it against the projection floor. + + `place_w` is the width the deck places this figure at. Passing it turns the + check on; the figure list at the bottom of this module keeps it in sync with + build_fused_deck.py, and tools/audit prints both sides. + """ + path = OUT / f"{name}.png" + texts = [(t.get_text(), t.get_fontsize()) + for t in fig.findobj(matplotlib.text.Text) + if t.get_text().strip() and t.get_visible()] + fig.savefig(path, dpi=DPI, transparent=False, bbox_inches="tight", pad_inches=0.08) plt.close(fig) - print("wrote", name) + + if place_w is None: + place_w = PLACE_W.get(name) + if place_w is None: + print("wrote", name) + return + from PIL import Image + pw = Image.open(path).size[0] / DPI + scale = place_w / pw + bad = sorted({(round(sz * scale, 2), round(sz, 1), txt[:44].replace("\n", " ")) + for txt, sz in texts if sz * scale < MIN_EFF_PT - 0.05}) + print(f"wrote {name:22s} {pw:5.2f} in -> {place_w:5.2f} in " + f"(x{scale:.3f}) min eff {min((sz * scale for _, sz in texts), default=99):.1f} pt") + for eff, raw, txt in bad: + _VIOLATIONS.append((name, eff, raw, txt)) + print(f" ILLEGIBLE {eff:5.2f} pt (set {raw:4.1f}) {txt!r}") + + +def legibility_report(): + if not _VIOLATIONS: + print(f"\nlegibility: every string in every figure renders at " + f">= {MIN_EFF_PT} pt on the slide.") + return + print(f"\nlegibility: {len(_VIOLATIONS)} strings below {MIN_EFF_PT} pt " + f"in {len({v[0] for v in _VIOLATIONS})} figures") + for name, eff, raw, txt in sorted(_VIOLATIONS): + print(f" {name:22s} {eff:5.2f} pt {txt!r}") def bare(ax, keep=("left", "bottom")): @@ -126,11 +220,11 @@ def fig_spectra_valid(): label=f"{name} ({d['totals'][name]:,})") ax.set_yscale("log") ax.set_ylim(3e2, 5e6) - ax.legend(frameon=False, fontsize=8, labelcolor=TEXT2, loc="upper right") - ax.set_ylabel("clusters / bin", fontsize=8) + ax.legend(frameon=False, fontsize=9, labelcolor=TEXT2, loc="upper right") + ax.set_ylabel("clusters / bin", fontsize=9) bare(ax, keep=("left", "bottom")) ax.set_title("cluster energy spectrum · 3×3, 10 000 frames, 23.2 M clusters", - color=MUTED, fontsize=8.5, loc="left", pad=6) + color=MUTED, fontsize=9, loc="left", pad=6) m = h["cpu"] > 0 axr.axhspan(0.999, 1.001, color=GREEN, alpha=0.18, zorder=1) @@ -140,13 +234,13 @@ def fig_spectra_valid(): dev = max(np.abs(h[n][m] / h["cpu"][m] - 1).max() for n in ("frozen", "cuda")) axr.set_ylim(0.9955, 1.0045) axr.set_yticks([0.996, 1.0, 1.004]) - axr.set_yticklabels(["−0.4 %", "0", "+0.4 %"], fontsize=7.5) - axr.set_xlabel("cluster sum [ADU]", fontsize=8) - axr.set_ylabel("vs CPU", fontsize=8) + axr.set_yticklabels(["−0.4 %", "0", "+0.4 %"], fontsize=9) + axr.set_xlabel("cluster sum [ADU]", fontsize=9) + axr.set_ylabel("vs CPU", fontsize=9) bare(axr, keep=("left", "bottom")) axr.text(0.985, 0.90, f"worst populated bin: {dev*100:.3f} % · band = ±0.1 %", transform=axr.transAxes, ha="right", va="top", color=GREEN, - fontsize=7.5) + fontsize=9) fig.subplots_adjust(hspace=0.10) save(fig, "fig_spectra_valid") @@ -253,7 +347,7 @@ def fig_arc(): ax.plot([x0 + 0.08, x1 - 0.08], [TOP * 0.955] * 2, color=col, lw=2.2, zorder=2) ax.text((x0 + x1) / 2, TOP * 0.965, label, ha="center", va="bottom", - color=col, fontsize=8.5, fontweight="bold") + color=col, fontsize=9, fontweight="bold") ax.bar(x, fps, width=0.62, color=colors, zorder=3, linewidth=0) @@ -263,7 +357,7 @@ def fig_arc(): ax.axhspan(61312, 61859, color=GREEN, alpha=0.20, zorder=1) ax.axhline(61859, color=GREEN, lw=1.2, ls="--", zorder=4) ax.text(-0.42, 63200, "H2D floor · 61–62 k FPS · the GPU cannot be fed faster", - color=GREEN, fontsize=8.5, ha="left", va="bottom") + color=GREEN, fontsize=9, ha="left", va="bottom") for xi, (f, s) in enumerate(zip(fps, spd)): ax.text(xi, f + 1100, f"{f:,}", ha="center", va="bottom", @@ -272,7 +366,7 @@ def fig_arc(): ha="center", va="top", color=BG, fontsize=9, fontweight="bold") ax.set_xticks(x) - ax.set_xticklabels(steps, fontsize=8.5, color=TEXT2) + ax.set_xticklabels(steps, fontsize=9, color=TEXT2) ax.set_ylim(0, TOP) ax.set_yticks([]) bare(ax, keep=("bottom",)) @@ -308,7 +402,7 @@ def fig_arc_9x9(): ax.plot([x0 + 0.08, x1 - 0.08], [TOP * 0.955] * 2, color=col, lw=2.2, zorder=2) ax.text((x0 + x1) / 2, TOP * 0.965, label, ha="center", va="bottom", - color=col, fontsize=8.5, fontweight="bold") + color=col, fontsize=9, fontweight="bold") ax.bar(x, fps, width=0.58, color=colors, zorder=3, linewidth=0) @@ -324,10 +418,10 @@ def fig_arc_9x9(): # completely that it handed the constraint to the result path. ax.hlines(33323, -0.5, 4.5, color=GREEN, lw=1.3, ls="--", zorder=4) ax.text(-0.42, 33900, "f64 KERNEL floor · 33 323 FPS (nsys estimates 30 621)", - color=GREEN, fontsize=8.5, va="bottom") + color=GREEN, fontsize=9, va="bottom") ax.hlines(39775, 4.5, 5.5, color=AMBER, lw=1.3, ls="--", zorder=4) ax.text(5.46, 45600, "f32 D2H floor · 39 775 FPS\n(nsys estimates 39 614)", - color=AMBER, fontsize=8.5, ha="right", va="bottom", linespacing=1.35) + color=AMBER, fontsize=9, ha="right", va="bottom", linespacing=1.35) ax.add_patch(FancyArrowPatch((4.62, 34100), (4.62, 39100), arrowstyle="-|>", mutation_scale=11, color=AMBER, lw=1.5, zorder=5)) ax.text(3.30, 37600, "−40 % kernel → D2H binds instead", color=AMBER, @@ -340,7 +434,7 @@ def fig_arc_9x9(): ha="center", va="top", color=BG, fontsize=9, fontweight="bold") ax.set_xticks(x) - ax.set_xticklabels(steps, fontsize=8.5, color=TEXT2) + ax.set_xticklabels(steps, fontsize=9, color=TEXT2) ax.set_ylim(0, TOP) ax.set_yticks([]) bare(ax, keep=("bottom",)) @@ -383,7 +477,7 @@ def fig_overhead(): fontsize=9.5, fontweight="bold") if e > ymax * 0.05: ax.text(xi, f + e / 2, f"+{e:.0f}", ha="center", va="center", - color=BG, fontsize=8.5, fontweight="bold") + color=BG, fontsize=9, fontweight="bold") ax.set_xticks(x) ax.set_xticklabels(steps, color=TEXT2, fontsize=9) ax.set_ylim(0, ymax) @@ -395,9 +489,9 @@ def fig_overhead(): # label the two segments in place — the floor takes the act colour, so a # colour-keyed legend would be wrong axes[0].text(0, 16.17 / 2, "GPU\nfloor", ha="center", va="center", color=BG, - fontsize=8, fontweight="bold") + fontsize=9, fontweight="bold") axes[0].annotate("host excess", xy=(0.30, 25), xytext=(1.15, 41), - color=TEXT2, fontsize=8.5, + color=TEXT2, fontsize=9, arrowprops=dict(arrowstyle="-", color=MUTED, lw=0.9)) axes[1].text(1.5, 108, "Act II removes the host bar", color=PALE, fontsize=9, @@ -481,7 +575,7 @@ def fig_streams(): for i in range(3): _frame_bars(ax, 1.0, i * FR) ax.set_ylim(0.4, 2.3) - ax.text(3 * FR + 6, 1.34, "one engine at a time", color=MUTED, fontsize=7.5, + ax.text(3 * FR + 6, 1.34, "one engine at a time", color=MUTED, fontsize=9, va="center") # --- opt2: 4 streams, barrier after each round. The drain is not drawn in; @@ -493,7 +587,7 @@ def fig_streams(): 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") + color=AMBER, fontsize=9, ha="center", va="bottom") ax.set_ylim(-0.4, 4.9) # --- opt3: no barriers, continuous @@ -502,7 +596,7 @@ def fig_streams(): _draw_schedule(ax, frames) ax.set_ylim(-1.5, 4.5) ax.text(0, -0.25, "streams never wait on each other — the H2D engine never goes idle", - color=ACCENT, fontsize=7.5, va="top") + color=ACCENT, fontsize=9, va="top") titles = ["opt1 · 1 stream, synchronous", "opt2 · 4 streams, sync barrier per round", "opt3 · 4 streams, barriers removed"] @@ -514,14 +608,10 @@ def fig_streams(): handles = [Rectangle((0, 0), 1, 1, color=c) for c in (AMBER, ACCENT, PALE)] axes[0].legend(handles, ["H2D copy", "kernel", "D2H copy"], frameon=False, - fontsize=8, labelcolor=TEXT2, ncol=3, loc="lower right", + fontsize=9, 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") + axes[2].set_xlabel("time →", color=MUTED, fontsize=9, loc="left") 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") @@ -534,68 +624,71 @@ def fig_pinning(): def box(x, y, w, h, label, sub=""): ax.add_patch(Rectangle((x, y), w, h, facecolor=PANEL, edgecolor=RULE, lw=1)) ax.text(x + w / 2, y + h / 2 + 0.26, label, ha="center", va="center", - color=PALE, fontsize=8.5, fontweight="bold") + color=PALE, fontsize=9.5, fontweight="bold") ax.text(x + w / 2, y + h / 2 - 0.34, sub, ha="center", va="center", - color=MUTED, fontsize=7) + color=MUTED, fontsize=9.5) def arrow(x0, x1, y, color, label): ax.add_patch(FancyArrowPatch((x0, y), (x1, y), arrowstyle="-|>", mutation_scale=10, color=color, lw=1.6)) ax.text((x0 + x1) / 2, y + 0.22, label, ha="center", va="bottom", - color=color, fontsize=7) + color=color, fontsize=9.5) - ax.text(0, 5.95, "PAGEABLE · before opt4", color=AMBER, fontsize=8.5, + ax.text(0, 5.95, "PAGEABLE · before opt4", color=AMBER, fontsize=9.5, fontweight="bold") box(0, 4.05, 2.5, 1.1, "numpy array", "pageable") box(4.0, 4.05, 2.4, 1.1, "driver staging", "hidden pinned buf") box(7.9, 4.05, 2.5, 1.1, "GPU", "device memory") arrow(2.5, 4.0, 4.60, AMBER, "memcpy") arrow(6.4, 7.9, 4.60, AMBER, "DMA") - ax.text(0, 3.62, "every transfer is copied twice", color=MUTED, fontsize=7) + ax.text(0, 3.62, "every transfer is copied twice", color=MUTED, fontsize=9.5) - ax.text(0, 2.75, "PINNED · opt4", color=ACCENT, fontsize=8.5, fontweight="bold") + ax.text(0, 2.75, "PINNED · opt4", color=ACCENT, fontsize=9.5, fontweight="bold") box(0, 0.85, 2.5, 1.1, "numpy array", "page-locked") box(7.9, 0.85, 2.5, 1.1, "GPU", "device memory") - arrow(2.5, 7.9, 1.40, ACCENT, "DMA — engine reads host RAM directly") + arrow(2.5, 7.9, 1.40, ACCENT, "DMA · reads host RAM directly") ax.text(0, 0.42, "no staging copy, no page faults, fully async", - color=MUTED, fontsize=7) + color=MUTED, fontsize=9.5) # the rule, in one inset: pinning pays only where H2D is the tallest bar ax2 = fig.add_axes([0.70, 0.16, 0.30, 0.62]) x = np.arange(2) - before = [34.26, 82.17] - after = [25.98, 80.44] + # 3x3 from 2026-08-18_f64, 9x9 from 2026-08-20_f64_cap1700 -- NOT the + # cap-1500 ladder, which read 82.17 -> 80.44 and put a stale x1.02 on the + # slide for a step that actually buys x1.03. + before = [34.26, 82.44] + after = [25.98, 79.83] ax2.bar(x - 0.19, before, width=0.36, color=AMBER, zorder=3, label="opt3") ax2.bar(x + 0.19, after, width=0.36, color=ACCENT, zorder=3, label="opt4") for xi, (b, a) in enumerate(zip(before, after)): - ax2.text(xi - 0.19, b + 2, f"{b:.0f}", ha="center", color=TEXT2, fontsize=7.5) - ax2.text(xi + 0.19, a + 2, f"{a:.0f}", ha="center", color=TEXT2, fontsize=7.5) - ax2.text(xi, 92, f"×{b / a:.2f}", ha="center", color=PALE, fontsize=9, + ax2.text(xi - 0.19, b + 2, f"{b:.0f}", ha="center", color=TEXT2, fontsize=9.5) + ax2.text(xi + 0.19, a + 2, f"{a:.0f}", ha="center", color=TEXT2, fontsize=9.5) + ax2.text(xi, 92, f"×{b / a:.2f}", ha="center", color=PALE, fontsize=9.5, fontweight="bold") ax2.set_xticks(x) ax2.set_xticklabels(["3×3\nH2D-bound", "9×9\nkernel-bound"], color=TEXT2, - fontsize=7.5) + fontsize=9.5) ax2.set_ylim(0, 104); ax2.set_yticks([]); bare(ax2, keep=("bottom",)) - ax2.set_title("µs / frame", color=MUTED, fontsize=7.5, pad=6) - ax2.legend(frameon=False, fontsize=7, labelcolor=TEXT2, loc="center left") + ax2.set_title("µs / frame", color=MUTED, fontsize=9.5, pad=6) + ax2.legend(frameon=False, fontsize=9.5, labelcolor=TEXT2, loc="center left") save(fig, "fig_pinning") # ----------------------------------------------- 6. graphs (rejected route) def fig_graphs(): - fig = plt.figure(figsize=(7.7, 3.0)) - ax = fig.add_axes([0, 0.30, 0.66, 0.70]) + fig = plt.figure(figsize=(7.6, 2.55)) + ax = fig.add_axes([0, 0.07, 0.66, 0.90]) ax.axis("off"); ax.set_xlim(0, 11.2); ax.set_ylim(0, 4.6) def node(x, y, w, h, t, fc): ax.add_patch(Rectangle((x, y), w, h, facecolor=fc, edgecolor="none")) ax.text(x + w / 2, y + h / 2, t, ha="center", va="center", - color=BG, fontsize=7.5, fontweight="bold") + color=BG, fontsize=9.5, fontweight="bold") ops = [("H2D", AMBER), ("kernel", ACCENT), ("D2H", PALE)] * 2 ax.text(0, 4.15, "WITHOUT GRAPHS · one driver call per operation, every frame", - color=AMBER, fontsize=8.5, fontweight="bold") + color=AMBER, fontsize=9.5, fontweight="bold") for i, (t, c) in enumerate(ops): x = 0.1 + i * 1.62 node(x, 2.85, 1.4, 0.6, t, c) @@ -603,10 +696,10 @@ def fig_graphs(): arrowstyle="-|>", mutation_scale=7, color=MUTED, lw=0.9)) ax.text(11.1, 3.15, "CPU cost\n≈ 6 launches", ha="right", va="center", - color=MUTED, fontsize=7.5) + color=MUTED, fontsize=9.5) ax.text(0, 2.18, "WITH GRAPHS · record once, replay with one launch", - color=ACCENT, fontsize=8.5, fontweight="bold") + color=ACCENT, fontsize=9.5, fontweight="bold") ax.add_patch(Rectangle((0.1, 0.72), 9.20, 1.15, facecolor=PANEL, edgecolor=ACCENT, lw=1.2)) for i, (t, c) in enumerate(ops): @@ -614,31 +707,26 @@ def fig_graphs(): ax.add_patch(FancyArrowPatch((0.8, 2.02), (0.8, 1.90), arrowstyle="-|>", mutation_scale=8, color=ACCENT, lw=1.3)) ax.text(11.1, 1.30, "CPU cost\n≈ 1 launch", ha="right", va="center", - color=ACCENT, fontsize=7.5, fontweight="bold") + color=ACCENT, fontsize=9.5, fontweight="bold") # the verdict - ax2 = fig.add_axes([0.72, 0.30, 0.28, 0.62]) + ax2 = fig.add_axes([0.73, 0.13, 0.27, 0.78]) x = np.arange(2) - opt4 = [25.98, 80.44] - graph = [25.16, 94.78] + opt4 = [25.98, 79.83] # 3x3 2026-08-18_f64 · 9x9 2026-08-20_f64_cap1700 + graph = [25.16, 90.32] # cap 1500 read 80.44 / 94.78 and said "18 % slower" ax2.bar(x - 0.19, opt4, width=0.36, color=ACCENT, zorder=3, label="opt4") ax2.bar(x + 0.19, graph, width=0.36, color=AMBER, zorder=3, label="graphs") for xi, (a, g) in enumerate(zip(opt4, graph)): - ax2.text(xi - 0.19, a + 2.5, f"{a:.0f}", ha="center", color=TEXT2, fontsize=7.5) - ax2.text(xi + 0.19, g + 2.5, f"{g:.0f}", ha="center", color=TEXT2, fontsize=7.5) - ax2.text(1, 108, "18 % SLOWER", ha="center", color=AMBER, fontsize=8, + ax2.text(xi - 0.19, a + 2.5, f"{a:.0f}", ha="center", color=TEXT2, fontsize=9.5) + ax2.text(xi + 0.19, g + 2.5, f"{g:.0f}", ha="center", color=TEXT2, fontsize=9.5) + ax2.text(1, 100, "−12 % THROUGHPUT", ha="center", color=AMBER, fontsize=9.5, fontweight="bold") ax2.set_xticks(x) - ax2.set_xticklabels(["3×3", "9×9"], color=TEXT2, fontsize=8) + ax2.set_xticklabels(["3×3", "9×9"], color=TEXT2, fontsize=9.5) ax2.set_ylim(0, 120); ax2.set_yticks([]); bare(ax2, keep=("bottom",)) - ax2.set_title("µs / frame", color=MUTED, fontsize=7.5, pad=6) - ax2.legend(frameon=False, fontsize=7, labelcolor=TEXT2, loc="upper left") + ax2.set_title("µs / frame", color=MUTED, fontsize=9.5, pad=6) + ax2.legend(frameon=False, fontsize=9.5, labelcolor=TEXT2, loc="upper left") - fig.text(0.02, 0.10, - "REJECTED — launch overhead stops binding one step later. The graph finder " - "never got the chunked pipeline of opt5,\nand its original advantage is " - "swamped by an overlap it does not have.", - color=AMBER, fontsize=8, fontweight="bold", va="top") save(fig, "fig_graphs") @@ -657,7 +745,7 @@ def fig_resultpath(): 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") + ax.text(-0.55, floor + 1.2, "GPU floor", color=GREEN, fontsize=10, ha="left") ax.text(0, floor + 1.4, f"{floor:.1f} µs", ha="center", color=PALE, fontsize=10, fontweight="bold") ax.text(1, copy_us + 1.4, f"≈{copy_us:.0f} µs", ha="center", color=PALE, @@ -665,13 +753,13 @@ def fig_resultpath(): ax.set_xticks([0, 1]) ax.set_xticklabels(["GPU per frame\n(H2D ∥ kernel ∥ D2H)", "host copy per frame\ncollect() memcpy + malloc"], - color=TEXT2, fontsize=8.5) + color=TEXT2, fontsize=10) ax.set_xlim(-0.6, 1.7); ax.set_ylim(0, 52); ax.set_yticks([]) bare(ax, keep=("bottom",)) ax.set_title(title, color=PALE, fontsize=10, pad=10, loc="left") ax.text(1.68, 46, gain, color=col, fontsize=16, fontweight="bold", ha="right") - ax.text(1.68, 40, "opt5 → opt6", color=MUTED, fontsize=8, ha="right") - ax.text(-0.55, -11, verdict, color=col, fontsize=8.5, fontweight="bold", + ax.text(1.68, 40, "opt5 → opt6", color=MUTED, fontsize=10, ha="right") + ax.text(-0.55, -11, verdict, color=col, fontsize=10, fontweight="bold", va="top") fig.subplots_adjust(bottom=0.30) save(fig, "fig_resultpath") @@ -700,8 +788,8 @@ def fig_f32_kernel(): "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") + ax.barh(y + h / 2, f64, height=h, color=ACCENT, zorder=3, label="f64 ped") + ax.barh(y - h / 2, f32, height=h, color=AMBER, 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)): @@ -710,87 +798,154 @@ def fig_f32_kernel(): # 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) + fontsize=10, fontweight="bold" if top else "normal") + # A bar pair spans y +- h, so the delta and the verdict need a full h of + # clearance beyond that or they sit on the bar they are labelling. The + # y limits carry that margin explicitly rather than relying on the + # default padding, which the larger type had eaten. + ax.text(f64[0] * 0.42, -h - 0.30, delta, ha="center", va="bottom", + color=PALE, fontsize=10.5, fontweight="bold") + ax.text(0.985, 2 + h + 0.34, verdict, transform=ax.get_yaxis_transform(), + ha="right", va="top", color=PALE, fontsize=10, fontweight="bold") + ax.set_yticks(y); ax.set_yticklabels(labels, color=TEXT2, fontsize=10) + ax.invert_yaxis(); ax.set_xlim(0, 60); ax.set_xticks([]) + ax.set_ylim(3.05, -0.98) bare(ax, keep=("left",)) - ax.set_title(title, color=PALE, fontsize=9, pad=10, loc="left") + ax.set_title(title, color=PALE, fontsize=10, 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) + ax2.axvline(25.24, color=MUTED, lw=1.0, ls="--", zorder=1) # under the labels + # parked between the D2H and H2D rows, the one region no label reaches + ax1.legend(frameon=False, fontsize=10, labelcolor=TEXT2, loc="center right", + bbox_to_anchor=(1.0, 0.40)) + fig.subplots_adjust(bottom=0.20, top=0.86, left=0.09, right=0.99, wspace=0.30) save(fig, "fig_f32_kernel") # --------------------------------------------------------- 9. cancellation 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", "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.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² · each operand on a grid of 2 ADU²", - color=MUTED, fontsize=8, pad=8) + """The trap and what it did to the physics, on one pair of axes. - # 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. + LEFT is the cancellation itself: two ~2.17e7 operands, an answer of ~2025 for + a normal pixel and ~9 for a quiet one, against a fixed +-3 ADU^2 f32 error + that does not shrink with the answer. + + RIGHT is the consequence. The f64/CPU curve is MEASURED -- validation_tiers.json, + 23.2 M clusters over 10 000 frames. The f32 curve is RECONSTRUCTED, not + measured: the excess is placed where SS9 of + docs/pedestal_precision_f32_cancellation.md derives it (a threshold that has + collapsed from ~225 ADU to ~0 admits the whole positive side of the pixel's + distribution, so the population is smeared upward from zero) and its AREA is + set to the measured +28.06 % from SS1. Shape schematic, area measured, and the + caption on the slide says so. Drawing a measured f32 spectrum would be better; + that build is not one anyone should keep around to re-run. + + The variance-vs-rms error floor that used to sit in this panel is now + fig_varfloor, in the annex: it is the quantitative version of the same claim + and it was competing with the physics for the audience's attention. + """ + import json + fig, (ax, ax2) = plt.subplots(1, 2, figsize=(7.9, 2.34), + gridspec_kw={"width_ratios": [1.0, 1.28]}) + + names = ["E[X²]\n2.17e7", "mean²\n2.17e7", "var\nbulk 2025", "var\nquiet 9"] + ax.bar([0, 1], [2.17e7, 2.17e7], width=0.52, color=[PALE, PALE], zorder=3) + ax.bar([2], [2025], width=0.52, color=AMBER, zorder=3) + ax.bar([3], [9], width=0.52, color=AMBER, zorder=3) + ax.set_yscale("log"); ax.set_ylim(1, 8e8) + ax.set_xticks([0, 1, 2, 3]) + ax.set_xticklabels(names, color=TEXT2, fontsize=9, linespacing=1.35) + ax.set_yticks([1e0, 1e2, 1e4, 1e6, 1e8]) + ax.tick_params(labelsize=9) + ax.axhline(3, color=ACCENT, lw=1.4, ls="--", zorder=4) + ax.text(3.52, 1.1e7, "±3 ADU², the same for\nevery pixel — so it is the\n" + "quiet ones it swallows", color=ACCENT, fontsize=9, ha="right", + va="center", linespacing=1.4) + ax.plot([3.0, 3.0], [3.6, 1.2e6], color=ACCENT, lw=0.7, ls=":", zorder=4) + bare(ax) + ax.set_title("var = E[X²] − mean²", color=MUTED, fontsize=9.5, pad=8, loc="left") + + d = json.loads((Path(__file__).resolve().parent + / "validation_tiers.json").read_text()) + e = np.array(d["edges"]) + ctr = 0.5 * (e[1:] + e[:-1]) + bw = e[1] - e[0] + good = np.array(d["hists"]["cpu"], dtype=float) + + # +28.06 % more clusters (SS1), smeared upward from ~0 by a gate that has + # collapsed to 0 sigma. Exponential with a 400 ADU scale: shape schematic, + # integral set to the measured excess. + LAM = 400.0 + excess = good.sum() * 0.2806 * (bw / LAM) * np.exp(-np.clip(ctr, 0, None) / LAM) + ax2.step(ctr, good, where="mid", color=ACCENT, lw=1.8, zorder=4, + label="f64 pedestal · measured") + ax2.step(ctr, good + excess, where="mid", color=AMBER, lw=1.5, zorder=3, + label="f32 pedestal · +28.06 %") + ax2.fill_between(ctr, good, good + excess, step="mid", color=AMBER, + alpha=0.22, lw=0, zorder=2) + ax2.set_yscale("log") + ax2.set_ylim(2e2, 6e7) + ax2.set_xlim(0, 2600) + ax2.set_xticks([0, 1000, 2000]) + ax2.tick_params(labelsize=9, colors=MUTED) + ax2.set_xlabel("cluster energy [ADU]", color=TEXT2, fontsize=9.5) + bare(ax2, keep=("left", "bottom")) + ax2.legend(frameon=False, fontsize=9, labelcolor=TEXT2, loc="lower right", + handlelength=1.2, borderaxespad=0.3) + ax2.annotate("clusters below the 5σ cut\nthat should not be there", + xy=(150, 1.1e6), xytext=(560, 8.0e6), color=AMBER, fontsize=9, + linespacing=1.4, zorder=6, + arrowprops=dict(arrowstyle="->", color=AMBER, lw=0.9)) + ax2.text(1195, 3.4e6, "Cu Kα", color=PALE, fontsize=9, ha="center") + ax2.set_title("cluster-energy spectrum, 3×3", color=MUTED, fontsize=9.5, + pad=8, loc="left") + fig.subplots_adjust(left=0.075, right=0.995, top=0.86, bottom=0.20, wspace=0.24) + save(fig, "fig_cancellation") + + +def fig_varfloor(): + """The quantitative version of the left panel: which pixels the error reaches. + + 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. + """ + fig, ax2 = plt.subplots(figsize=(4.35, 3.05)) 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, + fontsize=9.5, 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 (ADU²)", color=TEXT2, fontsize=8.5) + ax2.set_xlabel("pixel rms (ADU)", color=TEXT2, fontsize=9.5) + ax2.set_ylabel("variance (ADU²)", color=TEXT2, fontsize=9.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) + ax2.tick_params(labelsize=9.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) + color=AMBER, fontsize=9.5, va="top", linespacing=1.45, 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, + color=MUTED, fontsize=9.5, va="top", linespacing=1.45, zorder=6) + ax2.text(3.05, 31.0, "true variance = rms²", color=PALE, fontsize=9.5, zorder=6) + ax2.text(5.90, 5.6, "f32 error floor ±3–4 ADU²", color=ACCENT, fontsize=9.5, ha="right", zorder=6) - save(fig, "fig_cancellation") + fig.subplots_adjust(left=0.13, right=0.98, top=0.97, bottom=0.15) + save(fig, "fig_varfloor") + + +fig_varfloor() # ------------------------------- 10. the same kernel, measured five ways @@ -896,7 +1051,7 @@ def fig_overlap(): ax.add_patch(Rectangle((x, y), w, lane_h, facecolor=col, edgecolor=BG, linewidth=1.4, zorder=3)) ax.text(x + w / 2, y + lane_h / 2, txt, ha="center", va="center", - color=tcol, fontsize=8.5, fontweight="bold", zorder=4) + color=tcol, fontsize=10.5, fontweight="bold", zorder=4) # ---- serial: G H G H G H G H, strictly alternating ---------------------- yG, yH = 3.30, 2.78 @@ -918,13 +1073,13 @@ def fig_overlap(): for y, lbl in ((yG, "GPU"), (yH, "host"), (yG2, "GPU"), (yH2, "host")): ax.text(-0.25, y + lane_h / 2, lbl, ha="right", va="center", - color=TEXT2, fontsize=9) + color=TEXT2, fontsize=10.5) ax.text(-0.25, yG + lane_h + 0.30, "submit → collect, serialized (opt4)", - ha="left", va="bottom", color=TEXT2, fontsize=9.5, fontweight="bold") + ha="left", va="bottom", color=TEXT2, fontsize=10.5, fontweight="bold") ax.text(-0.25, yG2 + lane_h + 0.30, "submit(i+1) before collect(i) (opt5)", - ha="left", va="bottom", color=AMBER, fontsize=9.5, fontweight="bold") + ha="left", va="bottom", color=AMBER, fontsize=10.5, fontweight="bold") for x, y0, y1, col in ((serial_end, yH, yG + lane_h, MUTED), (pipe_end, yH2, yG2 + lane_h, AMBER)): @@ -935,12 +1090,12 @@ def fig_overlap(): arrowprops=dict(arrowstyle="<|-|>", color=AMBER, lw=1.5)) ax.text((pipe_end + serial_end) / 2, 0.20, "saved: min(GPU, host) per chunk", ha="center", va="top", - color=AMBER, fontsize=9.5, fontweight="bold") + color=AMBER, fontsize=10.5, fontweight="bold") ax.text(serial_end + 0.25, yG + lane_h / 2, "GPU + host per chunk", - ha="left", va="center", color=MUTED, fontsize=9) + ha="left", va="center", color=MUTED, fontsize=10.5) ax.text(pipe_end + 0.25, yG2 + lane_h / 2, "max(GPU, host) per chunk", - ha="left", va="center", color=AMBER, fontsize=9, fontweight="bold") + ha="left", va="center", color=AMBER, fontsize=10.5, fontweight="bold") ax.set_xlim(-1.6, serial_end + 4.0) ax.set_ylim(0, 4.25) @@ -1051,7 +1206,7 @@ def fig_mismatch147(): cmap = plt.cm.viridis.copy() cmap.set_bad(PANEL) - fig, axes = plt.subplots(1, 2, figsize=(11.2, 4.3)) + fig, axes = plt.subplots(1, 2, figsize=(9.0, 3.9)) for ax, name, col in ((axes[0], "frozen", PALE), (axes[1], "cuda", ACCENT)): w, m = sub[name], msk[name] ax.imshow(np.ma.masked_where(~m, w), cmap=cmap, vmin=0, vmax=vmax, @@ -1060,10 +1215,10 @@ def fig_mismatch147(): for j in range(w.shape[1]): if m[i, j]: ax.text(j, i, f"{w[i, j]:.0f}", ha="center", va="center", - fontsize=4.6, + fontsize=6.4, zorder=7, color="white" if w[i, j] < 0.55 * vmax else "black") for (cx, cy) in cen[name]: - ax.plot(cx - X0, cy - Y0, ".", color="#FF4B4B", ms=6, zorder=5) + ax.plot(cx - X0, cy - Y0, ".", color="#FF4B4B", ms=4.5, zorder=4) for (cx, cy) in only: if name == "cuda": ax.add_patch(plt.Circle((cx - X0, cy - Y0), 2.6, fill=False, @@ -1079,7 +1234,7 @@ def fig_mismatch147(): axes[0].text(-0.03, 0.5, f"frame {d['fid']}\nzoom on (202, 8)", transform=axes[0].transAxes, rotation=90, ha="right", - va="center", color=MUTED, fontsize=8, linespacing=1.4) + va="center", color=MUTED, fontsize=9, linespacing=1.4) # anchored in DATA coordinates: the window is no longer square (the cut is # clipped by the top edge of the frame), so axes fractions do not track the # pixel once matplotlib letterboxes the image to keep aspect. @@ -1087,7 +1242,7 @@ def fig_mismatch147(): axes[1].annotate("cuda keeps this one;\nfrozen does not", xy=(ox - X0 + 2.9, oy - Y0), xytext=(0.99, 0.93), xycoords="data", textcoords="axes fraction", - color=AMBER, fontsize=8.5, fontweight="bold", ha="right", + color=AMBER, fontsize=9, fontweight="bold", ha="right", linespacing=1.35, arrowprops=dict(arrowstyle="-|>", color=AMBER, lw=1.3)) fig.subplots_adjust(wspace=0.06) @@ -1111,7 +1266,7 @@ def fig_regpressure(): ("3×3", 38, 6, ACCENT), ("9×9", 128, 2, AMBER), ] - fig, ax = plt.subplots(figsize=(11.2, 2.55)) + fig, ax = plt.subplots(figsize=(7.4, 2.15)) y, ticks, labels = 0.0, [], [] for name, regs, blocks, col in rows: @@ -1123,7 +1278,7 @@ def fig_regpressure(): ax.barh(y, seg - 0.5, left=b * seg + 0.25, height=0.52, color=col, zorder=3, linewidth=0) full = abs(frac - 100.0) < 0.6 - ax.text(frac + 1.5, y, f"{frac:.0f} %", va="center", fontsize=9, + ax.text(frac + 1.5, y, f"{frac:.0f} %", va="center", fontsize=10, color=PALE if full else TEXT2, fontweight="bold" if full else "normal") ticks.append(y); labels.append(f"{name} · {kind}") @@ -1131,19 +1286,19 @@ def fig_regpressure(): y -= 0.45 ax.axvline(100, color=MUTED, lw=1.1, ls="--", zorder=4) - ax.text(99, 1.05, "capacity of one SM", color=MUTED, fontsize=8.5, + ax.text(99, 1.05, "capacity of one SM", color=MUTED, fontsize=10, ha="right") ax.set_yticks(ticks) - ax.set_yticklabels(labels, color=TEXT2, fontsize=9) + ax.set_yticklabels(labels, color=TEXT2, fontsize=10) ax.set_xlim(0, 152); ax.set_xticks([]) ax.set_ylim(y + 0.6, 1.5) bare(ax, keep=("left",)) ax.tick_params(axis="y", length=0) ax.text(115, ticks[0] - 0.5, "6 blocks resident\n38 × 256 = 9 728 regs each", - color=ACCENT, fontsize=8.5, va="center", linespacing=1.4) + color=ACCENT, fontsize=10, va="center", linespacing=1.4) ax.text(115, ticks[2] - 0.5, "2 blocks resident\n128 × 256 = 32 768 regs each", - color=AMBER, fontsize=8.5, va="center", linespacing=1.4) + color=AMBER, fontsize=10, va="center", linespacing=1.4) fig.subplots_adjust(left=0.16, right=0.99, top=0.88, bottom=0.06) save(fig, "fig_regpressure") @@ -1159,7 +1314,7 @@ fig_regpressure() 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", + fontsize=9.5, labelcolor=TEXT2, ncol=3, loc="lower right", bbox_to_anchor=(1.02, y), handlelength=1.1) @@ -1181,9 +1336,9 @@ def fig_opt1_timeline(): 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", + color=AMBER, fontsize=9, 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, + ax.text(3 * PER + 4, 1.34, "one engine\nat a time", color=TEXT2, fontsize=9, va="center", linespacing=1.5) ax.set_xlim(-4, 3 * PER + 40) ax.set_ylim(0.02, 2.45) @@ -1207,7 +1362,7 @@ def fig_opt2_timeline(): _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") + fontsize=9, 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 @@ -1222,20 +1377,16 @@ def fig_opt2_timeline(): 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") + fontsize=9, 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) + color=TEXT2, fontsize=9, va="center", linespacing=1.5) _engine_legend(ax, y=0.94) - ax.set_xlabel("time →", color=MUTED, fontsize=8.5, loc="left") + ax.set_xlabel("time →", color=MUTED, fontsize=9, 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") @@ -1293,7 +1444,7 @@ def fig_f32_absolute(): 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_yticklabels(["0", "25", "50", "75"], fontsize=8.5) 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)], @@ -1316,7 +1467,7 @@ def fig_f32_absolute(): 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_yticklabels(["0", "2", "4"], fontsize=8.5) 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} %" @@ -1386,3 +1537,110 @@ def fig_variance_rewrite(): fig_variance_rewrite() + + +# ---------------------------- 16. the measurement convention: s1, s4 and the floor +def fig_measure(): + """What "busy per frame" means, and why it is not the sum of durations. + + LEFT is a real schedule at 9x9 proportions, one lane per STREAM, produced by + _schedule() rather than drawn: the copy lanes are single FIFO resources -- + H2D_overlap and D2H_overlap are 1.000 in every row of probes.csv, because the + GPU has one copy engine per direction -- so those bars stagger, while kernels + from different streams sit on top of each other in time. Under the lanes, the + two union strips are computed from that same schedule, which is the whole + point: the engine's busy time is the union of its intervals. + + No microseconds on the left panel. The schedule reproduces the SHAPE of the + 9x9 pipeline but not its exact overlap factor, and putting numbers on a + schematic next to a panel of measured ones invites them to be read across. + + RIGHT is measured: s4 engine occupancy at 9x9 f64 cap 1700 from probes.csv + (20.77 / 32.66 / 25.25) and the best unprofiled sustained rate from + ladder_9x9.csv opt6 (30.01). + """ + H, K, D = 21, 43, 25 + NF, NS = 8, 4 + frames, _ = _schedule(NF, NS, H=H, K=K, D=D) + + fig, (ax, ax2) = plt.subplots(1, 2, figsize=(11.6, 2.95), + gridspec_kw={"width_ratios": [1.62, 1]}) + + # ---- left: four stream lanes, then the union each engine actually sees + LANE, TOP = 0.52, 5.0 + for st, h0, k0, d0 in frames: + y = TOP - st * 0.72 + for t0, dur, col in ((h0, H, AMBER), (k0, K, ACCENT), (d0, D, PALE)): + ax.broken_barh([(t0, dur)], (y, LANE), facecolors=col, + edgecolor=BG, linewidth=0.8, zorder=3) + for st in range(NS): + ax.text(-6, TOP - st * 0.72 + LANE / 2, f"stream {st}", color=MUTED, + fontsize=9.5, ha="right", va="center") + + def union(iv): + pts = sorted(iv) + out = [list(pts[0])] + for a, b in pts[1:]: + if a <= out[-1][1]: + out[-1][1] = max(out[-1][1], b) + else: + out.append([a, b]) + return out + + lanes = [("H2D engine busy", [(h0, h0 + H) for _, h0, _, _ in frames], AMBER), + ("kernels engine busy", [(k0, k0 + K) for _, _, k0, _ in frames], ACCENT)] + ys = 1.62 + for name, iv, col in lanes: + for a, b in union(iv): + ax.broken_barh([(a, b - a)], (ys, 0.34), facecolors=col, zorder=3) + ax.text(-6, ys + 0.17, name.split(" ")[0], color=col, fontsize=9.5, + ha="right", va="center", fontweight="bold") + ys -= 0.62 + + ax.text(frames[-1][3] + D + 8, TOP - 0.72, "one copy engine\nper direction,\n" + "so the H2D bars\nqueue", color=MUTED, fontsize=9.5, va="center", + linespacing=1.4) + ax.text(frames[-1][2] + K + 8, 1.45, + "UNION — what s4 reports.\nThe kernels overlap, so it is\n" + "shorter than their sum.", color=TEXT2, fontsize=9.5, va="center", + linespacing=1.5) + ax.plot([-2, frames[-1][3] + D + 2], [2.42, 2.42], color=RULE, lw=0.9, zorder=1) + + ax.set_xlim(-74, frames[-1][3] + D + 86) + ax.set_ylim(0.75, 6.05) + ax.set_xticks([]); ax.set_yticks([]) + bare(ax, keep=()) + ax.set_title("s4 · the shipped pipeline, four streams · schematic, 9×9 shape", + color=MUTED, fontsize=9.5, loc="left", pad=10) + + # ---- right: the measured occupancies, and the two estimates of the floor + vals = [("H2D", 20.77, AMBER), ("kernel", 32.66, ACCENT), ("D2H", 25.25, PALE)] + xs = np.arange(3) + ax2.bar(xs, [v for _, v, _ in vals], width=0.56, + color=[c for _, _, c in vals], zorder=3) + for x, (_, v, _) in zip(xs, vals): + ax2.text(x, v + 1.0, f"{v:.2f}", ha="center", color=TEXT2, fontsize=10) + ax2.axhline(32.66, color=MUTED, lw=1.0, ls="--", zorder=4) + ax2.text(2.42, 33.3, "engine max\nprofiled 32.66", color=MUTED, fontsize=9.5, + ha="left", va="bottom", linespacing=1.4) + ax2.axhline(30.01, color=GREEN, lw=1.8, zorder=5) + ax2.text(2.42, 24.4, "best sustained\nFLOOR 30.01 µs/frame\n= 33 323 FPS", + color=GREEN, fontsize=9.5, ha="left", va="center", linespacing=1.4) + ax2.set_xticks(xs) + ax2.set_xticklabels([n for n, _, _ in vals], color=TEXT2, fontsize=10) + ax2.set_xlim(-0.55, 4.35) + ax2.set_ylim(0, 40) + ax2.set_yticks([]) + bare(ax2, keep=("bottom",)) + ax2.set_title("busy µs / frame · 9×9 f64 · measured", + color=PALE, fontsize=9.5, loc="left", pad=6) + fig.subplots_adjust(left=0.085, right=0.99, top=0.86, bottom=0.09, wspace=0.20) + save(fig, "fig_measure") + + +fig_measure() + + + + +legibility_report() diff --git a/docs/figures/fig_arc.png b/docs/figures/fig_arc.png index 7bbe2341..85338339 100644 Binary files a/docs/figures/fig_arc.png and b/docs/figures/fig_arc.png differ diff --git a/docs/figures/fig_arc_9x9.png b/docs/figures/fig_arc_9x9.png index 830dda02..5ebb94ba 100644 Binary files a/docs/figures/fig_arc_9x9.png and b/docs/figures/fig_arc_9x9.png differ diff --git a/docs/figures/fig_bottleneck.png b/docs/figures/fig_bottleneck.png index 2b355d93..0470265f 100644 Binary files a/docs/figures/fig_bottleneck.png and b/docs/figures/fig_bottleneck.png differ diff --git a/docs/figures/fig_cancellation.png b/docs/figures/fig_cancellation.png index a3667984..e00117a2 100644 Binary files a/docs/figures/fig_cancellation.png and b/docs/figures/fig_cancellation.png differ diff --git a/docs/figures/fig_correctness.png b/docs/figures/fig_correctness.png index a7a4d99c..0e86a1cc 100644 Binary files a/docs/figures/fig_correctness.png and b/docs/figures/fig_correctness.png differ diff --git a/docs/figures/fig_f32_absolute.png b/docs/figures/fig_f32_absolute.png index 3cc9d553..4a710079 100644 Binary files a/docs/figures/fig_f32_absolute.png and b/docs/figures/fig_f32_absolute.png differ diff --git a/docs/figures/fig_f32_kernel.png b/docs/figures/fig_f32_kernel.png index 142cdff1..3fe07343 100644 Binary files a/docs/figures/fig_f32_kernel.png and b/docs/figures/fig_f32_kernel.png differ diff --git a/docs/figures/fig_first_run.png b/docs/figures/fig_first_run.png index 2576595c..0ad3017e 100644 Binary files a/docs/figures/fig_first_run.png and b/docs/figures/fig_first_run.png differ diff --git a/docs/figures/fig_graphs.png b/docs/figures/fig_graphs.png index 4144eebd..c5a9adc5 100644 Binary files a/docs/figures/fig_graphs.png and b/docs/figures/fig_graphs.png differ diff --git a/docs/figures/fig_measure.png b/docs/figures/fig_measure.png new file mode 100644 index 00000000..ec2480c1 Binary files /dev/null and b/docs/figures/fig_measure.png differ diff --git a/docs/figures/fig_mismatch147.png b/docs/figures/fig_mismatch147.png index 28eef876..50e526d5 100644 Binary files a/docs/figures/fig_mismatch147.png and b/docs/figures/fig_mismatch147.png differ diff --git a/docs/figures/fig_occupancy.png b/docs/figures/fig_occupancy.png index 38d057cc..af330310 100644 Binary files a/docs/figures/fig_occupancy.png and b/docs/figures/fig_occupancy.png differ diff --git a/docs/figures/fig_opt1_timeline.png b/docs/figures/fig_opt1_timeline.png index e995efa2..0c31d995 100644 Binary files a/docs/figures/fig_opt1_timeline.png and b/docs/figures/fig_opt1_timeline.png differ diff --git a/docs/figures/fig_opt2_timeline.png b/docs/figures/fig_opt2_timeline.png index 09b5c30f..c5e14d4e 100644 Binary files a/docs/figures/fig_opt2_timeline.png and b/docs/figures/fig_opt2_timeline.png differ diff --git a/docs/figures/fig_overhead.png b/docs/figures/fig_overhead.png index 7fec7ec3..2f1bdd17 100644 Binary files a/docs/figures/fig_overhead.png and b/docs/figures/fig_overhead.png differ diff --git a/docs/figures/fig_overlap.png b/docs/figures/fig_overlap.png index 39b58938..604cdbe9 100644 Binary files a/docs/figures/fig_overlap.png and b/docs/figures/fig_overlap.png differ diff --git a/docs/figures/fig_pedtiming.png b/docs/figures/fig_pedtiming.png index 91ed8e76..ebb178c3 100644 Binary files a/docs/figures/fig_pedtiming.png and b/docs/figures/fig_pedtiming.png differ diff --git a/docs/figures/fig_pinning.png b/docs/figures/fig_pinning.png index 4fc6b01a..df96c5d1 100644 Binary files a/docs/figures/fig_pinning.png and b/docs/figures/fig_pinning.png differ diff --git a/docs/figures/fig_regpressure.png b/docs/figures/fig_regpressure.png index 4a803e9e..4c8a1f64 100644 Binary files a/docs/figures/fig_regpressure.png and b/docs/figures/fig_regpressure.png differ diff --git a/docs/figures/fig_resultpath.png b/docs/figures/fig_resultpath.png index fe62fd54..1d696ac5 100644 Binary files a/docs/figures/fig_resultpath.png and b/docs/figures/fig_resultpath.png differ diff --git a/docs/figures/fig_streams.png b/docs/figures/fig_streams.png index c6eb1e32..0dbd369a 100644 Binary files a/docs/figures/fig_streams.png and b/docs/figures/fig_streams.png differ diff --git a/docs/figures/fig_varfloor.png b/docs/figures/fig_varfloor.png new file mode 100644 index 00000000..b61b33ba Binary files /dev/null and b/docs/figures/fig_varfloor.png differ diff --git a/docs/figures/fig_variance_rewrite.png b/docs/figures/fig_variance_rewrite.png index 9866dc62..7f9ef1a7 100644 Binary files a/docs/figures/fig_variance_rewrite.png and b/docs/figures/fig_variance_rewrite.png differ diff --git a/docs/figures/img_Intel_comet_lake.jpg b/docs/figures/img_Intel_comet_lake.jpg new file mode 100644 index 00000000..d19c21e8 Binary files /dev/null and b/docs/figures/img_Intel_comet_lake.jpg differ diff --git a/docs/figures/img_cpu_die.png b/docs/figures/img_cpu_die.png new file mode 100644 index 00000000..c046f1c5 Binary files /dev/null and b/docs/figures/img_cpu_die.png differ