diff --git a/docs/cf_cuda_performance.pdf b/docs/cf_cuda_performance.pdf index e0c32a24..97b5d2ec 100644 Binary files a/docs/cf_cuda_performance.pdf and b/docs/cf_cuda_performance.pdf differ diff --git a/docs/cf_cuda_performance.pptx b/docs/cf_cuda_performance.pptx index b6247ff9..3acef000 100644 Binary files a/docs/cf_cuda_performance.pptx and b/docs/cf_cuda_performance.pptx differ diff --git a/docs/deck/CHANGELOG_2026-08-27.md b/docs/deck/CHANGELOG_2026-08-27.md new file mode 100644 index 00000000..3c5a33be --- /dev/null +++ b/docs/deck/CHANGELOG_2026-08-27.md @@ -0,0 +1,161 @@ +# 2026-08-27 — 15 pt body, wider margins, and a layout checker that works + +Session log for `docs/cf_cuda_performance.pptx` (via `docs/deck/build_performance_deck.py` ++ `make_figs.py` + `make_figs_kernel.py`). Deck rebuilt at **36 numbered slides**, +56 pages. The rendered PDF now audits **clean**: zero overlapping text, zero text +past the footer. + +--- + +## 1. The content box widened to `[0.50, 12.90]` + +It was `[0.70, 12.70]`. Narrowing the two outer margins buys 0.4 in of line +length — about 3 % more characters per line at every size, and on the widest +slides a whole line back. + +It was applied as an **affine map** of x and w over every module-level layout +call, not by hand: nothing changes proportion, columns keep their relative gap, +and only the two margins give up space. `M`, `COL`, `RAIL_X` and `RAIL_W` moved +with it, along with the geometry inside `chrome()`, `annex_chrome()`, `rail()` +and `section()`, which the rewriter deliberately did not touch (inside a helper +every x is an offset from a caller-supplied origin). + +Three module-level loops carry a **stride** rather than an offset — the hero +stat row, the s1/s4/floor cards, the closing four-card grid — and a stride is +not an x. Those were the only sites that needed a hand edit, and skipping them +would have overlapped the cards rather than moved them. + +`deckgate.placements()` holds its own copy of the four tokens so it can evaluate +placement expressions. **It was still holding the old values**, so for one build +the legibility gate was checking widths no figure was placed at. + +## 2. The interstitials read at body size + +`section()`'s thesis went 13.5 → **16 pt** and its "coming up" list 12 → **15**. +The list is the reason the divider exists, so it is set at body size, not +caption size. + +That costs height, so each entry is now **measured** rather than given a fixed +step: a two-line entry gets two lines of room instead of sitting on the one +below it. The pitch is `15 × 1.45 / 72`, not the `1.22` the nominal line spacing +suggests — see §6. + +## 3. Slide-by-slide, what came off + +Roughly 30 specific cuts. The pattern in all of them: the numbers and the +remarks stay, the sentence explaining why they are true moves to `notes()`. + +- **3** — spectra caption cut to one clause; the closing line loses "not lit pixels". +- **4** — second bullet and the pedestal-timing rail note removed; the timing + argument (and its link to A7) is now a speaker note. +- **5** — the Comet Lake die caption, which the picture already said. +- **6** — "A stalled warp is simply replaced, never reordered." +- **7** — right column shortened; the last callout had been sitting past the footer. +- **8** — the shared-memory-size paragraph → notes; the tile diagram grew into it. +- **9** — the register callout says the two cases and stops; rail note removed. +- **10, 12, 16, 18, 20, 21, 23, 24, 26, 28, 29, 30, 31, 35, 36** — one box, one + note or one repetitive sentence each. Slide 12 lost its second callout entirely + (the per-engine numbers moved to slide 13, where they belong to a measured run + rather than to a diagram) and its DMA footnote. +- **21** — the s1 / s4 / floor cards are one-liners now, and the picture beneath + them is 11.6 in wide instead of 10.6. +- **35** — the guidance column halved; the header row had been *under* the warning + callout, not below it. + +## 4. Two code panels had been silently deleted + +Slides 33 and 34 each contained a bare list literal where a `code(...)` call +used to be — the `code(s, x, y, w,` prefix was gone and the list was evaluated +and discarded. Slide 33 had lost the frame-147 walkthrough, which **is** the +slide's argument, and 34 had lost "THE RECOMMENDED PATTERN", which is the API +the slide exists to teach. Both restored; slide 33's now also shows cuda's +`accept` branch, which had never been on the slide next to frozen's `reject`. + +A third instance, on slide 15, had lost the opt2/opt3 barrier comparison. + +## 5. The figures + +Every figure was re-laid where text met text. The ones worth naming: + +- **`fig_pinning`** was the worst: box labels wider than their boxes, arrow + labels in a gutter narrower than the words in it. Boxes widened, arrow labels + moved above the row. +- **`fig_tile`** — `set_aspect("equal")` shrinks the axes box to honour the + aspect, so the data-space room reserved under the tile for its key was being + squeezed and the key clipped. The key is drawn in **figure** coordinates now. + The two panels also sit 0.06 apart instead of 0.135, and the canvas is 6.9 in + rather than 7.6 — a narrower canvas at the same placement width is a *bigger* + projected figure. +- **`fig_resultpath`** — the verdict is the second line of the axes **title**. + As free-floating text it had nowhere to go: above the bars it met the value + labels, below the axis it met two-line tick labels, beside them it met the + gain. Bars also moved 1.35 apart, because the tick labels are ~0.95 x-units + wide each. +- **`fig_streams`** — the barrier label was in the axes title's row. An axes + title is anchored to the axes *box*, so it moves with the limits and not with + the data; the fix is headroom in y. "time →" came inside the axes, because as + an xlabel it sat on the sentence already written there. +- **`fig_arc`** — band labels shortened to `ACT I / II / III`. The gloss was + wider than the one-bar ACT III band, so ACT III's label was on ACT II's. +- **`fig_pagefault`** — the key moved below the two rows; the right-hand panel + owns that half of the figure and "first touch → fault" is 18 x-units wide, + which was exactly the clearance left between them. +- **`fig_pedtiming`** — "every decision uses the frame-start snapshot" was at + `y + 0.18`, sitting on the 2.2 pt trace it labels. +- **`fig_cancellation`**, **`fig_f32_kernel`**, **`fig_gpu_model`**, + **`fig_regpressure`**, **`fig_occupancy`**, **`fig_overlap_9x9`**, + **`fig_first_run`**, **`fig_measure`**, **`fig_mismatch147`** — same class of + fix: headroom, or a label moved off the mark it was naming. + +## 6. Line height is not what the nominal spacing says + +Two constants, both wrong in the direction that clips: + +- `code()` sets `lh = 0.0189 * size` in/line for Consolas. The old 0.0174 + under-counted by 8.6 %. +- LibreOffice sets a 15 pt UI line at about **0.31 in**, not the 0.26 that + `line=1.25` implies. Every stack computed against 0.26 came out ~20 % short, + which is one line on a five-line slide. + +## 7. `deckgate.py`, and the kernel figures joining the gate + +The projection floor, the placement table and the report moved into their own +module. There are **two** figure generators and only `make_figs.py` was ever +checked; `make_figs_kernel.py`'s three figures were consequently set at 7–8 pt +and nobody found out until they were projected. All font sizes there went up +×1.30 and the module now imports `deckgate`. + +The floor itself is **10.0 pt**, two thirds of the 15 pt body. + +A string may opt out with `gid="texture"` — for marks nobody is asked to READ, +where the pattern is the message and the digits are shading. Two figures use it: +`fig_mismatch147`'s per-pixel ADU values (a cell is ~0.18 in wide and a value can +be three digits, so nothing that clears the floor can also fit) and `fig_frame`'s +3×3 zoom. It is set per `Text` object, never per figure. + +## 8. `audit_layout.py` now finds what it was built to find + +Three changes, each of which had been hiding a real defect: + +1. **Lines split at a column gutter.** Grouping words by baseline alone merges a + rail row and a body bullet at the same height into one full-width "line". + Two such lines overlap in x by construction, so the audit was reporting a + collision on every two-column slide — 11 pages of noise that buried the four + real ones. +2. **Overlap threshold 25 % → 10 %.** At 25 % a bullet ran a whole word under a + code panel title without registering. It did, on slide 17. +3. **Two narrow exemptions**, so that a clean deck reports *clean* and the tool + stays worth running: page 1 is the PSI template's own title slide, and a + one-character "line" beside a real one is a superscript marker — which is what + a footnote mark is. + +Result: **0 collisions, 0 overruns**, from 34 pages with collisions at the start +of the 15 pt migration. + +## 9. Wrap-join repairs + +Eleven more glued words from earlier bulk rewrites — `over 128SMs`, +`thecount`, `opt1 andopt2`, `the hostadds`, `there,invisible`, `make thecolumns`, +`costsnothing`, `whosetrue variance`, `mmapthreshold`, `at tol= 0` — where two +adjacent Python string literals were joined with nothing between them. Python +concatenates them silently; only a render shows it. diff --git a/docs/deck/README.md b/docs/deck/README.md index f8951a78..d87ce982 100644 --- a/docs/deck/README.md +++ b/docs/deck/README.md @@ -12,6 +12,7 @@ are counted. python docs/deck/make_figs.py # figures -> docs/figures/*.png python docs/deck/make_figs_kernel.py # 3 more (fig_frame, fig_occupancy, fig_tile) python docs/deck/build_performance_deck.py +python docs/deck/audit_layout.py # must print "clean" # the shareable copy, regenerated from the .pptx after every rebuild libreoffice --headless --convert-to pdf --outdir /tmp/pdfout \ @@ -41,39 +42,87 @@ interpreter with all four is `/home/ferjao_k/.conda/envs/py/bin/python` — the the title slide, and every other slide of it is deleted at build time. Do not edit the generated `.pptx` by hand — it is overwritten on every build. Edit the script. +## The type scale + +One place, `build_performance_deck.py`, because 186 scattered `size=` arguments cannot +be reasoned about: + +| token | pt | what it is for | +|---|---|---| +| `PT_BODY` | 15 | bullets: what the slide is claiming | +| `PT_LEAD` | 13 | callouts: the sentence to remember, already boxed and bold | +| `PT_RAIL` | 13 | rail row values | +| `PT_TABLE` | 11.5 | table cells | +| `PT_META` | 11 | captions and rail notes: provenance, deliberately quieter | +| `PT_CODE` | 10.5 | code panels — read by token, not word by word | +| `PT_LABEL` | 9.5 | small-caps labels on rails, tables, statstrips | + +**Raising `PT_BODY` costs text.** 15 pt holds roughly 45 % of the characters 10.5 pt did +in the same box, because area scales with the square of point size. The scale is a +budget, not a preference: a slide that will not fit loses words, never type size. + +Two line-height constants are calibrated against a LibreOffice render, not against the +nominal spacing, and both were wrong in the direction that clips: + +- `code()` sets `lh = 0.0189 * size` in/line for Consolas. The old 0.0174 under-counted + by 8.6 %, which is invisible on a short panel and eats the last line or two of a long + one. +- LibreOffice sets a 15 pt UI line at about **0.31 in**, not the 0.26 that `line=1.25` + implies. Budget 0.31 per body line when deciding what a slide can hold. + ## Layout guarantees, and how they are enforced -Two invariants are checked mechanically, because both fail silently otherwise. +Three invariants are checked mechanically, because all three fail silently otherwise. -**Nothing renders below 9 pt on the projected slide.** A figure's on-screen type size +**Nothing renders below 10 pt on the projected slide.** A figure's on-screen type size is `raw_pt × (placement_width / figure_width)`, and neither factor is visible at the -point where the font size is written. `make_figs.py` closes that loop: `_placements()` +point where the font size is written. `deckgate.py` closes that loop: `placements()` parses the placement width of every figure **out of the deck script itself**, so the -gate cannot drift from the layout it checks. Every run ends with either +gate cannot drift from the layout it checks. Both generators import it — `make_figs.py` +and `make_figs_kernel.py` — and each run ends with either ``` -legibility: every string in every figure renders at >= 9.0 pt on the slide. +legibility: every string in every figure renders at >= 10.0 pt on the slide. ``` -or a list of offenders. Fix them; do not raise the floor. 9 pt on a 13.33 × 7.5 in -slide is about 1/60 of slide height, which is the conventional bound for readable -supporting detail at 6–7 m. +or a list of offenders. Fix them; do not raise the floor. 10 pt is two thirds of the +15 pt body, which is the usual lower bound for supporting type at 6–7 m. + +The kernel figures were outside this gate for a long time and were set at 7–8 pt as a +result. If you add a third generator, import `deckgate` from it on day one. + +A string may opt out with `gid="texture"`, and only for marks nobody is asked to READ — +a value printed into every cell of a pixel map, where the pattern is the message and +the digits are shading. Two figures use it: `fig_mismatch147`'s per-pixel ADU values +and `fig_frame`'s 3×3 zoom. Set it on the `Text` object, never on a figure. Note the feedback trap: `savefig(bbox_inches="tight")` grows the saved canvas to fit a long in-figure caption, which shrinks the placement scale, which shrinks the caption. Raising the font size can make text *smaller*. Shorten the string or re-lay the axes. +The corollary is useful: a NARROWER `figsize` at the same placement width is a bigger +projected figure and bigger projected type. -**No text runs past the footer line.** Convert and check: +**No text overlaps other text, and none runs past the footer.** Both are checked by +`audit_layout.py`, which reads the rendered PDF's word boxes: ```bash libreoffice --headless --convert-to pdf --outdir /tmp/deck docs/cf_cuda_performance.pptx -python scratch/overflow.py /tmp/deck/cf_cuda_performance.pdf +python docs/deck/audit_layout.py /tmp/deck/cf_cuda_performance.pdf # -> "clean" ``` -Only page 1 may be flagged — that is the PSI template's own title slide. The same -script counts unrendered `**` markup, which is the usual symptom of putting markup in a -helper that does not parse it: `bullets`, `callout`, `table` and `code` understand -`**bold**`; `caption` does not, and nothing understands backticks or `*italics*`. +Three things about it are load-bearing: + +- It groups words into lines by baseline and then **splits each line at a column + gutter**. Without that split a rail row and a body bullet at the same height merge + into one full-width "line", and every two-column slide in the deck reports a + collision. +- It asserts on the parsed word count. A broken bbox parser yields empty boxes and + therefore a *clean* report, which is the worst failure mode a checker can have. +- Two exemptions are deliberate and narrow: page 1 is the PSI template's own title + slide, and a one-character "line" beside a real one is a superscript marker. + +Overlap has to cover 10 % of the narrower line to register. At 25 % a bullet could run a +whole word under a code-panel title without being reported, which it did. ## Numbering @@ -89,6 +138,13 @@ corner and the same way `annex_chrome` draws `A`. The badge names the *rung*, slide, so opt3, opt5 and opt7 repeat theirs across two slides each. When a slide has a badge its eyebrow must not also say "optN" — that reads twice. +**The content box is `[0.50, 12.90]` on a 13.333 in slide.** It used to be +`[0.70, 12.70]`; narrowing the two outer margins bought 0.4 in of line length, which is +about 3 % more characters per line at every size. The tokens that carry it are `M`, +`COL`, `RAIL_X` and `RAIL_W`, and **`deckgate.placements()` holds copies of all four** +so it can evaluate placement expressions. Change one and change the other, or the +legibility gate starts checking widths no figure is placed at. + **Nothing on a slide may point outside the deck.** No notebook names, no `python/tests/…`, no result directories, no "see the write-up" — the slides get shared on their own, and a pointer to something the reader does not have is worse than no @@ -128,8 +184,10 @@ report: opt5 at 9×9, whose per-frame allocation never lets the fault count conv | file | role | |---|---| | `build_performance_deck.py` | the deck: tokens, helpers, every slide | -| `make_figs.py` | most figures, plus the legibility gate | +| `make_figs.py` | most figures | | `make_figs_kernel.py` | `fig_frame`, `fig_occupancy`, `fig_tile` | +| `deckgate.py` | the projection floor and the placement table, shared by both generators | +| `audit_layout.py` | reads the rendered PDF: overlapping text, and text past the footer | | `QA.md` | questions the room asks, with the answers and where they are settled | | `frame147.json`, `validation_tiers.json` | measured data two figures read | | `branch_site.json` | the A7 site dump; written by `python/tests/branch_site_dump.py` | diff --git a/docs/deck/audit_layout.py b/docs/deck/audit_layout.py new file mode 100644 index 00000000..b5f59161 --- /dev/null +++ b/docs/deck/audit_layout.py @@ -0,0 +1,119 @@ +"""Does any text on a slide physically overlap any other text? + +The overflow checker only sees text crossing the footer line. It cannot see a +bullet running into the code panel underneath it, which is exactly the failure a +type-size change produces. This reads the RENDERED pdf and reports every pair of +text lines whose bounding boxes actually intersect. + + python docs/deck/audit_layout.py docs/cf_cuda_performance.pdf + +No column heuristics: two lines are a collision only if their boxes overlap in +BOTH axes, so side-by-side columns and consecutive baselines never register. +""" +import re +import subprocess +import sys +from collections import defaultdict + +PDF = sys.argv[1] if len(sys.argv) > 1 else "docs/cf_cuda_performance.pdf" +FOOTER = 525.5 # the progress track; nothing may sit below it +LINE_TOL = 2.5 # words whose baselines are this close share a line +X_FRAC = 0.10 # overlap must cover this much of the narrower line + # (0.25 let a bullet run a whole word under a code + # panel title without registering) +Y_MIN = 2.0 # ... and this many points vertically +GUTTER = 24.0 # white space this wide separates two columns + +ATTR = re.compile(r'(\w+)="([\d.eE+-]+)"') + +xml = subprocess.run(["pdftotext", "-bbox", PDF, "-"], + capture_output=True, text=True).stdout + +pages, cur = [], None +for raw in xml.split("\n"): + if "")[0])) + if {"xMin", "yMin", "xMax", "yMax"} <= g.keys(): + cur.append((float(g["xMin"]), float(g["yMin"]), + float(g["xMax"]), float(g["yMax"]), + raw.split(">", 1)[1].split("")[0])) + + +def lines_of(words): + """Group words into lines, and SPLIT a line at a column gutter. + + Grouping by baseline alone merges a rail row and a body bullet that happen + to sit at the same height into one full-width "line". Two such lines then + overlap in x by construction, and the audit reports a collision on every + two-column slide in the deck -- which is most of them. A run of words + separated by more than a third of an inch of white space is a different + column, not the same line. + """ + rows = defaultdict(list) + for w in words: + rows[round(w[3] / LINE_TOL)].append(w) + out = [] + for k in sorted(rows): + run = [] + for w in sorted(rows[k]): + if run and w[0] - run[-1][2] > GUTTER: + out.append(_span(run)) + run = [] + run.append(w) + if run: + out.append(_span(run)) + return out + + +def _span(ws): + return (min(w[0] for w in ws), min(w[1] for w in ws), + max(w[2] for w in ws), max(w[3] for w in ws), + " ".join(w[4] for w in ws)) + + +# Page 1 is the PSI template's own title slide. Its layout is not ours to fix -- +# the deck imports it whole -- and it puts the date below our footer line. +TEMPLATE_PAGE = 1 + +overlaps, footer = [], [] +for pno, words in enumerate(pages, 1): + ls = lines_of(words) + for i, a in enumerate(ls): + if a[3] > FOOTER and pno != TEMPLATE_PAGE: + footer.append((pno, a[3], a[4])) + for b in ls[i + 1:]: + # A one-character "line" beside a real one is a superscript marker + # sharing the baseline it annotates. That is what a footnote mark IS. + if len(a[4]) == 1 or len(b[4]) == 1: + continue + ox = min(a[2], b[2]) - max(a[0], b[0]) + oy = min(a[3], b[3]) - max(a[1], b[1]) + if oy < Y_MIN or ox <= 0: + continue + if ox >= X_FRAC * min(a[2] - a[0], b[2] - b[0]): + overlaps.append((pno, ox, oy, a[4][:44], b[4][:44])) + +nw = sum(len(p_) for p_ in pages) +assert nw > 500, f"parsed only {nw} words -- the bbox parser is broken, not the deck" +print(f"pages: {len(pages)} words parsed: {nw}\n") +print(f"COLLIDING text: {len(overlaps)}") +seen = set() +for pno, ox, oy, a, b in overlaps: + if pno in seen and len([o for o in overlaps if o[0] == pno]) > 3: + continue + seen.add(pno) + print(f" p{pno:>3} x{ox:5.1f} y{oy:5.1f} {a!r}\n {b!r}") +n_pages = len({o[0] for o in overlaps}) +print(f" ({n_pages} pages affected)") + +print(f"\nBELOW THE FOOTER: {len(footer)}") +for pno, y, txt in footer: + print(f" p{pno:>3} y {y:6.1f} {txt[:70]!r}") + +print("\nclean" if not (overlaps or footer) else "\nFIX THE ABOVE") +sys.exit(1 if (overlaps or footer) else 0) diff --git a/docs/deck/build_performance_deck.py b/docs/deck/build_performance_deck.py index 499ff1a7..3f679d92 100644 --- a/docs/deck/build_performance_deck.py +++ b/docs/deck/build_performance_deck.py @@ -69,9 +69,9 @@ RED = RGBColor(0xE2, 0x54, 0x54) # pointing only, never a data colou UI, MONO = "Segoe UI", "Consolas" W, H = 13.333, 7.5 -M = 0.7 # left margin -COL = 7.9 # left column width -RAIL_X, RAIL_W = 9.2, 3.5 # right rail +M = 0.50 # left margin +COL = 8.16 # left column width +RAIL_X, RAIL_W = 9.28, 3.62 # right rail A = "{http://schemas.openxmlformats.org/drawingml/2006/main}" P = "{http://schemas.openxmlformats.org/presentationml/2006/main}" @@ -166,6 +166,22 @@ def para(tf, first=False, space_after=0, space_before=0, line=None, align=None): # floor applied by hand is a floor that one new caption silently drops through. MIN_PT = 9.0 +# ------------------------------------------------------------ the type scale +# One place, because 186 scattered `size=` arguments cannot be reasoned about. +# Sized for a 13.33 x 7.5 in slide read at 6-7 m: BODY is the argument, and +# everything else is explicitly subordinate to it. Raising BODY costs text -- +# 15 pt holds roughly 45 % of the characters 10.5 pt did in the same box -- so +# the scale is a budget, not a preference. Figures are gated separately, in +# make_figs.py, and their floor moves with BODY: type inside a plot that is much +# smaller than the body text reads as an afterthought from the back of a room. +PT_BODY = 15 # bullets: what the slide is claiming +PT_LEAD = 13 # callouts: the sentence to remember, already boxed and bold +PT_TABLE = 11.5 # table cells +PT_RAIL = 13 # rail row values +PT_META = 11 # captions and rail notes: provenance, deliberately quieter +PT_CODE = 10.5 # code panels -- read by token, not word by word +PT_LABEL = 9.5 # small-caps labels on rails, tables, statstrips + def run(p, text, size=11, color=TEXT2, font=UI, bold=False, italic=False, spc=None, sup=False): @@ -195,20 +211,20 @@ def chrome(s, idx, eyebrow, title, title_size=27, opt=None): if opt: tf = tb(s, M, 0.44, 1.2, 0.34) run(para(tf, True), f"OPT{opt}", 15, ACCENT, bold=True) - tf = tb(s, 1.62, 0.50, 9.88, 0.32) + tf = tb(s, 1.45, 0.50, 10.25, 0.32) else: rect(s, M, 0.60, 0.35, 0.035, ACCENT) - tf = tb(s, 1.17, 0.50, 10.33, 0.32) + tf = tb(s, 1.00, 0.50, 10.70, 0.32) run(para(tf, True), eyebrow.upper(), 9, MUTED, bold=True, spc=1.6) - tf = tb(s, M, 0.86, 11.9, 1.0) + tf = tb(s, M, 0.86, 12.29, 1.0) run(para(tf, True, line=1.05), title, title_size, PALE, bold=True) - span, n = 11.0, N_SLIDES + span, n = 11.37, N_SLIDES pitch = span / n; wseg = pitch * 0.90 for i in range(n): rect(s, M + i * pitch, 7.28, wseg, 0.045, ACCENT if i <= idx - 1 else RULE) - tf = tb(s, 12.0, 7.14, 0.9, 0.3) + tf = tb(s, 12.29, 7.14, 0.9, 0.3) run(para(tf, True, align=PP_ALIGN.RIGHT), f"{idx} / {n}", 8.5, MUTED) @@ -225,21 +241,21 @@ def annex_chrome(s, grp, eyebrow, title, part=None, nparts=None, title_size=27): """ tf = tb(s, M, 0.44, 1.0, 0.34) run(para(tf, True), f"A{grp}", 15, AMBER, bold=True) - tf = tb(s, 1.42, 0.50, 10.1, 0.32) + tf = tb(s, 1.25, 0.50, 10.45, 0.32) tag = f"ANNEX · {eyebrow}" + (f" · {part} of {nparts}" if nparts else "") run(para(tf, True), tag.upper(), 9, MUTED, bold=True, spc=1.6) - tf = tb(s, M, 0.86, 11.9, 1.0) + tf = tb(s, M, 0.86, 12.29, 1.0) run(para(tf, True, line=1.05), title, title_size, PALE, bold=True) - pitch = 11.0 / N_ANNEX + pitch = 11.37 / N_ANNEX for i in range(N_ANNEX): rect(s, M + i * pitch, 7.28, pitch * 0.90, 0.045, AMBER if i <= grp - 1 else RULE) - tf = tb(s, 11.7, 7.14, 1.2, 0.3) + tf = tb(s, 11.99, 7.14, 1.2, 0.3) foot = f"A{grp}" + (f" · {part}/{nparts}" if nparts else "") run(para(tf, True, align=PP_ALIGN.RIGHT), foot, 8.5, MUTED) -def table(s, x, y, w, header, rows, colw, size=9.5, rowh=0.62): +def table(s, x, y, w, header, rows, colw, size=PT_TABLE, rowh=0.74): """Minimal header + zebra table. colw are fractions of w.""" xs, acc = [], 0.0 for c in colw: @@ -248,7 +264,7 @@ def table(s, x, y, w, header, rows, colw, size=9.5, rowh=0.62): rect(s, x, y, w, 0.34, PANEL) for cx, h in zip(xs, header): tf = tb(s, cx + 0.16, y + 0.08, w, 0.24) - run(para(tf, True), _up(h), 8, MUTED, bold=True, spc=1.2) + run(para(tf, True), _up(h), PT_LABEL, MUTED, bold=True, spc=1.2) yy = y + 0.38 for i, row in enumerate(rows): if i % 2 == 0: @@ -264,7 +280,7 @@ def table(s, x, y, w, header, rows, colw, size=9.5, rowh=0.62): return yy -def bullets(s, x, y, w, items, size=11, gap=7): +def bullets(s, x, y, w, items, size=PT_BODY, gap=9): tf = tb(s, x, y, w, 0.3) for i, it in enumerate(items): color, txt = (it if isinstance(it, tuple) else (TEXT2, it)) @@ -283,7 +299,7 @@ def bullets(s, x, y, w, items, size=11, gap=7): CODE_MARK = re.compile(r"«([^»]*)»|‹([^›]*)›") -def code(s, x, y, w, lines, size=9, title=None): +def code(s, x, y, w, lines, size=PT_CODE, title=None): # Line height has to follow the font size, and the constant has to match what # the RENDERER does, not what python-pptx assumes. Measured off a LibreOffice # render at 9 pt Consolas with line=1.12: 0.170 in per line, i.e. 0.0189 in @@ -300,7 +316,7 @@ def code(s, x, y, w, lines, size=9, title=None): ty = y + 0.12 if title: tf = tb(s, x + 0.18, ty, w - 0.36, 0.2) - run(para(tf, True), title, 9, CODEDIM, bold=True, spc=1.2) + run(para(tf, True), title, PT_LABEL, CODEDIM, 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): @@ -322,7 +338,7 @@ def code(s, x, y, w, lines, size=9, title=None): return h -def callout(s, x, y, w, text, h=0.78, color=ACCENT, size=10.5): +def callout(s, x, y, w, text, h=0.78, color=ACCENT, size=PT_LEAD): rect(s, x + 0.045, y, w - 0.045, h, PANEL) rect(s, x, y, 0.045, h, color) tf = tb(s, x + 0.28, y + 0.10, w - 0.5, h - 0.2, MSO_ANCHOR.MIDDLE) @@ -340,32 +356,32 @@ def _up(txt): def rail(s, items, y0=2.0, divider=True): if divider: - rect(s, 8.95, 2.0, 0.012, 4.55, RULE) + rect(s, 9.03, 2.0, 0.012, 4.55, RULE) y = y0 for it in items: kind = it[0] if kind == "label": tf = tb(s, RAIL_X, y, RAIL_W, 0.26) - run(para(tf, True), _up(it[1]), 8.5, MUTED, bold=True, spc=1.4) - y += 0.28 + run(para(tf, True), _up(it[1]), PT_LABEL, MUTED, bold=True, spc=1.4) + y += 0.32 elif kind == "stat": _, lab, val, col = it tf = tb(s, RAIL_X, y, RAIL_W, 0.24) - run(para(tf, True), _up(lab), 8.5, MUTED, spc=1.2) + run(para(tf, True), _up(lab), PT_LABEL, MUTED, spc=1.2) tf = tb(s, RAIL_X, y + 0.24, RAIL_W, 0.6) run(para(tf, True), val, 26, col, bold=True) - y += 0.98 + y += 1.02 elif kind == "row": _, lab, val, col = it tf = tb(s, RAIL_X, y, RAIL_W, 0.24) - run(para(tf, True), _up(lab), 8.5, MUTED, spc=1.2) - tf = tb(s, RAIL_X, y + 0.22, RAIL_W, 0.3) - run(para(tf, True), val, 13, col, bold=True) - y += 0.66 + run(para(tf, True), _up(lab), PT_LABEL, MUTED, spc=1.2) + tf = tb(s, RAIL_X, y + 0.24, RAIL_W, 0.32) + run(para(tf, True), val, PT_RAIL, col, bold=True) + y += 0.72 elif kind == "note": tf = tb(s, RAIL_X, y, RAIL_W, 0.9) - run(para(tf, True, line=1.25), it[1], 9, TEXT2) - y += 0.30 + 0.17 * (len(it[1]) // 42 + 1) + run(para(tf, True, line=1.25), it[1], PT_META, TEXT2) + y += 0.30 + 0.21 * (len(it[1]) // 34 + 1) elif kind == "gap": y += it[1] return y @@ -411,7 +427,7 @@ def notes(s, text): s.notes_slide.notes_text_frame.text = text -def caption(s, x, y, w, text, size=9): +def caption(s, x, y, w, text, size=PT_META): tf = tb(s, x, y, w, 0.3) run(para(tf, True, line=1.25), text, size, MUTED) @@ -426,8 +442,8 @@ def flow(s, x, y, w, steps, h=0.62): rect(s, bx, y, 0.03, h, ACCENT) tf = tb(s, bx + 0.20, y + 0.05, bw - 0.32, h - 0.10, MSO_ANCHOR.MIDDLE) p = para(tf, True, line=1.1) - run(p, f"{i + 1} ", 9, ACCENT, bold=True, font=MONO) - run(p, t, 9.5, PALE) + run(p, f"{i + 1} ", PT_LABEL, ACCENT, bold=True, font=MONO) + run(p, t, 12.5, PALE) if i < n - 1: tf = tb(s, bx + bw, y + 0.05, gap, h - 0.10, MSO_ANCHOR.MIDDLE) run(para(tf, True, align=PP_ALIGN.CENTER), "›", 15, MUTED, bold=True) @@ -439,10 +455,10 @@ def statstrip(s, x, y, w, items, h=0.80): for i, (lab, val) in enumerate(items): bx = x + i * (bw + gap) rect(s, bx, y, bw, h, PANEL) - tf = tb(s, bx + 0.20, y + 0.12, bw - 0.4, 0.22) - run(para(tf, True), lab.upper(), 8, MUTED, bold=True, spc=1.2) - tf = tb(s, bx + 0.20, y + 0.37, bw - 0.4, 0.34) - run(para(tf, True), val, 15, PALE, bold=True) + tf = tb(s, bx + 0.20, y + 0.10, bw - 0.4, 0.22) + run(para(tf, True), lab.upper(), PT_LABEL, MUTED, bold=True, spc=1.2) + tf = tb(s, bx + 0.20, y + 0.42, bw - 0.4, 0.34) + run(para(tf, True), val, 17, PALE, bold=True) # ------------------------------------------------------------ interstitial @@ -499,44 +515,52 @@ def section(kicker, title, thesis, items, rng, col=ACCENT, carry=None, s = new_slide() rect(s, 0, 0, 0.16, H, col) - tf = tb(s, M + 0.3, 1.52, 5.8, 0.3) - run(para(tf, True), kicker.upper(), 9.5, MUTED, bold=True, spc=1.8) - tf = tb(s, M + 0.3, 1.90, 6.0, 1.35) - run(para(tf, True, line=1.03), title, _fit(title, 6.0), PALE, bold=True) - rect(s, M + 0.3, 3.42, 1.5, 0.03, col) - tf = tb(s, M + 0.3, 3.70, 5.7, 1.6) - run(para(tf, True, line=1.4), thesis, 13.5, TEXT2) + tf = tb(s, M + 0.31, 1.52, 5.99, 0.3) + run(para(tf, True), kicker.upper(), 10.5, MUTED, bold=True, spc=1.8) + tf = tb(s, M + 0.31, 1.90, 6.2, 1.35) + run(para(tf, True, line=1.03), title, _fit(title, 6.2), PALE, bold=True) + rect(s, M + 0.31, 3.42, 1.55, 0.03, col) + tf = tb(s, M + 0.31, 3.70, 5.89, 1.6) + run(para(tf, True, line=1.35), thesis, 16, TEXT2) if carry: lab, val, sub = carry - rect(s, M + 0.3, 5.55, 5.7, 1.12, PANEL) - rect(s, M + 0.3, 5.55, 0.035, 1.12, col) - tf = tb(s, M + 0.60, 5.72, 5.2, 0.24) - run(para(tf, True), _up(lab), 8.5, MUTED, bold=True, spc=1.4) - tf = tb(s, M + 0.60, 5.94, 5.2, 0.4) + rect(s, M + 0.31, 5.55, 5.89, 1.12, PANEL) + rect(s, M + 0.31, 5.55, 0.035, 1.12, col) + tf = tb(s, M + 0.62, 5.70, 5.37, 0.24) + run(para(tf, True), _up(lab), 9.5, MUTED, bold=True, spc=1.4) + tf = tb(s, M + 0.62, 5.92, 5.37, 0.4) run(para(tf, True), val, 24, col, bold=True) - tf = tb(s, M + 0.60, 6.40, 5.2, 0.24) - run(para(tf, True), sub, 9.5, MUTED) + tf = tb(s, M + 0.62, 6.38, 5.37, 0.26) + run(para(tf, True), sub, 11, MUTED) - rect(s, 7.15, 1.95, 0.012, 4.4, RULE) - step = 0.42 if len(items) > 7 else 0.46 if len(items) > 5 else 0.54 - y = 1.95 + (4.4 - len(items) * step) / 2 - tf = tb(s, 7.45, y - 0.42, 5.0, 0.26) - run(para(tf, True), "COMING UP", 8.5, MUTED, bold=True, spc=1.6) - for num, txt in items: - tf = tb(s, 7.45, y, 0.7, 0.3) - run(para(tf, True), str(num), 12, col, bold=True, font=MONO) - tf = tb(s, 8.15, y, 4.5, 0.3) - p = para(tf, True) + # The list is the reason the divider exists, so it is set at body size, not + # at caption size. That costs height, so each entry is measured rather than + # given a fixed step: a two-line entry gets two lines of room instead of + # silently sitting on the one below it. + rect(s, 7.17, 1.95, 0.012, 4.4, RULE) + ITEM_PT, TXT_W = 15, 4.85 + cap = TXT_W / (ITEM_PT / 72) + nlines = [max(1, -(-_em(t.replace("**", "")) // cap)) for _, t in items] + step = ITEM_PT * 1.45 / 72 + total = sum(n * step for n in nlines) + 0.20 * (len(items) - 1) + y = 1.95 + (4.4 - total) / 2 + tf = tb(s, 7.47, y - 0.46, 5.17, 0.26) + run(para(tf, True), "COMING UP", 10, MUTED, bold=True, spc=1.6) + for (num, txt), nl in zip(items, nlines): + tf = tb(s, 7.47, y, 0.6, 0.3) + run(para(tf, True), str(num), ITEM_PT, col, bold=True, font=MONO) + tf = tb(s, 8.05, y, TXT_W, 0.3) + p = para(tf, True, line=1.22) for j, part in enumerate(txt.split("**")): if part: - run(p, part, 12, col if j % 2 else PALE, bold=bool(j % 2)) - y += step + run(p, part, ITEM_PT, col if j % 2 else PALE, bold=bool(j % 2)) + y += nl * step + 0.20 # The annex divider sits on the annex's own track: the main arc is finished # behind it, so lighting main-track segments would misreport where we are. n_track = N_ANNEX if annex else N_SLIDES - pitch = 11.0 / n_track + pitch = 11.37 / n_track for i in range(n_track): n = i + 1 # the section ahead in its own colour, what is already behind us dimmed, @@ -560,21 +584,21 @@ set_para_texts(by_name["CustomShape 3"], # =========================================================== 2 · HERO s = new_slide() -rect(s, 0, 0, 0.16, H, ACCENT) -tf = tb(s, M + 0.3, 0.85, 11, 0.3) +rect(s, 0, 0, 0.17, H, ACCENT) +tf = tb(s, M + 0.31, 0.85, 11.37, 0.3) run(para(tf, True), "AARE · HYBRID PIXEL DETECTORS · CUDA CLUSTERFINDER", 9.5, MUTED, bold=True, spc=1.8) -tf = tb(s, M + 0.3, 1.28, 11.4, 1.45) +tf = tb(s, M + 0.31, 1.28, 11.78, 1.45) _p = para(tf, True, line=1.02) run(_p, "The kernel was never the bottleneck —", 42, PALE, bold=True) run(_p, "*", 17, MUTED, bold=True, sup=110000) # discreet: the joke is the note run(_p, " feeding it was", 42, PALE, bold=True) -tf = tb(s, M + 0.3, 2.80, 11.4, 0.45) +tf = tb(s, M + 0.31, 2.80, 11.78, 0.45) run(para(tf, True, line=1.05), "One kernel, one thread per pixel, and seven steps to keep it fed", 22, ACCENT) -tf = tb(s, M + 0.3, 3.40, 11.4, 0.9) +tf = tb(s, M + 0.31, 3.40, 11.78, 0.9) run(para(tf, True, line=1.3), "The stencil was fast almost immediately: at 3×3 the kernel needs 5.5 µs per " "frame while getting that frame across PCIe costs 16.6 µs, or 13.2 " @@ -582,7 +606,7 @@ run(para(tf, True, line=1.3), "seven steps get data in, get results back, and measure honestly.", 12.5, TEXT2) -tf = tb(s, M + 0.3, 4.40, 11.4, 0.26) +tf = tb(s, M + 0.31, 4.40, 11.78, 0.26) run(para(tf, True), "WHERE THIS ENDS UP · THE SHIPPED f32 BUILD AFTER ALL SEVEN STEPS · " "ENGINE TIMES [f32 · s4] · RECONCILED IN A1", @@ -591,19 +615,19 @@ run(para(tf, True), stats = [("×9.1", "VS 24-THREAD CPU", ACCENT), ("61,312", "FRAMES / SECOND", PALE), ("16.3 µs", "PER FRAME, END TO END", PALE), ("6 / 23 M", "CLUSTER MISMATCH VS CPU TWIN", AMBER)] for i, (v, l, c) in enumerate(stats): - x = M + 0.3 + i * 2.85 - rect(s, x, 4.78, 0.035, 0.95, c) - tf = tb(s, x + 0.22, 4.78, 2.5, 0.55) + x = M + 0.31 + i * 2.95 + rect(s, x, 4.78, 0.04, 0.95, c) + tf = tb(s, x + 0.22, 4.78, 2.58, 0.55) run(para(tf, True), v, 30, c, bold=True) - tf = tb(s, x + 0.22, 5.41, 2.5, 0.3) + tf = tb(s, x + 0.22, 5.41, 2.58, 0.3) run(para(tf, True), l, 8.5, MUTED, spc=1.2) -rect(s, M + 0.3, 6.05, 11.0, 0.012, RULE) -tf = tb(s, M + 0.3, 6.25, 11.4, 0.6) +rect(s, M + 0.31, 6.05, 11.37, 0.012, RULE) +tf = tb(s, M + 0.31, 6.25, 11.78, 0.6) run(para(tf, True, line=1.35), "RTX 4090 (Ada, sm_89) · PCIe 4.0 ×16 · Mönch 400×400 uint16 · 3×3 clusters · " "100 000 frames · Cu fluorescence, MAX IV", 10, MUTED) -tf = tb(s, M + 0.3, 6.72, 11.4, 0.3) +tf = tb(s, M + 0.31, 6.72, 11.78, 0.3) run(para(tf, True), "* This em dash was hand written (not AI generated ☺)", 9, MUTED) # U+263A, not an emoji smiley: the colour-emoji planes (U+1F600+) are dropped @@ -625,30 +649,26 @@ section("Context · what the code does", s = new_slide() chrome(s, 3, "The algorithm · what it is for", "A photon is not a pixel — it is a cluster") -bullets(s, M, 1.90, 7.9, [ - "Charge from one absorbed photon **spreads over neighbouring pixels**. " - "Summing that 3×3 patch recovers the photon energy; a single pixel does not.", - "The histogram of those cluster energies **is** the measurement: peak position " - "and width give the detector's gain and **energy resolution**.", -], size=10.5) -figure(s, "fig_frame", M, 2.98, 7.45) +bullets(s, M, 1.82, 8.16, [ + "One photon's charge **spreads over neighbouring pixels**; summing the 3×3 " + "patch recovers its energy.", + "The histogram of those cluster energies **is** the measurement: peak and " + "width give gain and **energy resolution**.", +], size=PT_BODY) +figure(s, "fig_frame", M, 3.20, 7.70) rail(s, [("label", "Why it matters")], y0=1.95) h = card_figure(s, "img_spectra", RAIL_X, 2.32, RAIL_W) caption(s, RAIL_X, 2.32 + h + 0.10, RAIL_W, - "Cluster-energy spectra from an energy scan, against allpix² simulation. " - "Each peak is one beam energy; its width is the resolution being measured.", - size=8.5) + "Cluster-energy spectra from an energy scan.", size=PT_META) rail(s, [ ("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), ], 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. The detector's frame rate is the " - "number every throughput figure in this deck should be read against.") +caption(s, M, 7.02, 8.16, + "Real MÖNCH data, Cu fluorescence, MAX IV. One cluster per local maximum, " + "so 2 330 counts photons.") notes(s, """The detector, and why its frame rate is the yardstick. MOENCH03 is a hybrid silicon pixel detector: charge integration with analog @@ -669,15 +689,13 @@ that is the number that sets how much of the frame the 3x3 windows cover.""") # =========================================================== 4 · PER FRAME s = new_slide() chrome(s, 4, "The algorithm · per frame", "Per pixel: subtract, threshold, update the pedestal") -bullets(s, M, 1.95, COL, [ - "Per pixel: subtract a **running pedestal** (mean ± rms), keep pixels above " - "**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.", +bullets(s, M, 1.82, COL, [ + "Per pixel: subtract a **running pedestal**, keep what clears **nσ · rms**, " + "cut a 3×3 cluster around each local maximum.", + "The pedestal is **updated by ~80 % of pixels every frame**, so arithmetic " + "and data movement are coupled.", ]) -code(s, M, 3.55, COL, [ +code(s, M, 3.36, COL, [ "// the whole algorithm, per pixel", "v = frame[i] - pedestal_mean[i]", "rms = pedestal rms at i", @@ -687,24 +705,32 @@ code(s, M, 3.55, COL, [ " else -> «nothing» // ... I am in its shadow", "else -> update pedestal // I saw nothing", ], title="THE WHOLE ALGORITHM · THREE OUTCOMES, NOT TWO") -callout(s, M, 5.55, COL, +callout(s, M, 5.86, COL, "**Thesis of this talk:** the compute was fast almost immediately. " "Six of the seven steps are about feeding it.") rail(s, [ ("label", "The shape of the work"), - ("gap", 0.15), + ("gap", 0.20), ("stat", "Work items per frame", "160 000", PALE), + ("gap", 0.10), ("row", "Operations on each", "~5, identical", TEXT2), - ("gap", 0.12), + ("gap", 0.18), ("row", "Communication between them", "none", ACCENT), + ("gap", 0.18), ("row", "Order they may run in", "any, once the pedestal is fixed", ACCENT), - ("gap", 0.22), - ("note", "A pixel reads what its neighbours MEASURED, never what they decided. " - "With the pedestal held at its frame-start value, order cannot " - "matter. That is what licenses one thread per pixel. The serial CPU " - "updates the pedestal mid-scan and so gives that property up; annex " - "A7 measures what it costs."), -]) + ("gap", 0.20), + ("row", "Bytes per frame", "312.5 kB · ~2 330 clusters", TEXT2), +], y0=2.10) +notes(s, """Why order does not matter, and the one place it does. + +A pixel reads what its neighbours MEASURED, never what they decided. With the +pedestal held at its frame-start value, order cannot matter, and that is what +licenses one thread per pixel. + +The serial CPU finder updates the pedestal mid-scan and so gives that property +up: a pixel scanned late sees a pedestal that earlier pixels have already moved. +Annex A7 measures what that costs -- it is the channel behind 11 of the 19 +cluster differences between the serial CPU and the frozen-pedestal reference.""") # ==================================================== 5 · THE CPU # The two machine slides. They exist because the audience is asked, from slide 7 @@ -716,21 +742,18 @@ 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.84, 6.5) +figure(s, "img_cpu_core", M, 1.82, 6.72) # 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. -figure(s, "img_cpu_die", M, 5.04, 4.55) -callout(s, 5.42, 5.04, 3.18, - "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=1.34, size=10) -caption(s, 5.42, 6.50, 3.18, - "Comet Lake · 10 cores boxed; nearly half the die is cache and I/O.", size=8) -caption(s, 5.42, 6.88, 3.18, - "Both diagrams after Stanford CS149, Fall 2025.", size=8) +figure(s, "img_cpu_die", M, 5.04, 4.70) +callout(s, 5.38, 5.04, 3.29, + "Count the boxes: **6 fetch/decode** and two levels of private cache, all " + "to keep **two** instruction streams fed. The ALUs are the small part.", h=1.34, size=PT_LEAD) +caption(s, 5.38, 6.58, 3.29, + "Both diagrams after Stanford CS149, Fall 2025.", size=PT_META) rail(s, [ ("label", "pc-moench-04 · AMD Ryzen 9 7950X"), ("gap", 0.10), @@ -777,21 +800,19 @@ speedup without lying about anything.""") s = new_slide() chrome(s, 6, "The machine we are moving to", "GPU: throughput-oriented, the whole frame at once") -figure(s, "img_gpu_die", M, 2.10, 2.75) -caption(s, M, 4.92, 2.75, - "AD102 · 144 blocks, 128 enabled on this card. One SM boxed.", size=7.5) -figure(s, "img_gpu_sm", M + 3.05, 1.92, 4.75) +figure(s, "img_gpu_die", M, 1.82, 2.84) +caption(s, M, 5.16, 2.84, + "AD102 · 144 blocks, 128 enabled on this card. One SM boxed.", size=PT_META) +figure(s, "img_gpu_sm", M + 3.15, 1.82, 4.91) # The colour key is a separate crop: in the source it spans the full slide width # while the diagram spans 60 % of it, so one rectangle cannot hold both. -figure(s, "img_gpu_legend", M + 3.05, 5.26, 4.75) -callout(s, M, 6.02, COL, +figure(s, "img_gpu_legend", M + 3.15, 5.16, 4.91) +callout(s, M, 5.86, COL, "Same grammar, inverted proportions: **4 fetch/decode** for **64 warp " - "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.82, COL, + "contexts**.", h=0.72, size=PT_LEAD) +caption(s, M, 6.74, 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) + "48 warp slots, 100 kB shared memory. After Stanford CS149, Fall 2025.", size=PT_META) rail(s, [ ("label", "NVIDIA GeForce RTX 4090"), ("gap", 0.12), @@ -830,30 +851,26 @@ That gap between 512 and 9 IS the talk.""") # =========================================================== 7 · THE KERNEL s = new_slide() chrome(s, 7, "The CUDA kernel · execution model", "One thread per pixel") -flow(s, M, 1.90, 11.9, ["load tile + halo", "__syncthreads", "stencil reduction", +flow(s, M, 1.82, 12.30, ["load tile + halo", "__syncthreads", "stencil reduction", "classify", "append or update pedestal"]) -bullets(s, M, 2.85, 7.5, [ - "A **16×16 block = 256 threads** covers 256 pixels; the grid tiles the " - "whole 400×400 frame, 625 blocks of identical work.", - "The output is **sparse**: only detections touch global memory, through one " - "atomic bump of a per-frame counter. The **decision work is dense**: every " - "pixel is tested, independently and identically.", - "That is exactly the shape a GPU wants: regular, independent, repeated " - "160 000 times per frame.", -], size=10.5) -code(s, M, 4.35, 7.5, [ +bullets(s, M, 2.62, 7.75, [ + "A **16×16 block = 256 threads**; the grid tiles the frame in **625 " + "identical blocks**.", + "The output is **sparse**, but the **decision work is dense**: every pixel " + "is tested, **160 000 times per frame**.", +], size=PT_BODY) +code(s, M, 3.98, 7.75, [ "block = dim3(BLOCK_X, BLOCK_Y); // 16 x 16", "grid = dim3((ncols + BLOCK_X - 1)/BLOCK_X,", " (nrows + BLOCK_Y - 1)/BLOCK_Y);", "device::find_clusters_in_single_frame<<>>(", " 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.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.") -code(s, 8.5, 2.85, 4.1, [ +], size=PT_CODE, title="LAUNCH CONFIGURATION · ClusterFinderCUDA.hpp") +caption(s, M, 5.78, 7.75, + "400×400 → a 25×25 grid of 16×16 blocks = 625 blocks per frame, over 128 " + "SMs. Nothing in the launch depends on the cluster count.") +code(s, 8.56, 2.62, 4.24, [ "// ClusterFinder.hpp · the serial CPU", "if (max > nSigma * rms) {", " if (value < max)", @@ -863,26 +880,25 @@ code(s, 8.5, 2.85, 4.1, [ "} else {", " pedestal.«push_fast»(iy, ix, ...);", "}", -], size=8, title="THREE OUTCOMES, NOT TWO") -callout(s, 8.5, 5.10, 4.1, - "Verbatim from the CPU source. A shadow pixel is **neither recorded nor fed " - "back**.", h=0.74, size=10) -callout(s, 8.5, 6.00, 4.1, - "**~80 % of threads update the pedestal**, ~18 % are shadow, ~1.5 % peaks. " - "The update dominates, not the write.", h=0.86, size=10, color=AMBER) +], size=PT_CODE, title="THREE OUTCOMES, NOT TWO") +callout(s, 8.56, 5.04, 4.24, + "A shadow pixel is **neither recorded nor fed back**.", + h=0.72, size=PT_LEAD) +callout(s, 8.56, 5.96, 4.24, + "**~80 % of threads update the pedestal**, ~18 % shadow, ~1.5 % peaks.", + h=0.72, size=PT_LEAD, color=AMBER) # =========================================================== 6 · TILING s = new_slide() chrome(s, 8, "The CUDA kernel · shared memory", "Load the tile once, reuse it nine times") -bullets(s, M, 1.92, 12.0, [ - "Neighbouring threads need **overlapping** 3×3 windows. Without shared memory " - "each pixel would be fetched from global memory up to nine times.", - "Each block stages a tile of (16 + 2r) × (16 + 2r) **pedestal-subtracted** " - "values, the halo is the price of the stencil, and it is loaded cooperatively " +bullets(s, M, 1.82, 12.40, [ + "Neighbouring threads need **overlapping** windows; without shared memory " + "each pixel is fetched up to **nine times**.", + "Each block stages one **pedestal-subtracted** tile plus its halo, loaded " "by the threads on the block edges.", -], size=10.5) -figure(s, "fig_tile", M, 3.05, 7.5) -code(s, 8.4, 3.05, 4.3, [ +], size=PT_BODY) +figure(s, "fig_tile", M, 2.70, 8.36) +code(s, 9.02, 2.70, 3.88, [ "extern __shared__ unsigned char smem[];", "auto *sh = (COMPUTE_TYPE*)smem;", "auto stride = blockDim.x + 2*col_radius;", @@ -890,19 +906,16 @@ code(s, 8.4, 3.05, 4.3, [ " + (threadIdx.x + col_radius);", "// 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.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, 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.") -caption(s, M, 6.70, 7.5, - "Halo cost falls with block size: 56 % of the tile at 8×8, 27 % at 16×16, " - "13 % at 32×32, which is the first half of the block-size argument. " - "The second half is registers, next slide.") +], size=PT_CODE, title="clusterfinder_kernel.cuh") +callout(s, 9.02, 4.72, 3.88, + "Only **odd** cluster sizes are supported, so the centre pixel is unique.", + h=0.92, size=PT_LEAD) +notes(s, """Shared memory is never the binding resource here. + +The tile is stored in COMPUTE_TYPE (float), not in the pedestal type: 1.3 KB for +3x3 and 2.3 KB for 9x9, against 100 KB of shared memory per SM. Even the old +double-precision tile only reached 4.5 KB. What runs out first is registers -- +the next slide.""") # ================================================= 7 · REGISTERS (the input) # Registers come BEFORE occupancy, not after. Occupancy is an output of the @@ -911,24 +924,23 @@ caption(s, M, 6.70, 7.5, s = new_slide() 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, [ +bullets(s, M, 1.82, COL, [ + "An SM has **65 536 registers** and **1 536 thread slots**; whichever runs " + "out first decides how many blocks fit at once.", + "Every thread keeps a private **clusterData[CSX × CSY]**, so register " + "demand grows with the **square** of the cluster size. Neither build " + "spills.", +], size=PT_BODY) +figure(s, "fig_regpressure", M, 3.20, 8.16) +callout(s, M, 5.04, COL, + "**9×9:** the register file is exactly full at two blocks " + "(2 × 128 × 256 = 65 536). **3×3:** the slots run out first, and the " + "registers still have room.", h=0.80, size=PT_LEAD) +code(s, M, 6.20, 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", -], size=8.5, title="MEASURED, NOT ESTIMATED · READ FROM THE BUILT .SO") +], size=PT_CODE, title="MEASURED, NOT ESTIMATED · READ FROM THE BUILT .SO") rail(s, [ ("label", "Per SM · sm_89 · f32 build"), ("gap", 0.15), @@ -937,9 +949,6 @@ rail(s, [ ("gap", 0.05), ("row", "Spills, either case", "0 bytes", TEXT2), ("row", "3×3 on the f64 build", "47 regs → 5 blocks", AMBER), - ("gap", 0.20), - ("note", "How many blocks fit is the whole hardware story. The next slide is " - "what that buys."), ]) notes(s, """The arithmetic, and the one build-dependent caveat. @@ -961,24 +970,23 @@ calls cudaOccupancyMaxActiveBlocksPerMultiprocessor on the built kernel.""") 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, [ +bullets(s, M, 1.82, 12.30, [ + "When a warp stalls, the SM switches to another **already-resident** warp. " + "**Occupancy = resident warps / the maximum** — how many alternatives it " + "has to switch to.", +], size=PT_BODY) +statstrip(s, M, 2.55, 12.30, [ ("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) +], h=0.84) +figure(s, "fig_occupancy", 1.19, 3.44, 10.95) +callout(s, M, 6.62, 12.30, + "**16×16 is the balance point.** 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=PT_LEAD) 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 @@ -1015,32 +1023,29 @@ rows = [ ("opt7", "FP32 pedestal + variance rewrite", "the only kernel change in the deck", "kernel −41%", AMBER), ] -y = 1.70 +y = 1.78 for i, (tag, name, sub, gain, col) in enumerate(rows): if tag == "act": - rect(s, M, y + 0.30, 11.9, 0.016, col) - tf = tb(s, M + 0.02, y, 8.0, 0.28) - run(para(tf, True), name.upper(), 11, col, bold=True, spc=1.4) - tf = tb(s, M + 5.2, y + 0.03, 6.7, 0.26) - run(para(tf, True, align=PP_ALIGN.RIGHT), f"{sub} {gain}", 9, MUTED) - y += 0.40 + rect(s, M, y + 0.31, 12.30, 0.016, col) + tf = tb(s, M + 0.02, y, 8.27, 0.28) + run(para(tf, True), name.upper(), 12, col, bold=True, spc=1.4) + tf = tb(s, M + 5.37, y + 0.04, 6.92, 0.26) + run(para(tf, True, align=PP_ALIGN.RIGHT), f"{sub} {gain}", 10, MUTED) + y += 0.42 continue - rect(s, M, y, 11.9, 0.48, PANEL) - rect(s, M, y, 0.035, 0.48, col) - tf = tb(s, M + 0.28, y + 0.09, 1.0, 0.32) - run(para(tf, True), tag, 12.5, col, bold=True, font=MONO) - tf = tb(s, M + 1.45, y + 0.04, 5.0, 0.28) - run(para(tf, True), name, 11.5, PALE, bold=True) - tf = tb(s, M + 1.45, y + 0.26, 5.6, 0.26) - run(para(tf, True), sub, 9, MUTED) - tf = tb(s, 9.4, y + 0.09, 3.1, 0.35) - 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. 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.") + rect(s, M, y, 12.30, 0.50, PANEL) + rect(s, M, y, 0.04, 0.50, col) + tf = tb(s, M + 0.29, y + 0.10, 1.03, 0.32) + run(para(tf, True), tag, 13, col, bold=True, font=MONO) + tf = tb(s, M + 1.50, y + 0.04, 5.17, 0.28) + run(para(tf, True), name, 12.5, PALE, bold=True) + tf = tb(s, M + 1.50, y + 0.27, 5.79, 0.26) + run(para(tf, True), sub, 10, MUTED) + tf = tb(s, 9.49, y + 0.09, 3.20, 0.35) + run(para(tf, True, align=PP_ALIGN.RIGHT), gain, 13.5, col, bold=True) + y += 0.55 +caption(s, M, 6.98, 12.30, + "Speedups are 3×3 against the best CPU configuration, 24 threads.") # --------------------------------------------------------- divider · ACT I section("Act I of III · feed the GPU", @@ -1064,24 +1069,16 @@ section("Act I of III · feed the GPU", s = new_slide() chrome(s, 12, "Act I · the two memories", "The GPU consumes 32× faster than PCIe delivers") -bullets(s, M, 1.92, 11.9, [ +bullets(s, M, 1.82, 12.30, [ "A CUDA kernel addresses **the GPU's own memory and nothing else**. So every " "frame is copied in — **H2D, host to device** — and every result copied back " "out — **D2H**. Those two names are what the next five slides are about.", -], size=11) -figure(s, "fig_gpu_model", M + 0.74, 2.42, 10.45) -callout(s, M, 6.10, 5.85, - "**The asymmetry is the whole talk.** Host DRAM runs at ~71 GB/s and VRAM at " - "1 008, but everything between them crawls through **31.5**. Arithmetic is " - "cheap once the data is there; arriving is not.", h=0.88, size=10.5) -callout(s, 6.75, 6.10, 5.85, - "**At 3×3, in numbers:** the kernel needs **5.5 µs** per frame. Getting that " - "frame across PCIe costs **13.2 µs** uncontended. The link wins before " - "the kernel has done anything.", h=0.88, size=10.5, color=AMBER) -caption(s, M, 7.04, 11.9, - "Both copies are performed by dedicated DMA engines, one per direction, so " - "in principle they can run while the kernel does. Making that actually " - "happen is opt2 through opt6.", size=9) +], size=PT_BODY) +figure(s, "fig_gpu_model", 0.95, 2.82, 11.20) +callout(s, M, 6.40, 12.30, + "**The asymmetry is the whole talk.** Host DRAM runs at ~71 GB/s and VRAM " + "at 1 008, but everything between them crawls through **31.5**.", + h=0.78, size=PT_LEAD) notes(s, """This slide is for the half of the room that has never written CUDA. If they are all GPU people, say the two sentences on the left and move on. @@ -1119,29 +1116,24 @@ at 3x3 f32, D2H payload 93 kB at 3x3.""") s = new_slide() chrome(s, 13, "Act I · the first CUDA port", "The first port runs at 26 % of the GPU's floor", opt=1) -bullets(s, M, 1.95, COL, [ +bullets(s, M, 1.82, COL, [ "Shared-memory tiling with **halo loading** for any cluster size; pedestal " "subtraction fused into the tile load.", - "Cluster geometry is a **compile-time template parameter** → the 3×3 stencil " - "is fully unrolled.", - "One cudaMemcpy in, one kernel, one cudaMemcpy out; **the host blocks " - "on every frame**.", + "One cudaMemcpy in, one kernel, one cudaMemcpy out; **the host blocks on " + "every frame**.", ]) -code(s, M, 3.18, COL, [ +code(s, M, 3.20, COL, [ "// one frame at a time, the host waits at every step", "cudaMemcpy(d_frame, h_frame, bytes, cudaMemcpyHostToDevice);", "find_clusters_in_single_frame", " <<>>(d_frame, d_pd_mean, ...);", "cudaMemcpy(h_out, d_out, out_bytes, cudaMemcpyDeviceToHost);", ], title="ClusterFinderCUDAOpt2.hpp · find_clusters()") -callout(s, M, 4.62, COL, - "**PCIe is full-duplex**: H2D, kernel and D2H run on independent engines and " - "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 21; it does not change this one.", - h=1.00, size=10.5) -figure(s, "fig_opt1_timeline", M, 5.80, COL) +callout(s, M, 4.82, 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. " + "At 3×3 that is **16.2 µs → 61 859 FPS**.", h=1.0, size=PT_LEAD) +figure(s, "fig_opt1_timeline", M, 5.97, COL) rail(s, [ ("label", "opt1 · 3×3 · 100 k frames · f64"), ("gap", 0.10), @@ -1149,6 +1141,8 @@ rail(s, [ ("stat", "vs 24-thread CPU", "×2.34", ACCENT), ("gap", 0.05), ("row", "Per frame", "63.3 µs", TEXT2), + ("gap", 0.10), + ("row", "H2D · kernel · D2H [s1]", "13.2 · 5.5 · 6.4 µs", TEXT2), ("row", "The GPU floor", "61 859 FPS", ACCENT), ("row", "% of floor", "26 %", AMBER), ]) @@ -1157,12 +1151,11 @@ rail(s, [ s = new_slide() chrome(s, 14, "Act I · streams and batching", "Four streams and 2 000-frame batches: ×1.56", opt=2) -bullets(s, M, 1.95, COL, [ - "A **stream** is an ordered queue of GPU work. Work in **different** streams may " - "overlap, so a copy can run while another stream computes.", - "Each stream gets its own **StreamContext**: device frame buffer, output buffer " - "and pedestal. Frames are handed out **round-robin**.", - "The host now submits **2 000 frames per call** instead of one.", +bullets(s, M, 1.82, COL, [ + "A **stream** is an ordered queue of GPU work. Work in **different** " + "streams may overlap, so a copy can run while another stream computes.", + "Each stream owns a **StreamContext** — frame buffer, output buffer, " + "pedestal — and frames go **round-robin**, **2 000 per call**.", ]) code(s, M, 3.30, COL, [ "struct StreamContext {", @@ -1172,10 +1165,10 @@ code(s, M, 3.30, COL, [ "};", "auto &sc = v_sc[frame_idx % «n_streams»]; // round-robin", ], title="ClusterFinderCUDA.hpp · per-stream state") -figure(s, "fig_opt2_timeline", M, 4.77, COL) -callout(s, M, 6.48, COL, +figure(s, "fig_opt2_timeline", M, 4.98, 7.60) +callout(s, M, 6.66, COL, "**Scaffolding, not yet the payoff.** The streams exist, but the host still " - "synchronises after every round: see opt3.", h=0.60) + "synchronises after every round: see opt3.", h=0.58) rail(s, [ ("label", "opt2 · 3×3 · 4 streams · batch 2 000"), ("gap", 0.10), @@ -1191,36 +1184,35 @@ rail(s, [ s = new_slide() chrome(s, 15, "Act I · remove the sync barriers", "One sync per batch, not one per round: ×1.18", opt=3) -bullets(s, M, 1.95, 7.4, [ - "opt2 synchronised **all streams after every round** of n_streams frames. " - "The GPU drained to empty each time.", - "opt3 submits every frame's H2D → kernel → D2H **asynchronously**, then " - "synchronises **once at the end of the batch**.", -], size=10.5) -figure(s, "fig_streams", M, 3.05, 6.55) -code(s, 8.35, 1.95, 4.25, [ +# The diagram is the argument on this slide, so it takes the whole left column +# and the two bullets go underneath it as the caption they effectively are. +figure(s, "fig_streams", M, 1.86, 8.05) +code(s, 8.80, 1.86, 4.10, [ "// opt2: barrier after every round", "for (round) {", " submit(n_streams frames);", " «cudaDeviceSynchronize»();", "}", - "", "// opt3: submit everything, sync once", "for (frame : batch) {", " cudaMemcpyAsync(..., sc.stream);", " kernel<<<..., sc.stream>>>(...);", " cudaMemcpyAsync(..., sc.stream);", "}", - "for (sc : streams)", - " «cudaStreamSynchronize»(sc.stream);", -], size=8, title="THE ONE-LINE IDEA") -callout(s, 8.35, 5.05, 4.25, - "**29 188 FPS · ×4.32**\n34.3 µs/frame · 47 % of floor (was 40)", h=0.86, size=11) -caption(s, 8.35, 6.15, 4.25, - "Each lane is one stream. Removing the barrier lets a stream start its next " - "frame while its neighbours are still copying. The three panels are scheduled, " - "not sketched: H2D and D2H are one FIFO engine each, so a stream waits for the " - "copy engine, never for another stream's copy to finish overlapping it.") + "«cudaStreamSynchronize»(...); // once", +], title="THE BARRIER, BEFORE AND AFTER") +callout(s, 8.80, 5.10, 4.10, + "**29 188 FPS · ×4.32**\n34.3 µs/frame · 47 % of floor (was 40)", + h=0.86, size=PT_LEAD) +bullets(s, M, 5.86, 12.30, [ + "opt2 synchronised **all streams after every round**, draining the GPU empty " + "each time.", + "opt3 submits every frame's H2D → kernel → D2H **asynchronously** and " + "synchronises **once per batch**.", +], size=PT_BODY) +caption(s, M, 6.90, 12.30, + "Each lane is one stream. The panels are scheduled, not sketched: H2D and " + "D2H are one FIFO engine each.") # ===================================================== 15 · OPT3b · ONE D2H # opt3's title has always said "barriers", plural, but the deck only ever told @@ -1230,35 +1222,31 @@ caption(s, 8.35, 6.15, 4.25, s = new_slide() chrome(s, 16, "Act I · the other barrier", "One D2H per frame, not two", opt=3) -bullets(s, M, 1.92, 11.9, [ +bullets(s, M, 1.82, 12.30, [ "opt2 asked the device **how many clusters**, blocked until the answer came " "back, then asked for **that many**. The size of the second copy was a " "function of data that had not arrived yet.", "opt3 gives every frame a **fixed envelope** — count, then room for **cap** " "clusters — so the copy's size is known at construction and can be queued " "with the kernel. The count is still read, but **afterwards**, on the host.", -], size=11) +], size=PT_BODY) -flow(s, M, 3.42, 11.9, +flow(s, M, 3.20, 12.30, ["kernel", "copy 4 B", "BLOCK", "read count", "copy N B", "BLOCK"], h=0.62) -caption(s, M, 4.12, 11.9, - "opt2 · two transfers and two stalls per frame, because the second one " - "cannot be issued until the first has landed.", size=9.5) +caption(s, M, 3.97, 12.30, + "opt2 · two transfers and two stalls per frame: the host must learn the " + "count before it knows how much to fetch.", size=PT_META) -flow(s, M, 4.62, 11.9, +flow(s, M, 4.31, 12.30, ["kernel", "copy the whole envelope", "→ next frame, host not involved"], h=0.62) -caption(s, M, 5.32, 11.9, - "opt3 · one transfer, no stall. Nothing in the loop waits on a value.", size=9.5) +caption(s, M, 5.08, 12.30, + "opt3 · one transfer, no stall. Nothing in the loop waits on a value.", size=PT_META) -callout(s, M, 5.86, 11.9, - "**You cannot stream a transfer whose length depends on the transfer " - "before it.** opt3 pays bytes to delete that dependency: the envelope is " - "sized by the **cap**, not by how many clusters were found, so an empty " - "frame costs the same D2H as a full one: 120 kB at 3×3 against 93 kB of " - "real clusters.", h=0.94, size=10.5) -caption(s, M, 6.94, 11.9, - "Everything downstream needs that fixed layout: opt6 could not hand out a " - "view into a buffer whose shape was not known in advance.", size=9) +callout(s, M, 5.50, 12.30, + "**You can only stream a transfer whose size is known in advance.** The " + "price is bytes: the envelope is sized by the **cap**, not by the clusters " + "actually found. **3×3: 120 kB shipped for 93 kB of clusters. 9×9: 558 " + "for 467.**", h=0.86, size=PT_LEAD) notes(s, """The point to say out loud: this is the one step in the ladder that is not a setting. It changed the kernel signature, the buffer ownership and the collection loop, and it is why the opt1/opt2 class is frozen in a separate header @@ -1291,15 +1279,14 @@ no single object to copy. Merging the two buffers is what created one.""") s = new_slide() chrome(s, 17, "Act I · pinned (page-locked) memory", "Pinning the input buys DMA-speed H2D: ×1.32", opt=4) -bullets(s, M, 1.95, 12.0, [ - "Normal host memory is **pageable**: the OS may move or swap it. A DMA engine " - "cannot safely read that, so the driver first copies your data into a **hidden " - "pinned staging buffer**. Every transfer is copied twice.", - "**Pinning** locks the pages in physical RAM. The GPU's DMA engine then reads " - "host memory **directly**, no staging copy, and the transfer can be truly asynchronous.", -], size=10.5) -figure(s, "fig_pinning", M, 3.15, 7.6) -code(s, 8.5, 3.15, 4.1, [ +bullets(s, M, 1.82, COL, [ + "Normal host memory is **pageable**, so the driver stages every transfer " + "through a **hidden pinned buffer**: copied twice.", + "**Pinning** locks the pages in RAM, so the DMA engine reads host memory " + "**directly**.", +], size=PT_BODY) +figure(s, "fig_pinning", M, 3.36, 7.30) +code(s, 8.80, 3.36, 4.10, [ "// pin the whole dataset once", "«cudaHostRegister»(ptr, bytes,", " cudaHostRegisterDefault);", @@ -1307,19 +1294,24 @@ code(s, 8.5, 3.15, 4.1, [ "// ... run the whole campaign ...", "", "«cudaHostUnregister»(ptr);", -], size=8, title="ClusterFinderCUDA.hpp") -callout(s, 8.5, 4.90, 4.1, +], size=PT_CODE, title="ClusterFinderCUDA.hpp") +callout(s, 8.80, 5.30, 4.10, "**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.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.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) + h=1.00, size=PT_LEAD) +callout(s, M, 6.44, 12.30, + "**Measured H2D:** one 312.5 KiB frame in **13.2 µs = 24.2 GB/s** — 77 % of " + "PCIe 4.0 ×16, i.e. true DMA speed. **Pageable staging manages ~15 GB/s** " + "for the same frame, because it is copied twice.", h=0.78, size=PT_LEAD, + color=AMBER) +notes(s, """Which H2D number, and why there are two. + +13.2 us is uncontended [s1]. In the shipped pipeline H2D reads 16.6 us [s4]; the +extra 26 % is H2D<->D2H contention on the link, and the full engine grid is in +annex A1. + +The rule this slide is the first sighting of: pinning attacks H2D, which is the +tallest bar at 3x3 and the shortest at 9x9. So the same change is worth x1.32 at +3x3 and only x1.03 at 9x9.""") # -------------------------------------------------------- divider · ACT II section("Act II of III · get the results back", @@ -1338,21 +1330,19 @@ section("Act II of III · get the results back", s = new_slide() chrome(s, 18, "Act II · host↔GPU overlap", "Overlapping host and GPU hides min(host, GPU): ×1.31", opt=5) -bullets(s, M, 1.92, COL, [ - "opt3 overlapped H2D ∥ kernel ∥ D2H **across streams, inside one batch**, but " - "never the **host** with the GPU: find_clusters_batched synchronised, then built " - "thousands of ClusterVectors with the GPU idle.", - "opt5 keeps **one batch in flight while materialising the previous one**: chunk " - "i+1 is submitted before chunk i is collected.", +bullets(s, M, 1.82, COL, [ + "opt3 overlapped the GPU's own engines, never the **host** with the GPU: " + "the batch synchronised, then built ClusterVectors with the GPU idle.", + "opt5 keeps **one batch in flight while materialising the previous one**.", ]) -figure(s, "fig_overlap", M - 0.15, 3.26, COL + 0.30) -code(s, M, 5.88, COL, [ +figure(s, "fig_overlap", 0.42, 2.94, 8.50) +code(s, M, 5.70, COL, [ "tok = cf.«submit_batch»(data[a0:b0], first_frame=a0)", "for a, b in bounds[1:]:", " nxt = cf.«submit_batch»(data[a:b], first_frame=a) // GPU starts i+1", " results.extend(cf.«collect»(tok)) // host unpacks i", " tok = nxt", -], size=9, title="you never write these: find_clusters_batched() wraps them") +], size=PT_CODE, title="you never write these: find_clusters_batched() wraps them") notes(s, """opt5 — host<->GPU overlap. Code and chunk sizing are on annex A3. tok = cf.submit_batch(data[a0:b0], first_frame=a0) @@ -1385,11 +1375,6 @@ rail(s, [ ("gap", 0.06), ("row", "CUDA API calls added", "none", TEXT2), ("row", "what moved", "the host loop", PALE), - ("gap", 0.16), - ("note", "Told at 3×3, where the host copy is SHORTER than the GPU floor and " - "tucks underneath it. 9×9 is the other case, and it is the next " - "slide. For clusters that must outlive the finder, 3×3 opt5 is the " - "endpoint: opt6 lends, it does not give."), ]) @@ -1402,22 +1387,20 @@ rail(s, [ s = new_slide() chrome(s, 19, "Act II · at 9×9 · why overlap runs out", "Overlap runs out: the host is the taller bar", opt=5) -bullets(s, M, 1.84, 11.9, [ - "At 3×3 the host copy is **shorter than the GPU floor** and hides underneath " - "it. At 9×9 it is **roughly twice the floor** — ~62 µs of malloc-and-copy " - "against 30.01 µs of GPU — so overlap still works, it just has less to hide. " - "That is why opt5 is worth **×1.20** here and ×1.31 at 3×3.", -], size=10.5) -figure(s, "fig_overlap_9x9", M + 0.15, 2.44, 11.3) -callout(s, M, 6.30, 11.9, +bullets(s, M, 1.82, 12.30, [ + "At 3×3 the host copy is **shorter than the GPU floor** and hides under it. " + "At 9×9 it is **roughly twice the floor**, so overlap has less to hide: " + "opt5 is worth **×1.20** here against ×1.31 at 3×3.", +], size=PT_BODY) +figure(s, "fig_overlap_9x9", 0.66, 2.50, 11.68) +callout(s, M, 6.30, 12.30, "A deeper buffer **relocates the GPU's idle, it does not close it**: the " "host lane is already back-to-back in both strips, so it alone sets the " "pace. **The only way down is to make the host term smaller.**", - h=0.66, size=10.5) -caption(s, M, 7.06, 11.9, - "Measured proportions: GPU 30.01 µs/frame, host ~62 µs steady-state. " - "Fault correction and the raw 66.4 µs are in the notes and annex A4.", - size=9) + h=0.66, size=PT_LEAD) +caption(s, M, 7.06, 12.30, + "Measured proportions: GPU 30.01 µs/frame, host ~62 µs steady-state [A4].", + size=PT_META) notes(s, """This slide exists because "add more slots" is the reliable guess here, and it is worth letting the room say it out loud before the second strip goes up. @@ -1452,22 +1435,17 @@ the host term by about half.""") s = new_slide() chrome(s, 20, "Act II · zero-copy collection", "Read the results in place: ×2.21 at 9×9", opt=6) -bullets(s, M, 1.92, COL, [ - "The D2H lands in a **pinned host buffer**. collect() then allocates one " - "ClusterVector per frame and memcpys into it; at 9×9 that is **467 kB per frame, " - "~9.3 GB per run**, single-threaded.", - "collect_view() returns a **BatchView**: strided numpy views straight onto the " - "pinned buffer. It withholds ownership past the chunk, **not access**, every " - "cluster's payload and coordinates are readable.", -], size=10.5) -figure(s, "fig_resultpath", M - 0.15, 3.05, COL + 0.30) -callout(s, M, 5.86, COL, - "The win is **max(0, host copy − GPU floor)**: at 3×3 the 8 µs copy hides under " - "a 16.2 µs floor and opt5 had already absorbed most of it; at 9×9 the ~62 µs " - "host term is **twice the floor** and cannot hide at any overlap.", h=0.80, size=10) -caption(s, M, 6.78, COL, - "The two bars are the competing costs, not the two steps; the step times are " - "on the right.", size=8.5) +bullets(s, M, 1.82, COL, [ + "The D2H lands in a **pinned host buffer**; collect() then allocates one " + "ClusterVector per frame and memcpys into it — **467 kB per frame at 9×9**.", + "collect_view() returns a **BatchView** onto that buffer instead: it " + "withholds **ownership** past the chunk, not access.", +], size=PT_BODY) +figure(s, "fig_resultpath", 0.42, 3.36, 8.50) +callout(s, M, 6.32, COL, + "The win is **max(0, host copy − GPU floor)**. At 3×3 the copy already hid " + "under the floor; at 9×9 it is **twice the floor** and cannot hide at any " + "overlap.", h=0.80, size=PT_LEAD) notes(s, """Where 93 kB and 467 kB come from: clusters/frame x sizeof(Cluster). sizeof — Cluster.hpp:28 is two CoordType coords then std::array data. @@ -1499,10 +1477,7 @@ rail(s, [ ("stat", "9×9 throughput", "33 323 FPS", PALE), ("row", "per frame · step · of floor", "30.0 µs · ×2.21 · 100 %", AMBER), ("gap", 0.12), - ("note", "opt5 → opt6: 19.8 → 17.1 and 66.4 → 30.0 µs, bit-identical, " - "0.2 % spread, zero warm faults. FLOOR = lower of the s4 engine max " - "and the best sustained rate: at 3×3 the 16.17 µs max sets it; at 9×9 " - "f64 the run BEATS the 32.66 µs max, so the 30.0 sustained sets it [A1]."), + ("row", "Results", "bit-identical to opt5", TEXT2), ]) # ------------------------------------------------------- divider · ACT III @@ -1527,35 +1502,28 @@ section("Act III of III · the kernel", s = new_slide() chrome(s, 21, "How the engine times are measured", "Two configurations, one floor") +# Three definitions, one line of gloss each. Anything longer competes with the +# picture underneath, which is the thing that actually makes s4 make sense. 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."), + "How long an operation actually takes."), ("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."), + "How busy each engine is: a union, not a sum."), ("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."), + "1 / the busiest engine, in µs/frame or FPS."), ] 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) + rect(s, x, 1.82, 3.96, 1.24, PANEL) + rect(s, x, 1.82, 3.96, 0.035, col) + tf = tb(s, x + 0.26, 1.98, 3.41, 0.30) + run(para(tf, True), tag, 15, col, bold=True, font=MONO) + tf = tb(s, x + 0.26, 2.32, 3.50, 0.30) + run(para(tf, True, line=1.1), title, 12, PALE, bold=True) + tf = tb(s, x + 0.26, 2.62, 3.50, 0.40) + run(para(tf, True, line=1.22), body, 11.5, TEXT2) + x += 4.16 +figure(s, "fig_measure", 0.87, 3.32, 11.60) notes(s, """Say the floor out loud, in this order. 1. There are three engines and they are independent: the H2D copy engine, the @@ -1593,36 +1561,34 @@ the floor is the busiest engine." The rest is in A1.""") s = new_slide() chrome(s, 22, "Act III · FP32 device pedestal", "FP32 halves pedestal traffic: −41 % kernel time", opt=7) -bullets(s, M, 1.95, 7.5, [ - "**~80 % of pixels** take the **pedestal-update** branch, which reads and " - "writes six accumulator values. All four pedestal arrays are DEVICE_PED_TYPE, " - "so one typedef halves every one of them: **48 bytes per updating pixel in " - "FP64, 24 in FP32**. The kernel is **bandwidth-bound**, so halving the traffic " - "nearly halves the time.", -], size=10.5) -figure(s, "fig_f32_kernel", M, 3.05, 7.5) -caption(s, M, 6.10, 7.5, - "Two more effects, both real, neither the reason it works: on a GeForce " - "part FP64 arithmetic runs at 1/64 of FP32, and the pedestal update was " - "paying that on every pixel; and the narrower accumulators free 9 registers " - "at 3×3 (47 → 38), buying back a block per SM.", size=9) -code(s, 8.5, 1.95, 4.1, [ +bullets(s, M, 1.82, 7.75, [ + "**~80 % of pixels** take the **pedestal-update** branch: six accumulator " + "values read and written.", + "One typedef halves all four pedestal arrays — **48 bytes per updating " + "pixel in FP64, 24 in FP32** — and the kernel is **bandwidth-bound**.", +], size=PT_BODY) +figure(s, "fig_f32_kernel", M, 3.20, 8.16) +caption(s, M, 5.72, 8.16, + "Two more effects, neither the reason it works: FP64 arithmetic runs at " + "1/64 of FP32 on a GeForce part, and the narrower accumulators free 9 " + "registers at 3×3.", size=PT_META) +code(s, 8.80, 1.82, 4.10, [ "// clusterfinder_kernel.cuh", "using COMPUTE_TYPE = float;", "using DEVICE_PED_TYPE = «float»;", "// was: double", -], size=9, title="ONE TYPEDEF") -callout(s, 8.5, 3.24, 4.1, +], size=PT_CODE, title="ONE TYPEDEF") +callout(s, 8.80, 3.22, 4.10, "Kernel, 9×9 **[s1 · cap 1700]**\n**39.86 → 23.70 µs (−40.5 %)**", - h=1.00, size=10.5) -callout(s, 8.5, 4.36, 4.1, + h=0.94, size=PT_LEAD) +callout(s, 8.80, 4.30, 4.10, "At 3×3 the same typedef is worth **−70.6 %** in the kernel and " "**4.6 %** end to end: there the kernel was never the tallest bar.", - h=1.00, size=10.5) -callout(s, 8.5, 5.48, 4.1, - "Naive FP32 is **wrong** (next slide).", h=0.56, size=10.5, color=AMBER) -caption(s, 8.5, 6.16, 4.1, - "Both builds, same git rev, 20 000 frames. Full engine grid: A1.", size=9) + h=1.06, size=PT_LEAD) +callout(s, 8.80, 5.50, 4.10, + "Naive FP32 is **wrong** (next slide).", h=0.56, size=PT_LEAD, color=AMBER) +caption(s, 8.80, 6.24, 4.10, + "Both builds, same git rev, 20 000 frames. Full engine grid: A1.", size=PT_META) notes(s, """Reading the two panels, and the number to quote. Left is s1: one stream, nothing else running, so those are true durations. The @@ -1654,38 +1620,30 @@ nothing end to end, because at 3x3 H2D is the floor.""") s = new_slide() chrome(s, 23, "Act III · catastrophic cancellation", "Accumulate what is small, not what is large", opt=7) -bullets(s, M, 1.90, COL, [ - (TEXT2, "**The trap.** **var = E[X²] − mean²** subtracts two numbers near " - "2 × 10⁷ to get one near 2 000. FP32 carries ~7 digits, so the answer " - "inherits an **absolute** error that **does not shrink as the answer does**."), - (TEXT2, "**What it cost.** On a quiet pixel that error is a third of the " - "variance; a little lower it goes negative, the rms **clamps to zero**, and a " - "5σ gate becomes a **0σ gate**. Those pixels fire every frame, below the " - "physical threshold."), - (TEXT2, "**The fix.** Freeze **X₀ = round(mean)** at the end of pedestal " - "training and accumulate the **centred** value **Y = X − X₀**. Both operands " - "become O(rms); 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) +bullets(s, M, 1.82, COL, [ + "**The trap.** **var = E[X²] − mean²** takes two numbers near 2 × 10⁷ to " + "get one near 2 000.", + "**What it cost.** The rms **clamps to zero**, so a 5σ gate becomes a " + "**0σ gate**.", + "**The fix.** Accumulate the **centred Y = X − X₀**: both operands become " + "O(rms).", +], size=PT_BODY) +figure(s, "fig_cancellation", M, 3.80, COL) +caption(s, M, 6.34, COL, + "Left: the operands and the answer against the ±3 ADU² error floor. " + "Right: the f32 curve is reconstructed — measured area, modelled shape. " + "The two-line patch is in annex A5.", size=PT_META) rail(s, [ ("label", "naive f32 · what it did"), ("gap", 0.12), ("stat", "Extra clusters", "+28.06 %", AMBER), ("row", "Pixels affected", "~1–2 % of the sensor", AMBER), - ("row", "Operands vs answer", "2.17×10⁷ vs ~2 000", TEXT2), ("row", "Error floor", "±3 ADU², absolute", AMBER), ("gap", 0.26), ("label", "after the rewrite"), ("gap", 0.12), ("row", "f32 vs f64 counts", "3 × 10⁻⁷", ACCENT), ("row", "vs the CPU baseline", "0.0039 %", ACCENT), - ("gap", 0.18), - ("note", "X₀ must never be updated: the accumulators are defined relative to it."), ]) notes(s, """The one sentence to leave the room with. @@ -1716,21 +1674,16 @@ change the update's arithmetic cost.""") s = new_slide() chrome(s, 24, "Act III · why this act comes last", "The saving never grew — the frame around it shrank", opt=7) -figure(s, "fig_f32_absolute", M, 1.95, 11.9) -callout(s, M, 4.90, 5.85, - "**The saving never grew: −4.63 µs at opt4, −4.87 µs at opt6.** What shrank " - "is the frame around it, 79.8 → 30.0 µs, so the same microseconds read " - "**−5.8 %, then −16.2 %**. Act II made the kernel **separable**, not bigger.", - h=1.10, size=10.5) -callout(s, 6.75, 4.90, 5.85, - "**And the act ends by handing the floor away.** At **s4** 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.", - h=1.10, size=10.5, color=AMBER) -caption(s, M, 6.16, 11.9, +figure(s, "fig_f32_absolute", M, 1.82, 12.30) +callout(s, M, 5.06, 12.30, + "**The act ends by handing the floor away.** At **s4** the f64 arm is " + "kernel-bound, **32.66 µs** against a **25.25 µs** D2H. The typedef puts " + "the kernel at **23.94**, below a D2H that never moved.", + h=0.80, size=PT_LEAD, color=AMBER) +caption(s, M, 6.02, 12.30, "9×9 · 20 000 frames · 4 streams · cap 1 700 · warm · both arms at the same " "git rev. opt3 excluded and opt5 marked †: their arms sat in different " - "allocator states. Worked through in annex A4.", size=9) + "allocator states. Worked through in annex A4.", size=PT_META) notes(s, """WHY opt3 IS EXCLUDED AND opt5 IS DAGGERED. Both steps' two arms sat in different allocator states, so the difference between them does not report the typedef -- it reports the heap. opt3 is dropped entirely; opt5's reading agrees @@ -1763,59 +1716,44 @@ section("Results · what came out of it", # =========================================================== 22 · RESULTS s = new_slide() chrome(s, 25, "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, +figure(s, "fig_arc", 1.79, 1.76, 9.71) +callout(s, M, 5.94, 6.04, "**×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; against the CPU twin that isolates the port, " - "**exact on the f64 pedestal** and **6 clusters in 23 M** on the shipped f32 " - "(slides 30–33).", 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 " - "allocator states) · each step in its own process · CPU baseline = " - "ClusterFinderMT at its best thread count, 24 here, first pass only.") +callout(s, 6.75, 5.94, 6.04, + "Every step is **monotonic**, and correctness held constant throughout: " + "**exact on the f64 pedestal**, **6 clusters in 23 M** on the shipped f32.", h=0.8, color=AMBER) +caption(s, M, 6.96, 12.30, + "3×3 · nσ = 5 · 100 000 frames · batch 2 000 · 4 streams · warm = best of " + "reps 1–4 · each step in its own process.") # =========================================================== 23 · RESULTS 9x9 s = new_slide() chrome(s, 26, "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, +figure(s, "fig_arc_9x9", 1.79, 1.76, 9.71) +callout(s, M, 5.94, 12.30, "**×26.5 over 32 CPU threads**; the kernel is the tallest bar for the whole " - "f64 arm, and opt7's −40 % drops it **below D2H**.", h=0.8) -callout(s, 6.75, 5.80, 5.85, - "opt4 buys **×1.03** here and **×1.32** at 3×3. Same code, opposite regimes, " - "the rule, twice.", h=0.8, color=AMBER) -caption(s, M, 6.62, 11.9, - "9×9 · cap 1 700 (lossless; 1 500 truncated 0.0095 % of clusters) · 20 000 frames · " - "opt1/opt2 are 3×3-only · CPU baseline ClusterFinderMT at 32 threads. The floor is the " - "lower of the nsys estimate and the best sustained rate; both arms are sustained-bound, " - "corroborated to 8.8 % (f64) and 0.4 % (f32). Why the arms differ: 544.5 kB of D2H " - "costs 25.2 µs on both [s4], hidden under the 32.7 µs f64 kernel, not under " - "the 23.9 µs f32 one.") + "f64 arm, and opt7's −40 % drops it **below D2H**.", h=0.72) +caption(s, M, 6.86, 12.30, + "9×9 · cap 1 700 · 20 000 frames · CPU baseline at 32 threads. opt1 and " + "opt2 are 3×3-only. Why the two arms differ is in the notes.") # =========================================================== 24 · WHERE TIME GOES s = new_slide() chrome(s, 27, "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, +figure(s, "fig_overhead", 1.79, 1.82, 9.71) +callout(s, M, 5.09, 6.04, "**Acts I and II never touch the arithmetic.** The GPU floor is a flat " "16.2 µs at 3×3 / 30.0 µs at 9×9; what collapses is everything stacked on it.", - h=0.86, size=10.5) -callout(s, 6.75, 5.30, 5.85, + h=0.86, size=PT_LEAD) +callout(s, 6.75, 5.09, 6.04, "**Act III is the only step that lowers the floor itself**, and it could not " - "have been seen until the stack above it was gone.", h=0.86, size=10.5, + "have been seen until the stack above it was gone.", h=0.86, size=PT_LEAD, 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 " - "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.") +caption(s, M, 6.24, 12.30, + "Coloured = the GPU floor for that act's build. Grey = everything the host " + "adds on top: +50 µs at opt3, nothing at opt6.") # ======================================================= 24 · MEASUREMENT AUDIT # Was a three-card slide that gave equal weight to page faults, CUDA-event @@ -1826,23 +1764,21 @@ caption(s, M, 6.45, 11.9, s = new_slide() chrome(s, 28, "Behind the numbers · the result heap, not the GPU", "Materialising clusters costs 2.6 M page faults") -bullets(s, M, 1.90, COL, [ +bullets(s, M, 1.82, COL, [ "Every run that keeps its results **allocates a fresh heap and touches it " - "once**. Linux only finds physical memory on **first touch**, so that first " - "pass is where the OS does its work: **inside the timer**, on a run whose GPU " - "work is under 2 s.", -], size=10.5) -figure(s, "fig_pagefault", M, 2.62, 7.9) -code(s, M, 4.90, 7.9, [ + "once**. Linux only finds physical memory on **first touch**, so that pass " + "is where the OS works — **inside the timer**.", +], size=PT_BODY) +figure(s, "fig_pagefault", M, 2.81, 8.16) +code(s, M, 4.94, 8.16, [ "# bracket every timed cell:", "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 PROTOCOL · BRACKET EVERY TIMED CELL") -callout(s, M, 6.24, 7.9, - "**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) +], size=PT_CODE, title="THE PROTOCOL · BRACKET EVERY TIMED CELL") +callout(s, M, 6.42, 8.16, + "**Re-run until minor faults plateau (< 200 k), and quote that run.**", + h=0.62, size=PT_LEAD) rail(s, [ ("label", "one 100 000-frame pass · 3×3"), ("gap", 0.12), @@ -1851,11 +1787,9 @@ rail(s, [ ("gap", 0.08), ("row", "Clusters materialised", "~10 GB", TEXT2), ("row", "Cost per fault, fitted", "0.68 µs", TEXT2), - ("row", "Major faults, all campaign", "0", ACCENT), - ("gap", 0.18), - ("note", "Not CUDA-specific: every anonymous allocation on Linux behaves this " - "way. Copying faster cannot avoid it; not allocating can. That is " - "opt6."), + ("gap", 0.14), + ("note", "Zero major faults all campaign. Copying faster cannot avoid this; " + "not allocating can. That is opt6."), ]) notes(s, """THE NUMBERS ARE 3x3, opt2, 100 000 frames, one process, results retained: 2 625 948 minor faults on run 1 at 6.110 s / 16 366 FPS, falling to @@ -1892,21 +1826,17 @@ allocations. glibc retains the arenas.""") s = new_slide() chrome(s, 29, "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, - "**Everything that materialises clusters loses a third of its throughput " - "on the first run**, +7 to +20 µs per frame, depending on how much of it " - "the GPU can hide behind its own work.", h=0.68, size=9.5) -callout(s, 6.75, 5.98, 5.85, - "**Only the two ends escape, for opposite reasons.** opt1 never grows the " - "heap: it discards each frame. opt6 never needs one, and reaches **98 % " - "of its peak on a cold process**.", h=0.68, size=9.5, color=AMBER) -caption(s, M, 6.76, 11.9, - "Single pass, one process, every ClusterVector retained, f32, the same " - "100 000 frames. Not the " - "campaign's \"cold\" rep, which discards results and so never grows the heap. Repeat " - "the run and the amber bars climb onto the blue ones; opt1 and opt6 never move, " - "because neither ever paid.", size=8.5) +figure(s, "fig_first_run", 1.57, 1.72, 10.20) +callout(s, M, 5.96, 6.04, + "**Everything that materialises clusters loses a third of its throughput on " + "the first run**: +7 to +20 µs per frame.", h=1.00, size=PT_LEAD) +callout(s, 6.75, 5.96, 6.04, + "**Only the two ends escape, for opposite reasons.** opt1 discards each " + "frame; opt6 never grows a heap, and reaches **98 % of its peak on a cold " + "process**.", h=1.00, size=PT_LEAD, color=AMBER) +caption(s, M, 7.06, 12.30, + "Single pass, one process, every ClusterVector retained · f32 · 100 000 " + "frames. Repeat the run and the amber bars climb onto the blue ones.", size=PT_META) # --------------------------------------------------- divider · VALIDATION section("Validation · does it find the same photons", @@ -1924,21 +1854,18 @@ section("Validation · does it find the same photons", s = new_slide() chrome(s, 30, "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, - "The serial CPU finder updates the pedestal **as the raster scan reaches each " - "pixel**. A CUDA thread cannot: 160 000 of them read the pedestal at once, so " - "the update is **applied at the frame boundary**.", h=1.02, size=10.5) -callout(s, 6.75, 5.30, 5.85, - "So a straight CPU↔CUDA comparison moves **two** things at once. " - "**ClusterFinderFrozen** is the serial finder with only the update moved to " - "the frame end: same arithmetic, same gates, same scan.", h=1.02, size=10.5, - color=AMBER) -caption(s, M, 6.52, 11.9, - "Frozen is a diagnostic twin, not a product: it exists so the next two slides " - "can attribute each disagreement to exactly one cause. cpu vs frozen = update " - "timing; frozen vs cuda = the port. Every finder on these slides is trained on " - "the same 1 000 pedestal frames and run over the same 10 000 data frames.") +figure(s, "fig_pedtiming", M - 0.15, 1.82, 12.61) +callout(s, M, 5.98, 6.04, + "The CPU finder updates the pedestal **as the raster reaches each pixel**. " + "160 000 CUDA threads read it at once, so the update is **applied at the " + "frame boundary**.", h=0.86, size=PT_LEAD) +callout(s, 6.75, 5.98, 6.04, + "A straight CPU↔CUDA comparison therefore moves **two** things at once. " + "**ClusterFinderFrozen** moves only the update: same arithmetic, same " + "gates, same scan.", h=0.86, size=PT_LEAD, color=AMBER) +caption(s, M, 7.02, 12.30, + "Frozen is a diagnostic twin, not a product. Every finder trains on the " + "same 1 000 pedestal frames and runs the same 10 000 data frames.") notes(s, """Why this slide is here. The obvious experiment - run the CPU finder and the CUDA finder over the same @@ -1970,13 +1897,13 @@ Frozen ships in the library as a diagnostic, not as the recommended finder.""") s = new_slide() chrome(s, 31, "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, [ +bullets(s, M, 1.82, 12.40, [ "**ClusterFinderFrozen** makes byte-for-byte the same decisions as ClusterFinder " "and differs in exactly one thing: **when** the pedestal is updated. Frozen per " "frame, pushed at frame end. That is the CUDA model, so comparing against it " "isolates everything else the port changes.", -], size=10.5) -table(s, M, 2.52, 12.0, +], size=PT_BODY) +table(s, M, 2.81, 12.40, ["comparison", "the one thing that differs", "A-only / B-only", "% of clusters"], [["serial CPU vs frozen CPU", "update timing alone: **a CPU-only effect**", @@ -1986,24 +1913,18 @@ table(s, M, 2.52, 12.0, ["frozen CPU vs CUDA [f64 ped]", "**nothing**", "**0 / 0**", "**0 %**"], ["frozen CPU vs CUDA [f32 ped]", "the float32 pedestal EMA drifting: " "see next slide", "0 / 6", "0.000026 %"]], - colw=[0.28, 0.40, 0.17, 0.15], size=9.5, rowh=0.60) + colw=[0.28, 0.40, 0.17, 0.15], size=PT_TABLE, rowh=0.60) # The two rows that measure the PORT with the timing held fixed. Ringed rather # than recoloured, and drawn after the table so the outline sits over the zebra. -frame_rect(s, M - 0.04, 4.13, 12.08, 1.34) -callout(s, M, 5.72, 11.9, - "**The ringed pair is the port on its own, and at f64 it is identical:** " - "23 244 605 clusters, not one disagreement. **cpu vs cuda** equals " - "**cpu vs frozen** exactly, so everything CUDA changes is worth zero and the " - "whole residual is a CPU-vs-CPU effect. The six at f32 are the next slide.", - h=0.80, size=10, color=RED) -caption(s, M, 6.58, 11.9, - "3×3 · 10 000 frames · 23.2 M clusters · same pedestal, same frames · exact " - "centre-set difference at tol = 0. [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. Either row is reproduced by rebuilding with that typedef and " - "re-running the comparison, which counts and localises every disagreement " - "rather than only totalling them.") +frame_rect(s, M - 0.04, 4.43, 12.48, 1.32) +callout(s, M, 5.94, 12.30, + "**The ringed pair is the port on its own.** At f64 the two finders agree " + "on all 23 244 605 clusters. Every difference in the rows above it is a " + "CPU-vs-CPU effect; the six at f32 are the next slide.", + h=0.76, size=PT_LEAD, color=RED) +caption(s, M, 7.02, 12.30, + "3×3 · 10 000 frames · 23.2 M clusters · exact centre-set difference at " + "tol = 0. [f64 ped] and [f32 ped] differ only in DEVICE_PED_TYPE.") notes(s, """Row 1 and row 2 are the same 8/11, which is the point: everything CUDA changes is worth zero, and the whole residual is a CPU-vs-CPU effect. @@ -2036,15 +1957,14 @@ finders freeze the pedestal there. That is the next slide.""") s = new_slide() chrome(s, 32, "Validation · the disagreement, seen", "The whole disagreement is one duplicate centre") -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, +figure(s, "fig_mismatch147", 1.19, 1.72, 10.95) +callout(s, M, 6.34, 6.04, + "Only each finder's **own** 3×3 footprints are drawn, so the panels differ " + "**exactly** where the finders do.", h=0.78, size=PT_LEAD) +callout(s, 6.75, 6.34, 6.04, "cuda's patch is **one row taller**: a second centre directly below the one " "both found. **The charge is already counted**: a duplicate, not a new " - "photon.", h=0.78, size=10.5, color=AMBER) + "photon.", h=0.78, size=PT_LEAD, 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. @@ -2071,7 +1991,9 @@ reproducible from ClusterFinderFrozen_vs_CUDA.ipynb directly.""") s = new_slide() chrome(s, 33, "Validation · the six residuals, dissected", "float32 cannot tell these two pixels apart") -code(s, M, 1.88, 6.35, [ +# The walkthrough IS the slide: the whole argument is that two printed numbers +# are the same number in float32, and that only lands if the room can read both. +code(s, M, 1.82, 6.56, [ "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", @@ -2079,26 +2001,21 @@ code(s, M, 1.88, 6.35, [ " [4763 4640 4858]] [ 1.1 70.7 136.4]] window", "", "the two contenders, at full precision:", - " rival (dy=-1) centre", - "frozen [f64 ped] 638.383019956 638.382773664", - "cuda [f32 ped] 638.382812500 «638.382812500»", + " rival (dy=-1) centre", + "frozen [f64 ped] 638.383019956 638.382773664", + "cuda [f32 ped] 638.382812500 «638.382812500»", "", "gate: accept if centre >= max(window)", " 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.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.00, size=10) -figure(s, "fig_spectra_valid", 7.30, 1.92, 5.40) -callout(s, M, 5.84, 11.9, - "**At f64 the residual is zero.**", h=0.52, size=10.5, color=AMBER) -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.") + " cuda 638.382812500 >= 638.382812500 -> «accept»", +], size=PT_CODE) +callout(s, M, 4.86, 6.56, + "Under the f64 pedestal the two pixels are **0.000246 ADU** apart; one " + "float32 ULP here is **0.000488**. **Half a ULP** — in float32, the same " + "number.", h=1.0, size=PT_LEAD) +figure(s, "fig_spectra_valid", 7.32, 1.82, 5.58) +callout(s, M, 6.10, 12.30, + "**At f64 the residual is zero.**", h=0.52, size=PT_LEAD, color=AMBER) notes(s, """Frame 147, centre (x=202, y=8). The full numbers. centre rival centre - rival @@ -2137,138 +2054,124 @@ Consequences worth stating out loud if asked: # =========================================================== 28 · API 1 s = new_slide() chrome(s, 34, "For users · Python API", "The fast path in eight lines") -code(s, M, 1.95, 7.6, [ +code(s, M, 1.82, 7.85, [ "from aare import File, ClusterFinderCUDA", - "", "cf = ClusterFinderCUDA(image_size=(400, 400), cluster_size=(3, 3),", " n_sigma=5, «n_streams»=4,", " «max_clusters_per_frame»=3000)", - "", "for _ in range(1000): «# 1. train the pedestal»", " cf.push_pedestal_frame(pd.read_frame())", - "", "data = f.read_n(100_000) «# 2. one contiguous array»", "cf.‹register_input_buffer›(data) «# 3. pin it once»", - "", "for s in range(0, len(data), 2000): «# 4. batch through it»", " clusters = cf.‹find_clusters_batched›(data[s:s+2000], first_frame=s)", - "", "cf.‹unregister_input_buffer›() «# 5. release the pages»", -], size=9, title="THE RECOMMENDED PATTERN") +], size=PT_CODE, title="THE RECOMMENDED PATTERN") # The zero-copy variant sits directly under the pattern it replaces, so the two # can be read against each other instead of across the slide. -code(s, M, 5.30, 7.6, [ +code(s, M, 4.62, 7.85, [ "for v in cf.‹find_cluster_views_batched_iter›(data, 2000): «# 4, zero-copy»", " hist.fill(v.sums()) # consume inside the loop:", " # the view dies with its chunk", -], size=9, title="ZERO-COPY · SAME PATTERN, STEP 4 SWAPPED") -caption(s, M, 6.52, 7.6, - "Steps 1, 2, 3 and 5 are unchanged. Only the call inside the loop differs, " - "and with it who owns the memory the clusters live in.", size=9) +], size=PT_CODE, title="ZERO-COPY · SAME PATTERN, STEP 4 SWAPPED") # Right column: one entry per numbered step, so the mapping needs no arrows. -bullets(s, 8.6, 2.0, 4.1, [ - "**3 · register_input_buffer()** is opt4 in one call. Pin **once**, outside the " - "loop; slices of a registered array inherit it. Applies to both patterns.", - "**4 · find_clusters_batched()** returns **one ClusterVector per frame**, in " - "order, and runs opt5's chunked overlap internally, so you get it for free.", -], size=10) -callout(s, 8.6, 4.30, 4.1, +bullets(s, 8.66, 1.82, 4.24, [ + "**3 · register_input_buffer()** — opt4 in one call. Pin **once**, outside " + "the loop.", + "**4 · find_clusters_batched()** — one ClusterVector per frame, with opt5's " + "overlap built in.", +], size=PT_BODY) +callout(s, 8.66, 3.50, 4.24, "**Zero-copy: ×1.16 at 3×3, ×2.21 at 9×9.** The views expose every " "cluster; they only withhold **ownership** past the chunk.", - h=1.05, size=10) -callout(s, 8.6, 5.52, 4.1, - "Reduce as you go and you never allocate: histogram, sum, or copy the few " - "you keep. Hold a view past its chunk and you stall the pipeline.", - h=1.05, size=10, color=AMBER) + h=1.05, size=PT_LEAD) +callout(s, 8.66, 4.70, 4.24, + "Reduce as you go and you never allocate. Hold a view past its chunk and " + "you stall the pipeline.", h=1.05, size=PT_LEAD, color=AMBER) # =========================================================== 27 · API 2 s = new_slide() chrome(s, 35, "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) +hdr = [("Parameter", 1.05), ("What it does", 3.72), ("Guidance", 5.37)] +# The ring and the warning are the same point, so they carry the same colour. +callout(s, M, 1.78, 12.30, + "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.62, size=PT_LEAD, color=RED) +y = 2.56 +rect(s, M, y, 12.30, 0.4, PANEL) for lab, dx in hdr: - tf = tb(s, M + dx - 0.85 if dx > 1.05 else M + 0.28, y + 0.09, 5.0, 0.3) - run(para(tf, True), lab.upper(), 9, MUTED, bold=True, spc=1.3) + tf = tb(s, M + dx - 0.85 if dx > 1.05 else M + 0.28, y + 0.09, 5.17, 0.3) + run(para(tf, True), lab.upper(), PT_LABEL, MUTED, bold=True, spc=1.3) y += 0.44 +# Guidance is what a user will act on, so it says the number and the consequence +# and stops. Everything that explains WHY is in the notes. params = [ - ("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×."), + ("n_streams", "Upper bound on frames in flight.", + "4 at both cluster sizes. 8 buys no kernel concurrency at 9×9."), ("max_clusters_per_frame", "Fixed size of the per-frame D2H transfer.", - "Must exceed the real maximum or clusters are silently dropped, and it sets the " - "D2H bar directly. Measured at 9×9: the maximum is 1 633, and the lossless " - "cap of 1 700 already makes D2H the bottleneck on the f32 build."), + "Must exceed the real maximum or clusters are silently dropped. At 9×9 the " + "maximum is 1 633, and a cap of 1 700 already makes D2H the bottleneck."), ("batch size", "Frames per find_clusters_batched call.", "2 000 amortises launch overhead without a large pinned footprint."), ("cluster_size", "Compile-time stencil geometry.", - "3×3 and 9×9 are registered; 9×9 moves the bottleneck off H2D and onto the " - "kernel on f64, and onto D2H once opt7 shortens it."), + "3×3 and 9×9 are registered. 9×9 moves the bottleneck off H2D."), ("register_input_buffer", "Page-locks the host array for DMA.", - "Always, if the data is already in RAM. Check the pinning budget first."), + "Always, if the data is already in RAM."), ] flagged = None for i, (p_, what, guide) in enumerate(params): if i % 2 == 0: - rect(s, M, y, 11.9, 0.82, PANEL) - tf = tb(s, M + 0.28, y + 0.14, 2.6, 0.5) - run(para(tf, True, line=1.1), p_, 9.5, ACCENT, font=MONO, bold=True) - tf = tb(s, M + 3.0, y + 0.14, 2.9, 0.6) - 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) + rect(s, M, y, 12.30, 0.82, PANEL) + tf = tb(s, M + 0.29, y + 0.16, 2.69, 0.5) + run(para(tf, True, line=1.1), p_, 11, ACCENT, font=MONO, bold=True) + tf = tb(s, M + 3.10, y + 0.16, 3.00, 0.6) + run(para(tf, True, line=1.2), what, 11, PALE) + tf = tb(s, M + 6.36, y + 0.16, 5.58, 0.6) + run(para(tf, True, line=1.2), guide, 11, TEXT2) if p_ == "max_clusters_per_frame": flagged = y - y += 0.80 + y += 0.86 # 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=RED) +frame_rect(s, M - 0.04, flagged - 0.05, 12.38, 0.96) # =========================================================== 28 · NEXT s = new_slide() chrome(s, 36, "Where this leaves us", "The bottleneck has walked from the host, to the GPU, to PCIe") +# Four conclusions, not four paragraphs: each card states the fact and the +# number that carries it. The reasoning behind each is in the notes. cards = [ ("DONE", ACCENT, "×9.1 at 3×3, ×26.5 at 9×9", - "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."), + "16.3 and 25.1 µs/frame end to end, both on their hardware floor. At 3×3 " + "that is 45× MÖNCH03's standard frame rate."), ("DONE", ACCENT, "FP32 pedestal, safely", - "−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."), + "−40.5 % kernel time, and correct: the variance accumulates on a frozen " + "per-pixel offset. 6 duplicate clusters 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]: " - "2 000 separate 320 kB descriptors, plus 26% of H2D↔D2H contention."), + "H2D sustains 16.31 µs against 13.15 uncontended — 2 000 separate 320 kB " + "descriptors, plus H2D↔D2H contention."), ("NEXT", AMBER, "9×9: the result path, not the kernel", - "At a lossless cap D2H already binds [f32 · s4]: 25.24 µs against a 23.94 µs " - "kernel. More " - "kernel work buys nothing until the D2H slot stops being cap-sized."), + "At a lossless cap D2H already binds: 25.24 µs against a 23.94 µs kernel. " + "More kernel work buys nothing."), ] -for i, (tag, col, title, body) in enumerate(cards): - cx = M + (i % 2) * 6.05 - cy = 2.05 + (i // 2) * 2.35 - rect(s, cx, cy, 5.85, 2.05, PANEL) - rect(s, cx, cy, 5.85, 0.035, col) - tf = tb(s, cx + 0.3, cy + 0.26, 1.4, 0.26) - run(para(tf, True), tag, 8.5, col, bold=True, spc=1.5) - tf = tb(s, cx + 0.3, cy + 0.60, 5.2, 0.4) - run(para(tf, True, line=1.1), title, 14, PALE, bold=True) - tf = tb(s, cx + 0.3, cy + 1.12, 5.2, 0.85) - run(para(tf, True, line=1.3), body, 10, TEXT2) -callout(s, M, 6.58, 11.2, +callout(s, M, 1.78, 12.30, "**Every number in this deck is measured on one machine, one dataset, one " "git revision** — and the annexes carry the reconciliation for each of them.", - h=0.60, size=10) + h=0.62, size=PT_LEAD) +for i, (tag, col, title, body) in enumerate(cards): + cx = M + (i % 2) * 6.25 + cy = 2.62 + (i // 2) * 2.30 + rect(s, cx, cy, 6.04, 2.00, PANEL) + rect(s, cx, cy, 6.04, 0.035, col) + tf = tb(s, cx + 0.3, cy + 0.26, 1.45, 0.26) + run(para(tf, True), tag, 9.5, col, bold=True, spc=1.5) + tf = tb(s, cx + 0.3, cy + 0.58, 5.37, 0.4) + run(para(tf, True, line=1.1), title, 15, PALE, bold=True) + tf = tb(s, cx + 0.3, cy + 1.08, 5.37, 0.85) + run(para(tf, True, line=1.3), body, 12, TEXT2) notes(s, """The headroom claim, stated carefully. WHERE THE MATERIAL IS, if anyone asks for it after the talk rather than from the @@ -2316,42 +2219,36 @@ section("", s = new_slide() annex_chrome(s, 1, "measurement convention · expands slide 21", "Uncontended, or as the pipeline runs it") -bullets(s, M, 1.90, 12.0, [ +bullets(s, M, 1.82, 12.40, [ "Every engine time in this deck is tagged **[build · s1|s4]**. **s1** is one " "stream with nothing else running: what an engine does **on its own**, which is " "the right number for a capability claim and for the headroom that remains. " "**s4** is the shipped four-stream pipeline: each engine's **busy time per " "frame**, the union of its intervals, the only number that can set a floor.", -], size=10.5) -table(s, M, 2.78, 5.8, +], size=PT_BODY) +table(s, M, 3.07, 5.99, ["3×3 · µs/frame", "s1 f64", "s1 f32", "s4 f64", "s4 f32"], [["H2D", "13.14", "13.15", "16.17", "16.63"], ["kernel", "14.72", "4.32", "15.17", "5.53"], ["D2H", "5.31", "5.27", "7.69", "7.57"], ["engine max [s4]", "—", "—", "16.17", "16.63"], ["FLOOR = lower of max, sustained", "—", "—", "**16.17**", "**16.31**"]], - colw=[0.28, 0.18, 0.18, 0.18, 0.18], size=9, rowh=0.44) -table(s, 7.0, 2.78, 5.8, + colw=[0.28, 0.18, 0.18, 0.18, 0.18], size=PT_TABLE, rowh=0.44) +table(s, 7.01, 3.07, 5.99, ["9×9 · cap 1700", "s1 f64", "s1 f32", "s4 f64", "s4 f32"], [["H2D", "13.20", "13.22", "20.77", "20.54"], ["kernel", "39.86", "23.70", "32.66", "23.94"], ["D2H", "21.97", "21.95", "25.25", "25.24"], ["engine max [s4]", "—", "—", "32.66", "25.24"], ["FLOOR = lower of max, sustained", "—", "—", "**30.01**", "**25.14**"]], - colw=[0.28, 0.18, 0.18, 0.18, 0.18], size=9, rowh=0.44) -callout(s, M, 5.66, 12.0, - "**Two traps this table closes.** (1) s4 is engine *occupancy*, not duration: " - "the 9×9 kernel row **falls** 39.86 → 32.66 while every transfer rises, and " - "nothing that measures a duration falls under load. (2) **the FLOOR is not the " - "engine max**: the max is profiled and runs 2–8 % high, so the floor is whichever is " - "lower, it or the best unprofiled sustained rate, which is why opt6 reports " - "**30.01 µs against a 32.66 max**.", h=0.92, size=10.5) -caption(s, M, 6.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, so any " - "claim about which engine binds must be read from the s4 columns. Bandwidth " - "arithmetic in the notes.") + colw=[0.28, 0.18, 0.18, 0.18, 0.18], size=PT_TABLE, rowh=0.44) +callout(s, M, 6.00, 12.40, + "**Two traps.** (1) s4 is engine *occupancy*, not duration — the 9×9 kernel " + "row **falls** under load. (2) the **floor is not the engine max**: the max " + "is profiled and runs 2–8 % high.", h=0.92, size=PT_LEAD) +caption(s, M, 7.00, 12.40, + "The D2H shift is an s4 phenomenon: at s1 the kernel binds in both arms. " + "Which engine binds must be read from the s4 columns.") 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 @@ -2371,14 +2268,14 @@ exists. A probe roofline is an estimate, never a hard denominator.""") 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=3) -bullets(s, M, 1.95, 12.0, [ +bullets(s, M, 1.95, 12.40, [ "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.", "A **CUDA Graph** captures the whole dependency DAG once. Replaying it is a " "**single** cudaGraphLaunch: the driver already knows every node and edge.", -], size=10.5) -figure(s, "fig_graphs", M, 3.15, 7.6) -code(s, 8.5, 3.15, 4.1, [ +], size=PT_BODY) +figure(s, "fig_graphs", M, 3.15, 7.85) +code(s, 8.56, 3.15, 4.24, [ "// record once, at setup", "cudaStreamBeginCapture(sc.stream, ...);", " submit_h2d_kernel_d2h(sc);", @@ -2387,42 +2284,37 @@ 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, 5.06, 4.1, +], size=PT_CODE, title="ClusterFinderCUDA_graph.hpp") +callout(s, 8.56, 5.06, 4.24, "**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, - "Its 3×3 edge was never established: the graph finder recorded no CUDA events " - "while the stream finder did, and that instrumentation tax (2.8 µs) is larger " - "than the gap (0.8 µs). More decisively, it never received the chunked pipeline " - "of opt5, so it is competing on ~2 µs of launch cost against a 24 µs floor. " - "Launch overhead stops binding one step later; the technique aimed at it can no " - "longer pay.") + h=1.10, size=PT_LEAD, color=AMBER) +caption(s, M, 6.62, 12.40, + "The 3×3 edge was never established: the instrumentation tax (2.8 µs) is" + "larger than the gap (0.8 µs). Detail in the notes.") # ---- 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=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** " - "cudaGraphLaunch, so the ceiling on what it can save is **¾ of the submission cost**.", -], size=10.5) -table(s, M, 2.72, COL, + "The stream path issues **four runtime calls per frame** — memset, H2D, " + "launch, D2H — and a graph replaces all four with **one** cudaGraphLaunch.", +], size=PT_BODY) +table(s, M, 2.80, COL, ["call", "per frame", "host cost", "µs/frame"], [["cudaMemcpyAsync", "2", "1.98 µs", "3.97"], ["cudaLaunchKernel", "1", "2.13 µs", "2.13"], ["cudaMemsetAsync", "1", "1.57 µs", "1.57"], ["**submission total**", "**4**", "", "**7.67**"]], - colw=[0.40, 0.18, 0.22, 0.20], size=9, rowh=0.44) -code(s, M, 5.42, COL, [ + colw=[0.40, 0.18, 0.22, 0.20], size=PT_TABLE, rowh=0.44) +code(s, M, 5.40, 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") +], size=PT_CODE, title="THE ARITHMETIC") 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) + "the GPU entirely.", h=0.80, size=PT_LEAD, color=AMBER) rail(s, [ ("label", "route A · measured verdict"), ("gap", 0.10), @@ -2456,68 +2348,63 @@ rows = [ "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 +y = 2.80 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) + rect(s, M, y, 12.30, 2.05, PANEL) + rect(s, M, y, 0.04, 2.05, AMBER) + tf = tb(s, M + 0.31, y + 0.22, 1.03, 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) + tf = tb(s, M + 1.34, y + 0.20, 6.20, 0.3) run(para(tf, True), name, 14, PALE, bold=True) - tf = tb(s, M + 1.30, y + 0.52, 6.0, 0.3) + tf = tb(s, M + 1.34, y + 0.52, 6.20, 0.3) run(para(tf, True), how, 10, MUTED, font=MONO) - tf = tb(s, M + 1.30, y + 0.90, 10.2, 1.0) + tf = tb(s, M + 1.34, y + 0.90, 10.54, 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) + tf = tb(s, M + 1.34, y + 1.68, 10.54, 0.3) run(para(tf, True), verdict, 10.5, AMBER, bold=True) y += 2.25 -callout(s, M, 6.38, 11.9, +callout(s, M, 1.82, 12.30, "**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) + "saying so, to stop the experiment being repeated.", h=0.80, size=PT_LEAD) # ---- A4 · THE OPT5 CODE -------------------------------------------------- s = new_slide() annex_chrome(s, 3, "opt5 · the overlap code · expands slide 20", "The overlap, in six lines, and why you never write them") -code(s, M, 1.86, 7.15, [ +code(s, M, 1.82, 7.39, [ "tok = cf.«submit_batch»(data[a0:b0], first_frame=a0)", "for a, b in bounds[1:]:", " nxt = cf.«submit_batch»(data[a:b], first_frame=a) # GPU starts N+1 …", " results.extend(cf.«collect»(tok)) # … host unpacks N", " tok = nxt", "results.extend(cf.«collect»(tok)) # drain the last one", -], size=8.5, title="THE WHOLE OF OPT5") -bullets(s, M, 3.66, 7.15, [ - "**You do not write this.** It is inside find_clusters_batched(), which chunks " - "the batch and runs the loop for you, opt5 arrived as a **speedup, not an API " - "change**, and every existing caller got it without editing a line.", - "submit_batch() and collect() stay public for anyone who wants the token by " - "hand, streaming from a detector, interleaving other work between chunks.", -], size=10) +], size=PT_CODE, title="THE WHOLE OF OPT5") +bullets(s, M, 3.64, 7.39, [ + "**You do not write this.** It is inside find_clusters_batched(), so opt5 " + "arrived as a **speedup, not an API change**.", + "submit_batch() and collect() stay public for streaming from a detector, or " + "interleaving other work between chunks.", +], size=PT_BODY) rail(s, [ ("label", "chunk sizing · the two constraints"), ("gap", 0.12), ("row", "multiple of", "n_streams", ACCENT), ("row", "capped at", "MAX_SLOT_BYTES", ACCENT), ("gap", 0.16), - ("note", "The chunk MUST be a multiple of n_streams. The device pedestal is " - "per-stream and advances once per frame the stream sees, so an uneven " - "chunk leaves the four pedestals at different ages, the finder would " - "stop being reproducible, and two runs of the same data would not agree."), + ("note", "A multiple of n_streams, because the device pedestal is per-stream: " + "an uneven chunk leaves the four pedestals at different ages and the " + "finder stops being reproducible."), ("gap", 0.14), - ("note", "It is also capped so the two pinned output slots stay bounded: the " - "slot is chunk × (4 + cap × sizeof(Cluster)), which at 9×9 and cap 1 700 " - "is 544.5 kB per frame. chunk_size_for(n) applies both rules; pass the " - "result to reserve_output_slots() to pre-pay the page-locking outside " - "any timed region."), + ("note", "Capped so the two pinned slots stay bounded — 544.5 kB per frame at " + "9×9, cap 1 700. chunk_size_for(n) applies both rules."), ]) -callout(s, M, 5.92, 7.15, +callout(s, M, 5.02, 7.39, "Two chunks in flight is enough. A third adds pinned memory and no overlap: " - "the host is already busy for the whole time the GPU is.", h=0.72, size=10) -caption(s, M, 6.80, 12.0, + "the host is already busy for the whole time the GPU is.", h=0.72, size=PT_LEAD) +caption(s, M, 5.98, 7.39, "The saving is min(GPU, host) per chunk, so opt5 pays most when the two terms " "are comparable, ×1.31 at 3×3, and least when one dominates, ×1.20 at 9×9, " "where the host term is roughly twice the GPU term. That gap is the diagnosis " @@ -2527,13 +2414,12 @@ caption(s, M, 6.80, 12.0, s = new_slide() annex_chrome(s, 4, "the fault model · expands slide 24", "The fault model, tested against every step") -bullets(s, M, 1.90, 12.0, [ - "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, +bullets(s, M, 1.82, 12.40, [ + "The **0.68 µs per fault** rate is fitted on the **3×3 f32** ladder. Every " + "row below is **9×9, f64 vs f32**, so the rate is applied **out of sample**, " + "never refitted.", +], size=PT_BODY) +table(s, M, 2.55, 12.40, ["step", "f64 warm (faults)", "f32 warm (faults)", "Δ wall", "Δ faults", "predicted", "verdict"], [["opt3", "82.44 µs (128 k)", "95.66 µs (521 k)", "**+13.22**", "+393 k", @@ -2543,20 +2429,15 @@ table(s, M, 2.78, 12.0, ["opt5", "66.39 µs (152 k)", "61.85 µs (10 k)", "−4.54", "−141 k", "**−4.81**", "**not separable**"], ["opt6", "30.01 µs (0)", "25.14 µs (0)", "−4.87", "0", "0.00", "clean"]], - colw=[0.08, 0.18, 0.18, 0.10, 0.10, 0.13, 0.23], size=8.5, rowh=0.50) -callout(s, M, 5.42, 12.0, - "**opt3 is the whole argument in one row.** A 393 k fault gap predicts +13.37 µs; " - "+13.22 was observed, agreement to **1 %**. The −40 % kernel is in there " - "somewhere, invisible under 13 µs of the OS zeroing pages. Only opt4 and opt6, " - "where the fault term is ~0, report the typedef at all.", h=0.90, size=10.5) -caption(s, M, 6.52, 12.0, - "the 9×9 ladder at cap 1 700, warm = best of reps " - "1–4, faults are that rep's own getrusage minor-fault count. Predicted = Δfaults " - "× 0.68 µs ÷ 20 000 frames, with 0.68 carried in unchanged from the 3×3 fit " - "(slide 28): 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.") + colw=[0.08, 0.18, 0.18, 0.10, 0.10, 0.13, 0.23], size=PT_TABLE, rowh=0.50) +callout(s, M, 5.24, 12.40, + "**opt3 is the whole argument in one row.** A 393 k fault gap predicts " + "**+13.37 µs**; **+13.22** was observed. The −40 % kernel is in there, " + "invisible under 13 µs of the OS zeroing pages.", h=0.9, size=PT_LEAD) +caption(s, M, 6.29, 12.40, + "9×9 · cap 1 700 · warm = best of reps 1–4. The 0.68 µs/fault rate is " + "carried in unchanged from the 3×3 fit: nothing here is tuned to make the " + "columns agree.") # ---- A5 · THE VARIANCE REWRITE IN FULL ----------------------------------- # Was main-arc slide 24, plus the error-floor panel that used to share @@ -2565,38 +2446,31 @@ caption(s, M, 6.52, 12.0, s = new_slide() annex_chrome(s, 5, "the variance rewrite · expands slide 23", "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, [ +bullets(s, M, 1.82, 7.65, [ + "Freeze **X₀ = round(mean)** once, at the end of pedestal training, and " + "never move it.", + "Accumulate the **centred** Y = X − X₀; report **X₀ + sum/n**. Both " + "operands are now **O(rms)-sized**.", +], size=PT_BODY) +code(s, M, 3.20, 7.65, [ "// 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, +], size=PT_CODE, title="clusterfinder_kernel.cuh") +callout(s, M, 5.02, 7.65, "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. " - "The naive f32 " - "build's +28.06 % excess is this shaded band, integrated over the sensor.", - size=9) + "70 clusters out of 233 million.", h=0.72, size=PT_LEAD) +callout(s, M, 5.88, 7.65, + "**X₀ must never be updated** — the accumulators are defined relative to " + "it. Welford's is the other correct answer; this one is two lines and " + "costs nothing.", h=0.86, size=PT_LEAD, color=AMBER) +h = figure(s, "fig_varfloor", 8.35, 1.82, 4.46) +caption(s, 8.35, 1.82 + h + 0.18, 4.46, + "Which pixels the ±3 ADU² floor actually reaches: the quiet ones, whose " + "true variance is smallest.", size=PT_META) # ---- A6·1 · THE THREE ARTEFACTS ------------------------------------------ # Main-arc slide 28 keeps only the first of these three, because it is the only @@ -2619,36 +2493,36 @@ items = [ ] 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) + rect(s, x, 2.0, 3.96, 3.15, PANEL) + rect(s, x, 2.0, 3.96, 0.035, col) + tf = tb(s, x + 0.26, 2.28, 3.41, 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) + tf = tb(s, x + 0.26, 3.02, 3.41, 1.5) run(para(tf, True, line=1.3), body, 10, TEXT2) - tf = tb(s, x + 0.26, 4.42, 3.3, 0.65) + tf = tb(s, x + 0.26, 4.42, 3.41, 0.65) run(para(tf, True, line=1.25), fix, 10, ACCENT) - x += 4.03 -code(s, M, 5.4, 11.9, [ + x += 4.16 +code(s, M, 5.32, 12.30, [ "# bracket every timed cell:", "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 · BRACKET EVERY TIMED CELL") -callout(s, M, 6.68, 11.9, +], size=PT_CODE, title="THE FAULT PROTOCOL · BRACKET EVERY TIMED CELL") +callout(s, M, 6.64, 12.30, "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) + "never the variable.", h=0.62, size=PT_LEAD) # ---- A6 · FAULTS --------------------------------------------------------- s = new_slide() annex_chrome(s, 6, "benchmark artefacts · expands slide 28", "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. " - "No disk I/O: ru_majflt stays 0 all campaign. At 4 kB/page, **1 GB = 262 144 faults**.", -], size=10.5) -table(s, M, 2.62, 12.0, +bullets(s, M, 1.9, 12.40, [ + "A page exists in the address space but has no physical frame yet. First " + "touch makes the kernel find one, **zero it** (mandatory), and map it. No " + "disk: ru_majflt stays 0 all campaign. At 4 kB/page, **1 GB = 262 144 faults**.", +], size=PT_BODY) +table(s, M, 2.62, 12.40, ["", "(a) result heap", "(b) pinned D2H slots"], [["allocator", "malloc → mmap, one ClusterVector per frame", "cudaMallocHost, in submit_batch"], @@ -2657,46 +2531,42 @@ table(s, M, 2.62, 12.0, "**no**: once per buffer, for its lifetime"], ["removed by", "**collect_view()**: it allocates nothing", "nothing; reserve_output_slots() only moves it out of the timer"]], - colw=[0.14, 0.43, 0.43], size=9) -callout(s, M, 5.62, 5.85, + colw=[0.14, 0.43, 0.43], size=PT_TABLE, rowh=0.50) +callout(s, M, 5.30, 6.04, "**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.42, 5.85, [ + "count from run 0 and changes nothing else.", h=0.72, size=PT_LEAD) +code(s, 6.75, 5.30, 6.04, [ "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.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 " - "independent argument for opt6.", h=0.72, size=10, color=AMBER) +], size=PT_CODE, title="NOT A CORRELATION, AN IDENTITY") +callout(s, M, 6.5, 12.30, + "**At 9×9 the heap never plateaus.** ~9.3 GB per pass exceeds glibc's mmap " + "threshold, so it is re-faulted every pass: **~10 µs/frame, permanently**.", h=0.72, size=PT_LEAD, color=AMBER) # ---- A7 · EVENTS --------------------------------------------------------- s = new_slide() annex_chrome(s, 6, "benchmark artefacts", "CUDA events measure the stream, not the kernel", part=3, nparts=4) -bullets(s, M, 1.90, COL, [ +bullets(s, M, 1.9, COL, [ "avg_kernel_time_ms() brackets the launch with **cudaEventRecord on the " - "kernel's own stream**. What it returns is elapsed time on **that stream's " - "timeline**, which includes time spent **queued behind other streams**.", - "So it is honest at 1 stream and inflates under saturation: up to **3.5×** at " - "8 streams. The tell is that the derived *PCIe + overhead* = wall/N − kernel_ms " - "**goes negative**: kernels overlap, so wall/frame < kernel/frame.", - "It is also not free: **~3.6 µs/frame**, i.e. 10–15 % of end-to-end throughput " - "at 3×3, to produce a number that is unusable exactly when it matters.", -], size=10.5) + "kernel's own stream**, so it returns time on **that stream's timeline** — " + "including time queued behind other streams.", + "Honest at 1 stream, inflated up to **3.5×** at 8. The tell: *PCIe + " + "overhead* = wall/N − kernel_ms **goes negative**.", + "It is not free either: **~3.6 µs/frame**, 10–15 % of throughput at 3×3.", +], size=PT_BODY) code(s, M, 4.75, COL, [ "if (m_time_kernels) // OFF by default, all 3 finders", " cudaEventRecord(start[slot][i], sc.stream);", "find_clusters_in_single_frame<<>>(...);", "if (m_time_kernels)", " cudaEventRecord(stop[slot][i], sc.stream); // <- queue-wait lands here", -], size=8, title="ClusterFinderCUDA.hpp") +], size=PT_CODE, title="ClusterFinderCUDA.hpp") 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) + h=0.72, size=PT_LEAD) rail(s, [ ("label", "9×9 kernel · f64 · nsys"), ("gap", 0.10), @@ -2714,13 +2584,12 @@ rail(s, [ s = new_slide() 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 " - "stamps its own start/end and the host reads those records **afterwards**, so " - "nothing is injected into the execution path.", -], size=10.5) -table(s, M, 2.60, 12.0, +bullets(s, M, 1.82, 12.40, [ + "Tracing a **host** call runs a callback on entry and exit, **inside the " + "interval being measured**. GPU work is stamped by the hardware and read " + "back **afterwards**, so nothing is injected into the execution path.", +], size=PT_BODY) +table(s, M, 2.55, 12.40, ["measurement", "sqlite table", "9×9 [f64 · s1 · cap 1500]", "verdict"], [["cudaLaunchKernel: the host call", "CUPTI_..._RUNTIME", "1.85 µs", "**inflated ~4×**"], @@ -2729,19 +2598,15 @@ table(s, M, 2.60, 12.0, "**inflated ~4×**"], ["the H2D / D2H transfer", "CUPTI_..._MEMCPY", "13.25 / 19.44 µs", "sound to ~2 %"]], - colw=[0.36, 0.24, 0.22, 0.18], size=9, rowh=0.52) -callout(s, M, 5.30, 11.9, + colw=[0.36, 0.24, 0.22, 0.18], size=PT_TABLE, rowh=0.52) +callout(s, M, 5.32, 12.30, "**Same cudaMemcpyAsync, two numbers: 1.65 µs to ask for the copy, 13.25 µs " - "for the copy to happen.** If the profiler must be present at the moment to " - "measure it, it distorts it; if the hardware records it anyway and the " - "profiler reads it later, it does not.", h=0.86, size=10.5) -caption(s, M, 6.40, 11.9, - "Every headline number in this deck (kernel times, transfer times, duty cycles, " - "overlap and every engine floor) comes from the _KERNEL and _MEMCPY tables (see " - "GPU-side timestamps). 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.") + "for the copy to happen.** A profiler that must be present to measure," + "distorts; one that reads a record later does not.", h=0.86, size=PT_LEAD) +caption(s, M, 6.33, 12.30, + "Every engine number in this deck comes from GPU-side timestamps. Proof " + "they are sound: opt7 sustains 25.14 µs unprofiled against a 25.24 µs" + "profiled estimate, 0.4 % apart.") 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 @@ -2766,12 +2631,12 @@ against a 24 us floor", and nothing finer than that.""") s = new_slide() annex_chrome(s, 7, "the three tests · completes slide 4", "There is a third test", part=1, nparts=2) -bullets(s, M, 1.80, 11.9, [ +bullets(s, M, 1.80, 12.30, [ "Slide 4 showed the decision with **two** tests, which is the right " "simplification for the arc. The finder has **three**, and the missing one " "is what the next slide is about.", -], size=10.5) -code(s, M, 2.48, 7.3, [ +], size=PT_BODY) +code(s, M, 2.48, 7.54, [ "v = frame[i] - ped_mean[i]; rms = ped_rms[i]", "if (v < -nSigma*rms) -> skip, no update", "m = «max» over the 3x3 window", @@ -2782,7 +2647,7 @@ code(s, M, 2.48, 7.3, [ "else if (total > «c3*nSigma»*rms) // TEST 3", " if (v == m) -> emit cluster", "else -> update pedestal", -], size=9, title="ClusterFinder.hpp:104-142 · all three branches") +], size=PT_CODE, title="ClusterFinder.hpp:104-142 · all three branches") rail(s, [ ("label", "what each test asks"), ("gap", 0.12), @@ -2795,19 +2660,14 @@ rail(s, [ ("note", "Test 1 reads the MAX of the window; Test 3 reads its SUM. That " "one difference is the whole of the next slide."), ], y0=2.15) -callout(s, M, 5.86, 11.9, - "**c3 = √(3×3) = 3 is not a fudge: it falls out of variance addition.** " - "For independent samples the variance of a sum is the sum of the " - "variances, so a 3×3 window of pixels each carrying noise σ gives " - "**Var(Σ v) = 9σ²**, i.e. a sum whose noise is **√9 · σ = 3σ**. Requiring " - "that sum to clear nSigma of ITS OWN noise is exactly **c3·nSigma·rms**, " - "the **same 5σ criterion as Test 1, asked of the window instead of the " - "pixel**. So Test 3 catches a photon whose charge is shared out so widely " - "that no single pixel reaches 5σ, but the nine together do.", - h=1.10, size=10.5) -caption(s, M, 7.02, 11.9, +callout(s, M, 5.86, 12.30, + "**c3 = √(3×3) = 3 falls out of variance addition.** Variances add, " + "standard deviations do not: nine pixels of noise σ give **Var(Σ v) = " + "9σ²**, so the sum's noise is **3σ**. **c3·nSigma·rms** is therefore the" + "*same 5σ criterion as Test 1*, asked of the window instead of the pixel.", h=1.1, size=PT_LEAD) +caption(s, M, 7.02, 12.30, "Roughly 80 % of pixels reach the last branch and push the pedestal; " - "~1.5 % are peaks and ~18 % sit in a peak's shadow.", size=9) + "~1.5 % are peaks and ~18 % sit in a peak's shadow.", size=PT_META) notes(s, """Why slide 4 leaves Test3 out: at 3x3 it is a small correction to the cluster count, and the arc's point there is the THREE OUTCOMES shape -- store, shadow, update -- not completeness. This slide is where completeness belongs. @@ -2847,21 +2707,21 @@ the SUM. That is the whole difference in sensitivity.""") s = new_slide() annex_chrome(s, 7, "which test causes the CPU/CPU gap · expands slide 31", "Test3 makes the clusters; Test1 only moves the pedestal", - part=2, nparts=2, title_size=25) -bullets(s, M, 1.76, 11.9, [ + part=2, nparts=2, title_size=PT_LEAD) +bullets(s, M, 1.76, 12.30, [ "**push_fast** touches only the pixel's **own** accumulators and never reads " "the stencil, so the update is **order-independent**: same set of updated " "pixels, bit-identical pedestal. Only a differing **decision** can diverge.", -], size=10.5) -figure(s, "fig_test3", M + 0.70, 2.34, 10.45) -callout(s, M, 6.22, 11.9, +], size=PT_BODY) +figure(s, "fig_test3", M + 0.72, 2.49, 10.80) +callout(s, M, 6.13, 12.30, "Measured two ways. **Instrumented**: of the 19 disagreeing clusters, the " "**11 that only frozen finds are all Test 3**; the 8 only serial finds are " "all the local-max gate. **Ablated**: compile Test 3 out and the 11 go to " - "**zero**.", h=0.72, size=10.5) -caption(s, M, 7.06, 11.9, + "**zero**.", h=0.72, size=PT_LEAD) +caption(s, M, 7.05, 12.30, "3×3 · 10 000 frames · 23 244 605 clusters · per-pixel branch codes " - "diffed frame by frame, then re-run with Test 3 compiled out.", size=9) + "diffed frame by frame, then re-run with Test 3 compiled out.", size=PT_META) notes(s, """TWO EXPERIMENTS, INDEPENDENT ROUTES. Instrumented (both finders shipped logic, branch codes recorded): 974 pixels out diff --git a/docs/deck/deckgate.py b/docs/deck/deckgate.py new file mode 100644 index 00000000..3124a4f8 --- /dev/null +++ b/docs/deck/deckgate.py @@ -0,0 +1,100 @@ +"""The projection floor, and the one thing both figure generators must share. + +A figure's on-screen type size is `raw_pt x (placement_width / figure_width)`, +and neither factor is visible where the font size is written. This closes that +loop by reading the placement width out of the deck script itself, so the gate +cannot drift from the layout it checks. + +It lives in its own module because there are TWO generators -- make_figs.py and +make_figs_kernel.py -- and for a long time only the first was checked. The +kernel figures were consequently set at 7-8 pt and nobody found out until they +were projected. +""" +import re +from pathlib import Path + +import matplotlib + +# The floor, in points, on the projected slide. Set for a room where the back +# row is 6-7 m away: on a 13.33 x 7.5 in slide, 10 pt is two thirds of the 15 pt +# body size, which is the usual lower bound for supporting type. Below that a +# label inside a plot reads as a footnote rather than part of the argument. +MIN_EFF_PT = 10.0 + +# A string may opt out by carrying gid="texture". That is for marks nobody is +# asked to READ -- a value printed into every cell of a pixel map, where the +# pattern is the message and the digits are shading. It is set per Text object, +# never per figure, so an exemption stays a decision about one label. +EXEMPT_GID = "texture" + +DECK = Path(__file__).resolve().parent / "build_performance_deck.py" +VIOLATIONS = [] + + +def placements(): + """{figure name: the narrowest width the deck places it at}. + + The token values must track build_performance_deck.py. They did not when the + margins were narrowed, and the gate then silently checked widths no figure + was ever placed at. + """ + if not DECK.exists(): + return {} + env = {"M": 0.50, "COL": 8.16, "RAIL_W": 3.62, "RAIL_X": 9.28, + "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 check(fig, path, name, dpi, place_w=None): + """Report every string in `fig` that will project below the floor. + + Call it BEFORE closing the figure and AFTER saving it: the saved width is + what sets the scale, and `bbox_inches="tight"` means that width is not the + figsize. + """ + from PIL import Image + + texts = [(t.get_text(), t.get_fontsize()) + for t in fig.findobj(matplotlib.text.Text) + if t.get_text().strip() and t.get_visible() + and t.get_gid() != EXEMPT_GID] + if place_w is None: + place_w = PLACE_W.get(name) + if place_w is None: + print("wrote", name) + return + 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 " + f"{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 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}") diff --git a/docs/deck/make_figs.py b/docs/deck/make_figs.py index 3d8a2eb4..b149c1d1 100644 --- a/docs/deck/make_figs.py +++ b/docs/deck/make_figs.py @@ -56,6 +56,7 @@ 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 sys import json import re import matplotlib @@ -66,6 +67,9 @@ import numpy as np from matplotlib.patches import FancyArrowPatch, FancyBboxPatch, Rectangle from pathlib import Path +sys.path.insert(0, str(Path(__file__).resolve().parent)) +import deckgate # noqa: E402 -- the projection floor, shared with make_figs_kernel.py + OUT = Path(__file__).resolve().parent.parent / "figures" OUT.mkdir(exist_ok=True) @@ -114,82 +118,30 @@ DPI = 220 # 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_performance_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() +# Raised from 9.0 when the body text went to 15 pt: 10 pt is two thirds of the +# body size, which is the usual floor for supporting type. Below that a label +# inside a plot reads as a footnote rather than part of the argument. +# The floor, the placement table and the report all live in deckgate.py, because +# there are two figure generators and for a long time only this one was checked. +MIN_EFF_PT = deckgate.MIN_EFF_PT +PLACE_W = deckgate.PLACE_W 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_performance_deck.py, and tools/audit prints both sides. + `place_w` overrides the width parsed out of the deck; it is only needed for a + figure the deck places through something the parser cannot evaluate. """ 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) + deckgate.check(fig, path, name, DPI, place_w) plt.close(fig) - 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}") + deckgate.report() def bare(ax, keep=("left", "bottom")): @@ -225,11 +177,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=9, labelcolor=TEXT2, loc="upper right") - ax.set_ylabel("clusters / bin", fontsize=9) + ax.legend(frameon=False, fontsize=10.6, labelcolor=TEXT2, loc="upper right") + ax.set_ylabel("clusters / bin", fontsize=10.6) bare(ax, keep=("left", "bottom")) ax.set_title("cluster energy spectrum · 3×3, 10 000 frames, 23.2 M clusters", - color=MUTED, fontsize=9, loc="left", pad=6) + color=MUTED, fontsize=10.6, loc="left", pad=6) m = h["cpu"] > 0 axr.axhspan(0.999, 1.001, color=GREEN, alpha=0.18, zorder=1) @@ -239,13 +191,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=9) - axr.set_xlabel("cluster sum [ADU]", fontsize=9) - axr.set_ylabel("vs CPU", fontsize=9) + axr.set_yticklabels(["−0.4 %", "0", "+0.4 %"], fontsize=10.6) + axr.set_xlabel("cluster sum [ADU]", fontsize=10.6) + axr.set_ylabel("vs CPU", fontsize=10.6) 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=9) + fontsize=10.6) fig.subplots_adjust(hspace=0.10) save(fig, "fig_spectra_valid") @@ -283,30 +235,30 @@ def fig_first_run(): # engine occupancy is 60 140 FPS, 1.9 % lower, and corroborates it. ax.axhline(61312, color=GREEN, lw=1.2, ls="--", zorder=4) ax.text(2.6, 62100, "H2D floor · 61 312 FPS (nsys estimates 60 140)", - color=GREEN, fontsize=8.5, ha="center", va="bottom") + color=GREEN, fontsize=10.6, ha="center", va="bottom") for xi, (a, b, f) in enumerate(zip(warm, cold, faults)): ax.text(xi - w / 2, a + 900, f"{a:,}", ha="center", va="bottom", - color=TEXT2, fontsize=9) + color=TEXT2, fontsize=10.6) ax.text(xi + w / 2, b + 900, f"{b:,}", ha="center", va="bottom", - color=PALE, fontsize=9.5, fontweight="bold") + color=PALE, fontsize=11.2, fontweight="bold") drop = 100 * (1 - b / a) big = drop > 10 ax.text(xi, -2600, f"{f:,} fault" + ("" if f == 1 else "s"), ha="center", va="top", color=AMBER if big else MUTED, - fontsize=8, fontweight="bold" if big else "normal") + fontsize=10.4, fontweight="bold" if big else "normal") ax.text(xi, -6300, ("−%.0f %%" % drop) if drop >= 1 else "—", ha="center", va="top", color=AMBER if big else MUTED, - fontsize=10 if big else 8.5, fontweight="bold") + fontsize=11.8 if big else 10.6, fontweight="bold") # the two steps that pay nothing, annotated just above their own bars for xi, yi, why in [(0, 21500, "discards every frame\nas it goes"), - (5, 65200, "allocates nothing\nat all")]: + (5, 67500, "allocates nothing\nat all")]: ax.text(xi, yi, why, ha="center", va="bottom", color=GREEN, - fontsize=8, linespacing=1.35) + fontsize=10.4, linespacing=1.35) ax.set_xticks(x) - ax.set_xticklabels(steps, fontsize=8.5, color=TEXT2) + ax.set_xticklabels(steps, fontsize=10.6, color=TEXT2) ax.set_ylim(0, TOP) ax.set_yticks([]) ax.tick_params(axis="x", pad=34) @@ -317,11 +269,11 @@ def fig_first_run(): Rectangle((0, 0), 1, 1, color=AMBER)] ax.legend(handles, ["achievable · warm, 5-rep campaign", "first run · results retained, one process"], - frameon=False, fontsize=9, labelcolor=TEXT2, loc="upper left", + frameon=False, fontsize=10.6, labelcolor=TEXT2, loc="upper left", bbox_to_anchor=(0.0, 1.055), ncol=2, handlelength=1.1) ax.set_title("frames / second · 3×3, 100 000 frames, f32 · " "minor faults and the throughput they cost, per step", - color=MUTED, fontsize=9, loc="left", pad=8) + color=MUTED, fontsize=10.6, loc="left", pad=8) save(fig, "fig_first_run") @@ -345,14 +297,16 @@ def fig_arc(): TOP = 78000 # act bands, behind the bars, labelled along the top - for x0, x1, label, col in [(-0.5, 4.5, "ACT I · feed the GPU", ACCENT), - (4.5, 6.5, "ACT II · get results back", PALE), - (6.5, 7.5, "ACT III · kernel", AMBER)]: + # Short labels only. The gloss ("feed the GPU") is wider than the one-bar + # ACT III band, so spelling it out here puts ACT III's label on ACT II's. + for x0, x1, label, col in [(-0.5, 4.5, "ACT I", ACCENT), + (4.5, 6.5, "ACT II", PALE), + (6.5, 7.5, "ACT III", AMBER)]: ax.axvspan(x0, x1, color=col, alpha=0.05, zorder=0) 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=9, fontweight="bold") + color=col, fontsize=10.6, fontweight="bold") ax.bar(x, fps, width=0.62, color=colors, zorder=3, linewidth=0) @@ -362,23 +316,23 @@ 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=9, ha="left", va="bottom") + color=GREEN, fontsize=10.6, ha="left", va="bottom") for xi, (f, s) in enumerate(zip(fps, spd)): ax.text(xi, f + 1100, f"{f:,}", ha="center", va="bottom", - color=PALE, fontsize=10.5, fontweight="bold") + color=PALE, fontsize=12.4, fontweight="bold") ax.text(xi, f - 1600, ("base" if s == 1.0 else f"×{s:.2f}"), - ha="center", va="top", color=BG, fontsize=9, fontweight="bold") + ha="center", va="top", color=BG, fontsize=10.6, fontweight="bold") ax.set_xticks(x) - ax.set_xticklabels(steps, fontsize=9, color=TEXT2) + ax.set_xticklabels(steps, fontsize=10.6, color=TEXT2) ax.set_ylim(0, TOP) ax.set_yticks([]) bare(ax, keep=("bottom",)) ax.spines["bottom"].set_color(RULE) ax.set_title("frames / second · 3×3 clusters, 100 000 frames, warm run · " "every bar against the best CPU configuration (24 threads)", - color=MUTED, fontsize=9, loc="left", pad=8) + color=MUTED, fontsize=10.6, loc="left", pad=8) save(fig, "fig_arc") @@ -398,7 +352,7 @@ def fig_arc_9x9(): fig, ax = plt.subplots(figsize=(11.4, 3.8)) x = np.arange(len(steps)) - TOP = 56000 + TOP = 62000 for x0, x1, label, col in [(-0.5, 2.5, "ACT I", ACCENT), (2.5, 4.5, "ACT II", PALE), @@ -407,7 +361,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=9, fontweight="bold") + color=col, fontsize=10.6, fontweight="bold") ax.bar(x, fps, width=0.58, color=colors, zorder=3, linewidth=0) @@ -423,30 +377,30 @@ 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=9, va="bottom") + color=GREEN, fontsize=10.6, 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=9, ha="right", va="bottom", linespacing=1.35) + ax.text(5.46, 45000, "f32 D2H floor · 39 775 FPS\n(nsys estimates 39 614)", + color=AMBER, fontsize=10.6, 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, - fontsize=9, ha="center", va="center", fontweight="bold") + ax.text(2.55, 51000, "−40 % kernel → D2H binds instead", color=AMBER, + fontsize=10.6, ha="center", va="center", fontweight="bold") for xi, (f, s) in enumerate(zip(fps, spd)): ax.text(xi, f + 800, f"{f:,}", ha="center", va="bottom", - color=PALE, fontsize=10.5, fontweight="bold") + color=PALE, fontsize=12.4, fontweight="bold") ax.text(xi, f - 1100, ("base" if s == 1.0 else f"×{s:.2f}"), - ha="center", va="top", color=BG, fontsize=9, fontweight="bold") + ha="center", va="top", color=BG, fontsize=10.6, fontweight="bold") ax.set_xticks(x) - ax.set_xticklabels(steps, fontsize=9, color=TEXT2) + ax.set_xticklabels(steps, fontsize=10.6, color=TEXT2) ax.set_ylim(0, TOP) ax.set_yticks([]) bare(ax, keep=("bottom",)) ax.spines["bottom"].set_color(RULE) ax.set_title("frames / second · 9×9, 20 000 frames, cap 1700 (lossless) · " "best CPU configuration (32 threads) · opt1/opt2 are 3×3 only", - color=MUTED, fontsize=9, loc="left", pad=8) + color=MUTED, fontsize=10.6, loc="left", pad=8) save(fig, "fig_arc_9x9") @@ -479,32 +433,32 @@ def fig_overhead(): linewidth=0, alpha=0.55) for xi, (f, e, m) in enumerate(zip(floor, excess, meas)): ax.text(xi, m + ymax * 0.022, f"{m:.1f}", ha="center", color=PALE, - fontsize=9.5, fontweight="bold") + fontsize=11.2, fontweight="bold") if e > ymax * 0.05: ax.text(xi, f + e / 2, f"+{e:.0f}", ha="center", va="center", - color=BG, fontsize=9, fontweight="bold") + color=BG, fontsize=10.6, fontweight="bold") ax.set_xticks(x) - ax.set_xticklabels(steps, color=TEXT2, fontsize=9) + ax.set_xticklabels(steps, color=TEXT2, fontsize=10.6) ax.set_ylim(0, ymax) ax.set_yticks([]) bare(ax, keep=("bottom",)) - ax.set_title(title, color=PALE, fontsize=10, pad=8, loc="left") + ax.set_title(title, color=PALE, fontsize=11.8, pad=8, loc="left") - axes[0].set_ylabel("µs / frame", color=TEXT2) + axes[0].set_ylabel("µs / frame", color=TEXT2, fontsize=10.6) # 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=9, fontweight="bold") + fontsize=10.6, fontweight="bold") axes[0].annotate("host excess", xy=(0.30, 25), xytext=(1.15, 41), - color=TEXT2, fontsize=9, + color=TEXT2, fontsize=10.6, arrowprops=dict(arrowstyle="-", color=MUTED, lw=0.9)) - axes[1].text(1.5, 108, "Act II removes the host bar", color=PALE, fontsize=9, + axes[1].text(1.5, 108, "Act II removes the host bar", color=PALE, fontsize=10.6, ha="center", fontweight="bold") axes[1].add_patch(FancyArrowPatch((1.5, 104), (3.25, 34), arrowstyle="-|>", mutation_scale=10, color=PALE, lw=1.3, connectionstyle="arc3,rad=-0.22")) - axes[1].text(4.0, 52, "Act III lowers\nthe floor", color=AMBER, fontsize=9, + axes[1].text(4.0, 52, "Act III lowers\nthe floor", color=AMBER, fontsize=10.6, ha="center", fontweight="bold") axes[1].add_patch(FancyArrowPatch((4.0, 46), (4.0, 27), arrowstyle="-|>", mutation_scale=10, color=AMBER, lw=1.3)) @@ -572,7 +526,7 @@ def _draw_schedule(ax, frames, top_lane=3.0, H=None, K=None, D=None): def fig_streams(): - fig, axes = plt.subplots(3, 1, figsize=(7.7, 3.9)) + fig, axes = plt.subplots(3, 1, figsize=(7.7, 2.85)) FR = H_ + K_ + D_ # --- opt1: one stream, strictly serial @@ -580,7 +534,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=9, + ax.text(3 * FR + 6, 1.34, "one engine at a time", color=MUTED, fontsize=10.6, va="center") # --- opt2: 4 streams, barrier after each round. The drain is not drawn in; @@ -591,9 +545,12 @@ def fig_streams(): _draw_schedule(ax, frames) for h_idle, t_end in drains: ax.axvspan(h_idle, t_end, color=AMBER, alpha=0.13, zorder=1) - ax.text((drains[0][0] + drains[0][1]) / 2, 4.15, "barrier — H2D starves", - color=AMBER, fontsize=9, ha="center", va="bottom") - ax.set_ylim(-0.4, 4.9) + # The label has to clear the top lane AND stay out of the axes title, which + # is anchored to the axes box and so moves with the LIMITS, not with the + # data. Extra headroom in y is what separates the two. + ax.text((drains[0][0] + drains[0][1]) / 2, 4.05, "barrier — H2D starves", + color=AMBER, fontsize=10.6, ha="center", va="bottom") + ax.set_ylim(-0.4, 5.6) # --- opt3: no barriers, continuous ax = axes[2] @@ -601,7 +558,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=9, va="top") + color=ACCENT, fontsize=10.6, va="top") titles = ["opt1 · 1 stream, synchronous", "opt2 · 4 streams, sync barrier per round", "opt3 · 4 streams, barriers removed"] @@ -609,13 +566,16 @@ def fig_streams(): ax.set_xlim(-2, 178) ax.set_yticks([]); ax.set_xticks([]) bare(ax, keep=()) - ax.set_title(t, color=TEXT2, fontsize=9, loc="left", pad=4) + ax.set_title(t, color=TEXT2, fontsize=10.6, loc="left", pad=4) 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=9, labelcolor=TEXT2, ncol=3, loc="lower right", + fontsize=10.6, 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=9, loc="left") + # "time →" inside the axes, at the far end: as an xlabel it sat under the + # left edge, on top of the sentence already written there. + axes[2].text(176, -0.25, "time →", color=MUTED, fontsize=10.6, + ha="right", va="top") fig.subplots_adjust(hspace=0.75, bottom=0.10) save(fig, "fig_streams") @@ -631,7 +591,7 @@ def fig_gpu_model(): attack on one of the two arrows, so the audience needs the picture before the ladder starts, not after. """ - fig, ax = plt.subplots(figsize=(11.6, 3.9)) + fig, ax = plt.subplots(figsize=(11.6, 3.4)) def panel(x0, x1, y0, y1, title, col): ax.add_patch(FancyBboxPatch((x0, y0), x1 - x0, y1 - y0, @@ -639,16 +599,16 @@ def fig_gpu_model(): facecolor=PANEL, edgecolor=col, linewidth=1.6, zorder=2)) ax.text((x0 + x1) / 2, y1 - 2.2, title, ha="center", va="center", - color=col, fontsize=11, fontweight="bold") + color=col, fontsize=13.0, fontweight="bold") def chip(x0, x1, y0, y1, top, bottom, col): ax.add_patch(Rectangle((x0, y0), x1 - x0, y1 - y0, facecolor=BG, edgecolor=col, linewidth=1.3, zorder=3)) ax.text((x0 + x1) / 2, (y0 + y1) / 2 + 1.6, top, ha="center", - va="center", color=PALE, fontsize=11, fontweight="bold", + va="center", color=PALE, fontsize=13.0, fontweight="bold", zorder=4) ax.text((x0 + x1) / 2, (y0 + y1) / 2 - 1.9, bottom, ha="center", - va="center", color=TEXT2, fontsize=10, zorder=4) + va="center", color=TEXT2, fontsize=11.8, zorder=4) panel(1, 33, 3, 44, "HOST", MUTED) chip(4, 30, 31, 40, "CPU", "24 threads", ACCENT) @@ -664,25 +624,28 @@ def fig_gpu_model(): arrowprops=dict(arrowstyle="<|-|>", color=col, lw=1.4)) if lab: ax.text(x + 1.4, 24.5, lab, ha="left", va="center", color=col, - fontsize=10.5, fontweight="bold") + fontsize=12.4, fontweight="bold") ax.text(17, 5.2, "off-limits to the kernel", - ha="center", va="center", color=MUTED, fontsize=10.5) + ha="center", va="center", color=MUTED, fontsize=12.4) # the slow wire, between them ax.annotate("", xy=(69, 16.5), xytext=(31, 16.5), arrowprops=dict(arrowstyle="-|>", color=GREEN, lw=2.4)) ax.text(50, 18.4, "H2D the frame, 312.5 kB", ha="center", va="bottom", - color=GREEN, fontsize=11, fontweight="bold") + color=GREEN, fontsize=13.0, fontweight="bold") ax.annotate("", xy=(31, 10.5), xytext=(69, 10.5), arrowprops=dict(arrowstyle="-|>", color=PALE, lw=2.4)) ax.text(50, 8.6, "D2H the clusters, 93 kB at 3×3", ha="center", va="top", - color=PALE, fontsize=11, fontweight="bold") - ax.text(50, 31, "PCIe 4.0 ×16", ha="center", va="center", color=TEXT2, - fontsize=11, fontweight="bold") - ax.text(50, 27.5, "31.5 GB/s", ha="center", va="center", color=TEXT2, - fontsize=11) - ax.text(50, 23.6, "the narrowest link\nin the chain", ha="center", - va="center", color=MUTED, fontsize=10.5) + color=PALE, fontsize=13.0, fontweight="bold") + # The middle column carries four separate strings between the two panels. + # They are spaced off the arrows below them, not off each other: at 13 pt in + # a 34-unit gap, two of them a row apart is already a collision. + ax.text(50, 37.5, "PCIe 4.0 ×16", ha="center", va="center", color=TEXT2, + fontsize=13.0, fontweight="bold") + ax.text(50, 33.0, "31.5 GB/s", ha="center", va="center", color=TEXT2, + fontsize=13.0) + ax.text(50, 28.0, "the narrowest link in the chain", ha="center", + va="center", color=MUTED, fontsize=12.4) ax.set_xlim(-1, 101); ax.set_ylim(1, 46) ax.axis("off") @@ -691,41 +654,45 @@ def fig_gpu_model(): # ------------------------------------------------------------- 5. pinning def fig_pinning(): - fig = plt.figure(figsize=(7.7, 3.0)) - ax = fig.add_axes([0, 0.05, 0.60, 0.95]); ax.axis("off") - ax.set_xlim(0, 10.4); ax.set_ylim(0, 6.4) + fig = plt.figure(figsize=(7.7, 2.6)) + ax = fig.add_axes([0, 0.04, 0.655, 0.96]); ax.axis("off") + # 11.6 units across 5.04 in is 0.435 in per unit; every box below is sized + # so its longest string fits INSIDE it at 11.2 pt, and every arrow label + # sits above the boxes rather than in the gap between them, because the gap + # is narrower than the words that were being put in it. + ax.set_xlim(0, 11.6); ax.set_ylim(0, 7.0) 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=9.5, fontweight="bold") - ax.text(x + w / 2, y + h / 2 - 0.34, sub, ha="center", va="center", - color=MUTED, fontsize=9.5) + ax.text(x + w / 2, y + h / 2 + 0.24, label, ha="center", va="center", + color=PALE, fontsize=11.2, fontweight="bold") + ax.text(x + w / 2, y + h / 2 - 0.32, sub, ha="center", va="center", + color=MUTED, fontsize=11.2) - def arrow(x0, x1, y, color, label): + def arrow(x0, x1, y, color, label, ytext): 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=9.5) + ax.text((x0 + x1) / 2, ytext, label, ha="center", va="bottom", + color=color, fontsize=11.2) - ax.text(0, 5.95, "PAGEABLE · before opt4", color=AMBER, fontsize=9.5, + ax.text(0, 6.52, "PAGEABLE · before opt4", color=AMBER, fontsize=11.2, 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=9.5) + box(0, 4.10, 3.0, 1.05, "numpy array", "pageable") + box(4.3, 4.10, 3.0, 1.05, "driver staging", "hidden buffer") + box(8.6, 4.10, 3.0, 1.05, "GPU", "device memory") + arrow(3.0, 4.3, 4.62, AMBER, "memcpy", 5.35) + arrow(7.3, 8.6, 4.62, AMBER, "DMA", 5.35) + ax.text(0, 3.64, "every transfer is copied twice", color=MUTED, fontsize=11.2) - 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 · reads host RAM directly") + ax.text(0, 2.86, "PINNED · opt4", color=ACCENT, fontsize=11.2, fontweight="bold") + box(0, 0.90, 3.0, 1.05, "numpy array", "page-locked") + box(8.6, 0.90, 3.0, 1.05, "GPU", "device memory") + arrow(3.0, 8.6, 1.42, ACCENT, "DMA · reads host RAM directly", 2.15) ax.text(0, 0.42, "no staging copy, no page faults, fully async", - color=MUTED, fontsize=9.5) + color=MUTED, fontsize=11.2) # 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]) + ax2 = fig.add_axes([0.745, 0.16, 0.255, 0.62]) x = np.arange(2) # 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 @@ -735,16 +702,16 @@ def fig_pinning(): 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=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, + ax2.text(xi - 0.19, b + 2, f"{b:.0f}", ha="center", color=TEXT2, fontsize=11.2) + ax2.text(xi + 0.19, a + 2, f"{a:.0f}", ha="center", color=TEXT2, fontsize=11.2) + ax2.text(xi, 92, f"×{b / a:.2f}", ha="center", color=PALE, fontsize=11.2, fontweight="bold") ax2.set_xticks(x) ax2.set_xticklabels(["3×3\nH2D-bound", "9×9\nkernel-bound"], color=TEXT2, - fontsize=9.5) + fontsize=11.2) ax2.set_ylim(0, 104); ax2.set_yticks([]); bare(ax2, keep=("bottom",)) - ax2.set_title("µs / frame", color=MUTED, fontsize=9.5, pad=6) - ax2.legend(frameon=False, fontsize=9.5, labelcolor=TEXT2, loc="center left") + ax2.set_title("µs / frame", color=MUTED, fontsize=11.2, pad=6) + ax2.legend(frameon=False, fontsize=11.2, labelcolor=TEXT2, loc="center left") save(fig, "fig_pinning") @@ -757,12 +724,12 @@ def fig_graphs(): 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=9.5, fontweight="bold") + color=BG, fontsize=11.2, 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=9.5, fontweight="bold") + color=AMBER, fontsize=11.2, 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) @@ -770,10 +737,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=9.5) + color=MUTED, fontsize=11.2) ax.text(0, 2.18, "WITH GRAPHS · record once, replay with one launch", - color=ACCENT, fontsize=9.5, fontweight="bold") + color=ACCENT, fontsize=11.2, 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): @@ -781,7 +748,7 @@ 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=9.5, fontweight="bold") + color=ACCENT, fontsize=11.2, fontweight="bold") # the verdict ax2 = fig.add_axes([0.73, 0.13, 0.27, 0.78]) @@ -791,15 +758,15 @@ def fig_graphs(): 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=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, + ax2.text(xi - 0.19, a + 2.5, f"{a:.0f}", ha="center", color=TEXT2, fontsize=11.2) + ax2.text(xi + 0.19, g + 2.5, f"{g:.0f}", ha="center", color=TEXT2, fontsize=11.2) + ax2.text(1, 100, "−12 % THROUGHPUT", ha="center", color=AMBER, fontsize=11.2, fontweight="bold") ax2.set_xticks(x) - ax2.set_xticklabels(["3×3", "9×9"], color=TEXT2, fontsize=9.5) + ax2.set_xticklabels(["3×3", "9×9"], color=TEXT2, fontsize=11.2) ax2.set_ylim(0, 120); ax2.set_yticks([]); bare(ax2, keep=("bottom",)) - ax2.set_title("µs / frame", color=MUTED, fontsize=9.5, pad=6) - ax2.legend(frameon=False, fontsize=9.5, labelcolor=TEXT2, loc="upper left") + ax2.set_title("µs / frame", color=MUTED, fontsize=11.2, pad=6) + ax2.legend(frameon=False, fontsize=11.2, labelcolor=TEXT2, loc="upper left") save(fig, "fig_graphs") @@ -812,30 +779,37 @@ def fig_resultpath(): for ax, title, floor, copy_us, gain, verdict, col in [ (a1, "3×3 · host copy 93 kB / frame", 16.17, 8.0, "×1.16", - "copy hides under the GPU\n→ small win", ACCENT), + "the copy hides under the GPU floor → small win", ACCENT), (a2, "9×9 · host copy 467 kB / frame", 30.01, 62.0, "×2.21", - "copy is larger than the GPU\n→ cannot hide at any overlap", AMBER), + "the copy is larger than the floor → cannot hide", AMBER), ]: + # 1.35 apart, not 1.0: at 11.8 pt each two-line tick label is about + # 0.95 x-units wide, so adjacent bars put them into each other. 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.bar([1.35], [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.8, "GPU floor", color=GREEN, fontsize=10, ha="left") + # ABOVE the line and to the right of both bars. On the line it read as + # struck through; at the left edge it met a bar's own value label. + ax.text(2.55, floor + 1.2, "GPU floor", color=GREEN, fontsize=11.8, + ha="right", va="bottom") ax.text(0, floor + 2.1, f"{floor:.1f} µs", ha="center", color=PALE, - fontsize=10, fontweight="bold") - ax.text(1, copy_us + 2.1, f"≈{copy_us:.0f} µs", ha="center", color=PALE, - fontsize=10, fontweight="bold") - 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=10) - ax.set_xlim(-0.6, 1.7); ax.set_ylim(0, 78); ax.set_yticks([]) + fontsize=11.8, fontweight="bold") + ax.text(1.35, copy_us + 2.1, f"≈{copy_us:.0f} µs", ha="center", color=PALE, + fontsize=11.8, fontweight="bold") + ax.set_xticks([0, 1.35]) + ax.set_xticklabels(["GPU per frame\nH2D ∥ kernel ∥ D2H", + "host copy per frame\ncollect() memcpy"], + color=TEXT2, fontsize=11.8) + ax.set_xlim(-0.7, 2.6); ax.set_ylim(0, 82); ax.set_yticks([]) bare(ax, keep=("bottom",)) - ax.set_title(title, color=PALE, fontsize=10, pad=10, loc="left") - ax.text(1.68, 75, gain, color=col, fontsize=16, fontweight="bold", ha="right") - ax.text(1.68, 68.5, "opt5 → opt6", color=MUTED, fontsize=10, ha="right") - ax.text(-0.55, -16, verdict, color=col, fontsize=10, fontweight="bold", - va="top") - fig.subplots_adjust(bottom=0.30) + # The verdict is the second line of the TITLE. As free-floating text it + # had nowhere to go: above the bars it met the value labels, below the + # axis it met two-line tick labels, and beside them it met the gain. + ax.set_title(f"{title}\n{verdict}", color=PALE, fontsize=11.8, pad=10, + loc="left", linespacing=1.5) + ax.text(2.55, 76, gain, color=col, fontsize=18.9, fontweight="bold", ha="right") + ax.text(2.55, 69, "opt5 → opt6", color=MUTED, fontsize=11.8, ha="right") + fig.subplots_adjust(bottom=0.24, top=0.78) save(fig, "fig_resultpath") @@ -850,7 +824,7 @@ def fig_f32_kernel(): number is occupancy. Only at s4 does the f32 kernel fall below D2H, which is why the left panel cannot be used to argue the handover. """ - fig, (ax1, ax2) = plt.subplots(1, 2, figsize=(8.0, 3.15)) + fig, (ax1, ax2) = plt.subplots(1, 2, figsize=(7.4, 2.8)) labels = ["kernel", "D2H", "H2D"] y = np.arange(3); h = 0.34 panels = [ @@ -872,27 +846,29 @@ 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=10, fontweight="bold" if top else "normal") + fontsize=11.8, 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") + color=PALE, fontsize=12.4, 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([]) + ha="right", va="top", color=PALE, fontsize=11.8, fontweight="bold") + ax.set_yticks(y); ax.set_yticklabels(labels, color=TEXT2, fontsize=11.8) + # "39.86 ◀ binds" is ~22 x-units of text starting at the bar's end, so + # the axis has to be long enough to hold the label as well as the bar. + ax.invert_yaxis(); ax.set_xlim(0, 80); ax.set_xticks([]) ax.set_ylim(3.05, -0.98) bare(ax, keep=("left",)) - ax.set_title(title, color=PALE, fontsize=10, pad=10, loc="left") + ax.set_title(title, color=PALE, fontsize=11.8, 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=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", + ax1.legend(frameon=False, fontsize=11.8, 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) + fig.subplots_adjust(bottom=0.20, top=0.86, left=0.10, right=0.99, wspace=0.34) save(fig, "fig_f32_kernel") @@ -926,18 +902,22 @@ def fig_cancellation(): 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) + # Headroom above the 2.17e7 bars, so the note can sit over empty axes + # instead of over the two operands it is talking about. + ax.set_yscale("log"); ax.set_ylim(1, 3e11) ax.set_xticks([0, 1, 2, 3]) - ax.set_xticklabels(names, color=TEXT2, fontsize=9, linespacing=1.35) + ax.set_xticklabels(names, color=TEXT2, fontsize=10.6, linespacing=1.35) ax.set_yticks([1e0, 1e2, 1e4, 1e6, 1e8]) - ax.tick_params(labelsize=9) + ax.set_xlim(-0.6, 3.7) + ax.tick_params(labelsize=10.6) 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) + ax.text(1.5, 3.0e9, "±3 ADU² of f32 error —\nit swallows the quiet pixels", + color=ACCENT, fontsize=10.6, ha="center", va="center", + linespacing=1.4) + ax.annotate("", xy=(3.34, 3.4), xytext=(3.34, 6e7), + arrowprops=dict(arrowstyle="->", color=ACCENT, lw=0.8)) bare(ax) - ax.set_title("var = E[X²] − mean²", color=MUTED, fontsize=9.5, pad=8, loc="left") + ax.set_title("var = E[X²] − mean²", color=MUTED, fontsize=11.2, pad=8, loc="left") d = json.loads((Path(__file__).resolve().parent / "validation_tiers.json").read_text()) @@ -958,20 +938,23 @@ def fig_cancellation(): 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_ylim(2e2, 2e9) 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) + ax2.tick_params(labelsize=10.6, colors=MUTED) + ax2.set_xlabel("cluster energy [ADU]", color=TEXT2, fontsize=11.2) bare(ax2, keep=("left", "bottom")) - ax2.legend(frameon=False, fontsize=9, labelcolor=TEXT2, loc="lower right", + ax2.legend(frameon=False, fontsize=10.6, labelcolor=TEXT2, loc="lower right", handlelength=1.2, borderaxespad=0.3) + # Both labels live in the decade of headroom above the 1.9e6 peak, so + # neither can land on the curve at any zoom. 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, + xy=(170, 3.0e4), xytext=(700, 2.4e8), color=AMBER, fontsize=10.6, + linespacing=1.4, zorder=6, ha="center", 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, + ax2.text(1188, 6.0e6, "Cu Kα", color=PALE, fontsize=10.6, ha="center", + va="bottom") + ax2.set_title("cluster-energy spectrum, 3×3", color=MUTED, fontsize=11.2, 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") @@ -999,21 +982,21 @@ def fig_varfloor(): 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=9.5, va="center", zorder=6, + fontsize=11.2, 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=9.5) - ax2.set_ylabel("variance (ADU²)", color=TEXT2, fontsize=9.5) + ax2.set_xlabel("pixel rms (ADU)", color=TEXT2, fontsize=11.2) + ax2.set_ylabel("variance (ADU²)", color=TEXT2, fontsize=11.2) 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=9.5, colors=MUTED) + ax2.tick_params(labelsize=10.4, 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=9.5, va="top", linespacing=1.45, zorder=6) + color=AMBER, fontsize=11.2, va="top", linespacing=1.45, zorder=6) ax2.text(2.20, 38.8, "rms 2–5: threshold\ncorrupted, not clamped", - 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, + color=MUTED, fontsize=11.2, va="top", linespacing=1.45, zorder=6) + ax2.text(3.05, 31.0, "true variance = rms²", color=PALE, fontsize=11.2, zorder=6) + ax2.text(5.90, 5.6, "f32 error floor ±3–4 ADU²", color=ACCENT, fontsize=11.2, ha="right", zorder=6) fig.subplots_adjust(left=0.13, right=0.98, top=0.97, bottom=0.15) save(fig, "fig_varfloor") @@ -1052,25 +1035,25 @@ def fig_bottleneck(): ax.plot([xi - 0.13, xi + 0.13], [p_, p_], color=c, lw=2.6, zorder=3) if h_ - l_ < 1: ax.text(xi, p_ - 5.0, f"{p_:.1f}%", ha="center", color=PALE, - fontsize=11.5, fontweight="bold") + fontsize=13.6, fontweight="bold") ax.text(xi, p_ + 2.0, "resolvable to 0.0 pts", ha="center", - color=AMBER, fontsize=8) + color=AMBER, fontsize=9.4) else: ax.text(xi, h_ + 1.6, f"{l_:+.0f} … {h_:+.0f}%", ha="center", - color=TEXT2, fontsize=9.5) + color=TEXT2, fontsize=11.2) - ax.set_xticks(x); ax.set_xticklabels(steps, color=TEXT2, fontsize=9) + ax.set_xticks(x); ax.set_xticklabels(steps, color=TEXT2, fontsize=10.6) ax.set_xlim(-0.6, 3.6) ax.set_ylim(-34, 30) ax.set_yticks([-20, 0, 20]) - ax.set_yticklabels(["−20 %", "0", "+20 %"], fontsize=8) + ax.set_yticklabels(["−20 %", "0", "+20 %"], fontsize=9.4) bare(ax, keep=("left", "bottom")) ax.spines["bottom"].set_color(RULE) ax.set_title("end-to-end change from the SAME −40 % kernel, 9×9 · " "bar = measurement, band = what the reps allow", - color=MUTED, fontsize=9, pad=10, loc="left") + color=MUTED, fontsize=10.6, pad=10, loc="left") ax.text(1.5, -29.5, "through an allocating result path the effect is not measurable " - "— the band straddles zero", color=MUTED, fontsize=8.5, + "— the band straddles zero", color=MUTED, fontsize=10.0, ha="center", va="center") fig.subplots_adjust(bottom=0.24) save(fig, "fig_bottleneck") @@ -1086,15 +1069,15 @@ def fig_correctness(): ax.bar(x, diff, width=0.5, color=colors, zorder=3) for xi, d in enumerate(diff): ax.text(xi, d + 0.00022, ("reference" if d == 0 else f"{d:.4f}%"), - ha="center", color=PALE if d else MUTED, fontsize=8.5) + ha="center", color=PALE if d else MUTED, fontsize=10.0) ax.axhline(0.01, color=PALE, lw=1.2, ls="--") ax.text(4.4, 0.0104, "0.01% — well inside statistical noise", color=PALE, - fontsize=8, ha="right") - ax.set_xticks(x); ax.set_xticklabels(names, color=TEXT2, fontsize=8.5) + fontsize=9.4, ha="right") + ax.set_xticks(x); ax.set_xticklabels(names, color=TEXT2, fontsize=10.0) ax.set_ylim(0, 0.0125); ax.set_yticks([]) bare(ax, keep=("bottom",)) ax.set_title("cluster-count difference vs CPU · 233 M clusters, 3×3", - color=MUTED, fontsize=8.5, pad=8) + color=MUTED, fontsize=10.0, pad=8) save(fig, "fig_correctness") @@ -1130,7 +1113,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=10.5, fontweight="bold", zorder=4) + color=tcol, fontsize=12.4, fontweight="bold", zorder=4) # ---- serial: G H G H G H G H, strictly alternating ---------------------- yG, yH = 3.30, 2.78 @@ -1152,13 +1135,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=10.5) + color=TEXT2, fontsize=12.4) ax.text(-0.25, yG + lane_h + 0.30, "submit → collect, serialized (opt4) · 3×3", - ha="left", va="bottom", color=TEXT2, fontsize=10.5, fontweight="bold") + ha="left", va="bottom", color=TEXT2, fontsize=12.4, 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=10.5, fontweight="bold") + ha="left", va="bottom", color=AMBER, fontsize=12.4, fontweight="bold") for x, y0, y1, col in ((serial_end, yH, yG + lane_h, MUTED), (pipe_end, yH2, yG2 + lane_h, AMBER)): @@ -1169,12 +1152,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=10.5, fontweight="bold") + color=AMBER, fontsize=12.4, fontweight="bold") ax.text(serial_end + 0.25, yG + lane_h / 2, "GPU + host per chunk", - ha="left", va="center", color=MUTED, fontsize=10.5) + ha="left", va="center", color=MUTED, fontsize=12.4) ax.text(pipe_end + 0.25, yG2 + lane_h / 2, "max(GPU, host) per chunk", - ha="left", va="center", color=AMBER, fontsize=10.5, fontweight="bold") + ha="left", va="center", color=AMBER, fontsize=12.4, fontweight="bold") ax.set_xlim(-1.6, serial_end + 4.0) ax.set_ylim(0, 4.25) @@ -1206,7 +1189,7 @@ def fig_overlap_9x9(): 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=BG, fontsize=10, fontweight="bold", zorder=4) + color=BG, fontsize=11.8, fontweight="bold", zorder=4) # The host is saturated in both cases, so its lane is the same schedule twice: # chunk i is collected as soon as the host is free, never before G. @@ -1241,9 +1224,12 @@ def fig_overlap_9x9(): block(host_start[i], yH, H, PALE, f"host {i + 1}") for lbl, y in (("GPU", yG), ("host", yH)): ax.text(-6, y + lane_h / 2, lbl, ha="right", va="center", - color=TEXT2, fontsize=10.5) - ax.text(-6, yG + lane_h + 7.0, tag, ha="left", va="bottom", - color=col, fontsize=10.5, fontweight="bold") + color=TEXT2, fontsize=12.4) + # The tag is left-aligned and 70 x-units long; the idle marker sits at + # the middle of the stall, which is inside that span. They can only be + # separated vertically, so the tag clears the marker by a full line. + ax.text(-6, yG + lane_h + 13.0, tag, ha="left", va="bottom", + color=col, fontsize=12.4, fontweight="bold") # every gap the GPU sits through, marked where it happens for i in range(1, n): gap = gpu[i] - (gpu[i - 1] + G) @@ -1253,14 +1239,14 @@ def fig_overlap_9x9(): arrowprops=dict(arrowstyle="<|-|>", color=AMBER, lw=1.3)) ax.text((gpu[i] + gpu[i - 1] + G) / 2, yG + lane_h + 1.0, f"idle {gap:.0f} µs", ha="center", va="bottom", - color=AMBER, fontsize=10.5, fontweight="bold") + color=AMBER, fontsize=12.4, fontweight="bold") - ax.plot([finish, finish], [1, 65], color=GREEN, lw=1.6, ls="--", zorder=6) + ax.plot([finish, finish], [1, 70], color=GREEN, lw=1.6, ls="--", zorder=6) ax.text(finish + 5, 33, "same finish\nboth ways", ha="left", va="center", - color=GREEN, fontsize=11, fontweight="bold") + color=GREEN, fontsize=13.0, fontweight="bold") ax.set_xlim(-32, finish + 62) - ax.set_ylim(0, 76) + ax.set_ylim(0, 86) ax.axis("off") save(fig, "fig_overlap_9x9") @@ -1340,7 +1326,7 @@ def fig_test3(): ax.annotate("", xy=(-0.85, 0.0), xytext=(-0.85, n), arrowprops=dict(arrowstyle="-|>", color=MUTED, lw=1.2)) ax.text(n / 2, n + 0.85, "raster order", ha="center", va="bottom", - color=MUTED, fontsize=10.5) + color=MUTED, fontsize=12.4) # the window Test3 summed, and inside it the one pixel the models disagree on ax.add_patch(Rectangle((R - 1, n - R - 2), 3, 3, facecolor="none", edgecolor=PALE, linewidth=1.4, ls=(0, (3, 2)), @@ -1349,7 +1335,7 @@ def fig_test3(): edgecolor=GREEN, linewidth=2.4, zorder=7)) ax.annotate("the pixel they\ndisagree on", xy=(R + 1.6, n - R - 0.5), xytext=(n + 1.4, n - R - 0.5), ha="left", va="center", - color=PALE, fontsize=10.5, + color=PALE, fontsize=12.4, arrowprops=dict(arrowstyle="-", color=PALE, lw=1.0, shrinkA=2, shrinkB=2)) # Headroom for the title ABOVE the raster-order label, not on top of it. @@ -1357,7 +1343,7 @@ def fig_test3(): ax.set_aspect("equal"); ax.axis("off") ax.text(n / 2, n + 3.5, f"the finished frame · 21 × 21 around " f"({site['iy']}, {site['ix']})", - ha="center", va="top", color=TEXT2, fontsize=10.5) + ha="center", va="top", color=TEXT2, fontsize=12.4) # ---- right: the 3x3, with each finder's reading named ------------------- for r in range(3): @@ -1369,14 +1355,14 @@ def fig_test3(): linewidth=2.2 if past else 1.6, zorder=2)) if past: bx.text(c + 0.5, 2 - r + 0.66, f"{vf[r, c]:.3f}", ha="center", - va="center", color=RED, fontsize=11, + va="center", color=RED, fontsize=13.0, fontweight="bold", zorder=3) bx.text(c + 0.5, 2 - r + 0.32, f"{vc[r, c]:.3f}", ha="center", - va="center", color=ACCENT, fontsize=11, + va="center", color=ACCENT, fontsize=13.0, fontweight="bold", zorder=3) else: bx.text(c + 0.5, 2 - r + 0.50, f"{vf[r, c]:.3f}", ha="center", - va="center", color=PALE, fontsize=11, + va="center", color=PALE, fontsize=13.0, fontweight="bold", zorder=3) bx.add_patch(Rectangle((1, 1), 1, 1, facecolor="none", edgecolor=GREEN, linewidth=2.6, zorder=4)) @@ -1384,7 +1370,7 @@ def fig_test3(): bx.set_aspect("equal"); bx.axis("off") bx.text(1.5, 3.60, f"the moment ({site['iy']}, {site['ix']}) was tested · " f"only the 4 already-scanned neighbours differ", - ha="center", va="top", color=TEXT2, fontsize=10.5) + ha="center", va="top", color=TEXT2, fontsize=12.4) # The verdict goes in figure coords: inside the axes it would be laid out # against an equal-aspect box and squeeze the grid into a column. Three @@ -1395,7 +1381,7 @@ def fig_test3(): facecolor=col, edgecolor=edge or BG, linewidth=2.0 if edge else 0.8, transform=fig.transFigure)) - fig.text(0.046, y + 0.017, txt, color=TEXT2, fontsize=10.5, + fig.text(0.046, y + 0.017, txt, color=TEXT2, fontsize=12.4, va="center") key(0.245, GREEN, "stored: this pixel IS the window max, and it clears 5σ") @@ -1409,16 +1395,16 @@ def fig_test3(): # Keep this line no longer than the sigma rows below it: bbox_inches="tight" # widens the canvas to the widest string, which shrinks every other one. fig.text(0.545, 0.250, "frozen above serial · amber ring = pushed pedestal", - color=TEXT2, fontsize=10.5) + color=TEXT2, fontsize=12.4) fig.text(0.545, 0.175, f"serial Σ = {site['total_cpu']:.3f} < " f"{thr:.3f} → pedestal sample", - color=ACCENT, fontsize=10.5, fontweight="bold") + color=ACCENT, fontsize=12.4, fontweight="bold") fig.text(0.545, 0.100, f"frozen Σ = {site['total_frz']:.3f} > " f"{thr:.3f} → CLUSTER", - color=RED, fontsize=10.5, fontweight="bold") + color=RED, fontsize=12.4, fontweight="bold") fig.text(0.545, 0.025, f"max = {site['max_frz']:.3f} in both → " f"Test1 never moved", - color=GREEN, fontsize=10.5, fontweight="bold") + color=GREEN, fontsize=12.4, fontweight="bold") fig.subplots_adjust(left=0.015, right=0.985, top=0.97, bottom=0.31, wspace=0.06) @@ -1439,7 +1425,7 @@ def fig_pagefault(): mapped" and "you write here" between them, where they collided with the very arrows they were naming. """ - fig, ax = plt.subplots(figsize=(11.4, 3.15)) + fig, ax = plt.subplots(figsize=(11.4, 2.85)) PW, GAP, N = 9.0, 1.6, 4 xs = [2 + i * (PW + GAP) for i in range(N)] cx = [x + PW / 2 for x in xs] @@ -1454,9 +1440,9 @@ def fig_pagefault(): box(x, 32.0, RED if i == 2 else ACCENT, False) box(x, 9.0, [AMBER, AMBER, RED, MUTED][i], i >= 2) ax.text(2, 38.6, "VIRTUAL · the ClusterVector malloc just handed you", - ha="left", va="bottom", color=TEXT2, fontsize=10.5) + ha="left", va="bottom", color=TEXT2, fontsize=12.4) ax.text(2, 7.8, "PHYSICAL · 4 kB frames the kernel actually owns", - ha="left", va="top", color=TEXT2, fontsize=10.5) + ha="left", va="top", color=TEXT2, fontsize=12.4) for i in (0, 1): ax.annotate("", xy=(cx[i], 13.8), xytext=(cx[i], 31.9), @@ -1465,32 +1451,34 @@ def fig_pagefault(): arrowprops=dict(arrowstyle="-|>", color=RED, lw=1.8, linestyle=(0, (3, 2)))) - # key, to the right of the rows: nothing is written between them - for y, col, ls, lab in ((25.0, GREEN, "-", "already mapped"), - (16.0, RED, (0, (3, 2)), "first touch → fault")): - ax.plot([44, 52], [y, y], color=col, lw=1.8, ls=ls) - ax.text(44, y + 1.8, lab, ha="left", va="bottom", color=col, - fontsize=10.5) + # Key UNDER the two rows, not beside them: the right-hand half of the figure + # belongs to the fault panel, and "first touch → fault" is 18 x-units wide, + # which is exactly the clearance that was left between the two. + for x0, col, ls, lab in ((2, GREEN, "-", "already mapped"), + (26, RED, (0, (3, 2)), "first touch → fault")): + ax.plot([x0, x0 + 4.5], [2.4, 2.4], color=col, lw=1.8, ls=ls) + ax.text(x0 + 5.6, 2.4, lab, ha="left", va="center", color=col, + fontsize=12.4) - ax.add_patch(FancyBboxPatch((61, 6), 39, 32, + ax.add_patch(FancyBboxPatch((58, 4), 42, 34, boxstyle="round,pad=0,rounding_size=1.2", facecolor=PANEL, edgecolor=RED, linewidth=1.5, zorder=3)) - ax.text(63.5, 34.2, "MINOR FAULT · what the kernel does", ha="left", - va="center", color=RED, fontsize=10.5, fontweight="bold") + ax.text(60.5, 34.4, "MINOR FAULT · in the kernel", ha="left", + va="center", color=RED, fontsize=12.4, fontweight="bold") for i, (n, txt, hot) in enumerate([("1", "traps into the kernel", False), ("2", "finds a free frame", False), ("3", "zeroes all 4 kB of it", True), ("4", "maps it; the write proceeds", False)]): - y = 28.6 - i * 5.0 - ax.text(64.0, y, n, ha="left", va="center", color=MUTED, fontsize=10.5) - ax.text(67.0, y, txt, ha="left", va="center", - color=AMBER if hot else TEXT2, fontsize=10.5, + y = 28.8 - i * 5.2 + ax.text(61.0, y, n, ha="left", va="center", color=MUTED, fontsize=12.4) + ax.text(64.0, y, txt, ha="left", va="center", + color=AMBER if hot else TEXT2, fontsize=12.4, fontweight="bold" if hot else "normal") - ax.text(63.5, 8.8, "no disk: that would be a MAJOR fault", ha="left", - va="center", color=MUTED, fontsize=10.5) + ax.text(60.5, 6.6, "no disk: that would be a MAJOR fault", ha="left", + va="center", color=MUTED, fontsize=12.4) - ax.set_xlim(-1, 101); ax.set_ylim(3, 42) + ax.set_xlim(-1, 101); ax.set_ylim(0, 42) ax.axis("off") save(fig, "fig_pagefault") @@ -1514,13 +1502,13 @@ def fig_pedtiming(): """ fig, ax = plt.subplots(figsize=(11.6, 3.6)) XF, XE = 6.0, 7.4 # frame end, dotted continuation end - rows = [("ClusterFinder\nserial CPU", 2.75, ACCENT, True), - ("ClusterFinderFrozen\nCPU twin", 1.55, PALE, False), - ("ClusterFinderCUDA", 0.35, AMBER, False)] + rows = [("ClusterFinder\nserial CPU", 3.00, ACCENT, True), + ("ClusterFinderFrozen\nCPU twin", 1.70, PALE, False), + ("ClusterFinderCUDA", 0.40, AMBER, False)] for label, y, col, stair in rows: ax.text(-0.35, y + 0.26, label, ha="right", va="center", color=TEXT2, - fontsize=9.5, linespacing=1.35) + fontsize=11.2, linespacing=1.35) ax.plot([0, XE], [y, y], color=RULE, lw=1.0, zorder=1) if stair: @@ -1530,42 +1518,42 @@ def fig_pedtiming(): ax.plot([XF, XE], [y + 0.52, y + 0.62], color=col, lw=2.0, ls=":", zorder=3) ax.text(XF / 2, y + 0.64, "the pedestal moves DURING the scan", - ha="center", va="bottom", color=col, fontsize=9, + ha="center", va="bottom", color=col, fontsize=10.6, fontweight="bold") ax.text(XF / 2, y - 0.13, - "a pixel late in the frame is judged against a pedestal that\n" - "already contains this frame's earlier pixels", - ha="center", va="top", color=MUTED, fontsize=8, - linespacing=1.4) + "a late pixel is judged against a pedestal that already moved", + ha="center", va="top", color=MUTED, fontsize=10.0) else: ax.plot([0, XF], [y + 0.10] * 2, color=col, lw=2.2, zorder=3) ax.plot([XF, XF], [y + 0.10, y + 0.52], color=col, lw=2.2, zorder=3) ax.plot([XF, XE], [y + 0.52] * 2, color=col, lw=2.0, ls=":", zorder=3) - ax.text(XF / 2, y + 0.18, "every decision uses the frame-start snapshot", - ha="center", va="bottom", color=col, fontsize=9, + # ABOVE the riser, not on the trace: at y + 0.18 the label sat + # directly on the 2.2 pt line it was labelling. + ax.text(XF / 2, y + 0.62, "every decision uses the frame-start snapshot", + ha="center", va="bottom", color=col, fontsize=10.6, fontweight="bold") ax.axvline(XF, color=MUTED, lw=1.1, ls="--", zorder=2, ymin=0.02, ymax=0.93) - ax.text(XF / 2, 3.80, "one frame · 160 000 pixels in raster order", - ha="center", va="bottom", color=MUTED, fontsize=9.5) - ax.text(XF + 0.12, 3.80, "frame ends →\nupdates applied", ha="left", - va="bottom", color=MUTED, fontsize=8.5, linespacing=1.35) + ax.text(XF / 2, 4.10, "one frame · 160 000 pixels in raster order", + ha="center", va="bottom", color=MUTED, fontsize=11.2) + ax.text(XF + 0.12, 4.10, "frame ends →\nupdates applied", ha="left", + va="bottom", color=MUTED, fontsize=10.0, linespacing=1.35) # what the middle row buys: two comparisons, one variable each XA = 8.35 - ax.annotate("", xy=(XA, 1.72), xytext=(XA, 2.92), + ax.annotate("", xy=(XA, 1.87), xytext=(XA, 3.17), arrowprops=dict(arrowstyle="<|-|>", color=PALE, lw=1.5)) - ax.text(XA + 0.18, 2.32, "update TIMING\narithmetic held fixed", ha="left", - va="center", color=PALE, fontsize=9, fontweight="bold", + ax.text(XA + 0.18, 2.52, "update TIMING\narithmetic held fixed", ha="left", + va="center", color=PALE, fontsize=10.6, fontweight="bold", linespacing=1.4) - ax.annotate("", xy=(XA, 0.52), xytext=(XA, 1.72), + ax.annotate("", xy=(XA, 0.57), xytext=(XA, 1.87), arrowprops=dict(arrowstyle="<|-|>", color=AMBER, lw=1.5)) - ax.text(XA + 0.18, 1.12, "the PORT\ntiming held fixed", ha="left", - va="center", color=AMBER, fontsize=9, fontweight="bold", + ax.text(XA + 0.18, 1.22, "the PORT\ntiming held fixed", ha="left", + va="center", color=AMBER, fontsize=10.6, fontweight="bold", linespacing=1.4) ax.set_xlim(-2.9, 12.0) - ax.set_ylim(-0.35, 4.30) + ax.set_ylim(-0.30, 4.60) ax.axis("off") save(fig, "fig_pedtiming") @@ -1611,8 +1599,13 @@ def fig_mismatch147(): for i in range(w.shape[0]): for j in range(w.shape[1]): if m[i, j]: + # A cell is ~0.18 in wide in the saved figure and a value can + # be three digits, so anything above ~5.5 pt collides with + # the cell next to it. These are texture, not readings: the + # slide's argument is which pixels are lit, and the numbers + # are on annex A7 and in the notes for anyone who wants them. ax.text(j, i, f"{w[i, j]:.0f}", ha="center", va="center", - fontsize=6.4, zorder=7, + fontsize=5.4, zorder=7, gid="texture", 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=4.5, zorder=4) @@ -1624,14 +1617,14 @@ def fig_mismatch147(): ax.add_patch(plt.Circle((cx - X0, cy - Y0), 2.6, fill=False, ec=AMBER, lw=1.4, ls=":", zorder=6)) ax.set_title(f"{name} — {d['n_' + name]:,} clusters in this frame", - color=col, fontsize=10, fontweight="bold", pad=7) + color=col, fontsize=11.8, fontweight="bold", pad=7) ax.set_xticks([]); ax.set_yticks([]) for sp in ax.spines.values(): sp.set_color(RULE) 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=9, linespacing=1.4) + va="center", color=MUTED, fontsize=10.6, 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. @@ -1639,7 +1632,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=9, fontweight="bold", ha="right", + color=AMBER, fontsize=10.6, fontweight="bold", ha="right", linespacing=1.35, arrowprops=dict(arrowstyle="-|>", color=AMBER, lw=1.3)) fig.subplots_adjust(wspace=0.06) @@ -1675,7 +1668,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=10, + ax.text(frac + 1.5, y, f"{frac:.0f} %", va="center", fontsize=11.8, color=PALE if full else TEXT2, fontweight="bold" if full else "normal") ticks.append(y); labels.append(f"{name} · {kind}") @@ -1683,19 +1676,20 @@ 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=10, + ax.text(99, 1.05, "capacity of one SM", color=MUTED, fontsize=11.8, ha="right") ax.set_yticks(ticks) - ax.set_yticklabels(labels, color=TEXT2, fontsize=10) - ax.set_xlim(0, 152); ax.set_xticks([]) + ax.set_yticklabels(labels, color=TEXT2, fontsize=11.8) + ax.set_xlim(0, 172); 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=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=10, va="center", linespacing=1.4) + # Clear of the "100 %" value labels, which are bold and ~10 x-units wide. + ax.text(122, ticks[0] - 0.5, "6 blocks resident\n38 × 256 = 9 728 regs each", + color=ACCENT, fontsize=11.8, va="center", linespacing=1.4) + ax.text(122, ticks[2] - 0.5, "2 blocks resident\n128 × 256 = 32 768 regs each", + color=AMBER, fontsize=11.8, va="center", linespacing=1.4) fig.subplots_adjust(left=0.16, right=0.99, top=0.88, bottom=0.06) save(fig, "fig_regpressure") @@ -1711,7 +1705,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=9.5, labelcolor=TEXT2, ncol=3, loc="lower right", + fontsize=11.2, labelcolor=TEXT2, ncol=3, loc="lower right", bbox_to_anchor=(1.02, y), handlelength=1.1) @@ -1735,9 +1729,9 @@ def fig_opt1_timeline(): ax.axvline(t0 + PER, color=AMBER, lw=0.9, alpha=0.55, zorder=1) ax.annotate("host blocks: every engine idle", xy=(WORK + HOST / 2, 0.98), xytext=(WORK + HOST / 2, 0.56), - color=AMBER, fontsize=9, ha="center", va="top", + color=AMBER, fontsize=10.6, 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=9, + ax.text(3 * PER + 4, 1.34, "one engine\nat a time", color=TEXT2, fontsize=10.6, va="center", linespacing=1.5) ax.set_xlim(-4, 3 * PER + 40) ax.set_ylim(0.02, 2.45) @@ -1761,7 +1755,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=9, ha="right", va="center") + fontsize=10.6, 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 @@ -1776,15 +1770,15 @@ 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=9, ha="center", va="bottom") + fontsize=10.6, 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=9, va="center", linespacing=1.5) + color=TEXT2, fontsize=10.6, va="center", linespacing=1.5) _engine_legend(ax, y=0.94) - ax.set_xlabel("time →", color=MUTED, fontsize=9, loc="left") + ax.set_xlabel("time →", color=MUTED, fontsize=10.6, loc="left") fig.subplots_adjust(left=0.01, right=0.99, top=0.90, bottom=0.19) save(fig, "fig_opt2_timeline") @@ -1832,30 +1826,30 @@ def fig_f32_absolute(): axA.bar(x + w / 2 + 0.015, f32v, width=w, color=AMBER, zorder=3, linewidth=0) for i, (a, b) in enumerate(zip(f64v, f32v)): axA.text(i - w / 2 - 0.015, a + 1.6, f"{a:.1f}", ha="center", - color=TEXT2, fontsize=8.5) + color=TEXT2, fontsize=10.0) axA.text(i + w / 2 + 0.015, b + 1.6, f"{b:.1f}", ha="center", - color=TEXT2, fontsize=8.5) + color=TEXT2, fontsize=10.0) axA.annotate("opt6 + f32 = opt7\nthe shipped build", xy=(2.19, 19.0), xytext=(2.36, 56), arrowprops=dict(arrowstyle="->", color=AMBER, lw=1.2), - color=AMBER, fontsize=8.5, fontweight="bold", linespacing=1.5, + color=AMBER, fontsize=10.0, fontweight="bold", linespacing=1.5, ha="center") axA.set_xlim(-0.62, 3.02) axA.set_ylim(0, 94) axA.set_yticks([0, 25, 50, 75]) - axA.set_yticklabels(["0", "25", "50", "75"], fontsize=8.5) - axA.set_ylabel("end-to-end µs / frame", color=MUTED, fontsize=8.5) + axA.set_yticklabels(["0", "25", "50", "75"], fontsize=10.0) + axA.set_ylabel("end-to-end µs / frame", color=MUTED, fontsize=10.0) axA.set_xticks(x) axA.set_xticklabels([f"{s}\n{t}" for s, t in zip(steps, sub)], - fontsize=8.5, linespacing=1.6, color=TEXT2) + fontsize=10.0, linespacing=1.6, color=TEXT2) axA.tick_params(axis="x", length=0, pad=7) bare(axA, keep=("left", "bottom")) axA.spines["bottom"].set_color(RULE) handles = [Rectangle((0, 0), 1, 1, color=c) for c in (ACCENT, AMBER)] axA.legend(handles, ["f64 pedestal", "f32 pedestal"], frameon=False, - fontsize=8.5, labelcolor=TEXT2, ncol=2, loc="upper right", + fontsize=10.0, labelcolor=TEXT2, ncol=2, loc="upper right", bbox_to_anchor=(1.02, 1.10), handlelength=1.1) - axA.set_title("the frame shrinks by 2.7×", color=MUTED, fontsize=9, + axA.set_title("the frame shrinks by 2.7×", color=MUTED, fontsize=10.6, pad=14, loc="left") # ---- B: the saving, which does not @@ -1863,20 +1857,20 @@ def fig_f32_absolute(): axB.axhline(np.mean(dv), color=PALE, lw=0.9, ls="--", zorder=4) for i, (d, p_) in enumerate(zip(dv, pct)): axB.text(i, d + 0.24, f"−{d:.2f} µs", ha="center", color=PALE, - fontsize=10.5, fontweight="bold") + fontsize=12.4, fontweight="bold") axB.set_ylim(0, 7.6) axB.set_yticks([0, 2, 4]) - 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_yticklabels(["0", "2", "4"], fontsize=10.0) + axB.set_ylabel("µs saved by the f32 pedestal", color=MUTED, fontsize=10.0) axB.set_xticks(x) axB.set_xticklabels([f"{s}{' †' if d else ''}\n{p:+.1f} %" for s, d, p in zip(steps, dagger, pct)], - fontsize=9, linespacing=1.7, color=TEXT2) + fontsize=10.6, linespacing=1.7, color=TEXT2) axB.tick_params(axis="x", length=0, pad=7) bare(axB, keep=("left", "bottom")) axB.spines["bottom"].set_color(RULE) axB.set_title(f"the saving does not · dashed = {np.mean(dv):.2f} µs mean", - color=MUTED, fontsize=9, pad=14, loc="left") + color=MUTED, fontsize=10.6, pad=14, loc="left") fig.subplots_adjust(bottom=0.30, top=0.84, left=0.055, right=0.985, wspace=0.26) @@ -1909,19 +1903,19 @@ def fig_variance_rewrite(): solid_capstyle="butt", zorder=3) ax.plot([operand], [y], "o", color=col, ms=9, zorder=5) ax.text(operand * 2.2, y + 0.17, f"operands {mag}", color=TEXT2, - fontsize=8, va="center") - ax.text(operand * 2.2, y - 0.19, note, color=col, fontsize=8, + fontsize=9.4, va="center") + ax.text(operand * 2.2, y - 0.19, note, color=col, fontsize=9.4, va="center", fontweight="bold") ax.axvline(ANS, color=PALE, lw=1.2, ls="--", zorder=4) ax.text(ANS * 1.25, -0.44, "the answer: variance ≈ 2025", color=PALE, - fontsize=7.5, ha="left", va="center") + fontsize=8.8, ha="left", va="center") ax.annotate("", xy=(2.17e7, 1.44), xytext=(ANS, 1.44), arrowprops=dict(arrowstyle="<->", color=MUTED, lw=1.0)) ax.text(2.1e5, 1.52, "4 decades of common term to cancel", color=MUTED, - fontsize=7.5, ha="center") + fontsize=8.8, ha="center") ax.set_yticks([1, 0]) ax.set_yticklabels(["before\naccumulate X", "after\naccumulate Y = X − X₀"], - fontsize=8.5, linespacing=1.5) + fontsize=10.0, linespacing=1.5) for lab, c in zip(ax.get_yticklabels(), (TEXT2, PALE)): lab.set_color(c) ax.tick_params(axis="y", length=0) @@ -1974,7 +1968,7 @@ def fig_measure(): 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") + fontsize=11.2, ha="right", va="center") def union(iv): pts = sorted(iv) @@ -1992,16 +1986,16 @@ def fig_measure(): 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, + ax.text(-6, ys + 0.17, name.split(" ")[0], color=col, fontsize=11.2, 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", + "so the H2D bars\nqueue", color=MUTED, fontsize=11.2, 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", + "shorter than their sum.", color=TEXT2, fontsize=11.2, va="center", linespacing=1.5) ax.plot([-2, frames[-1][3] + D + 2], [2.42, 2.42], color=RULE, lw=0.9, zorder=1) @@ -2010,7 +2004,7 @@ def fig_measure(): 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) + color=MUTED, fontsize=11.2, 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)] @@ -2018,21 +2012,21 @@ def fig_measure(): 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.text(x, v + 1.0, f"{v:.2f}", ha="center", color=TEXT2, fontsize=11.8) 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, + ax2.text(2.42, 33.3, "engine max\nprofiled 32.66", color=MUTED, fontsize=11.2, 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) + color=GREEN, fontsize=11.2, 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_xticklabels([n for n, _, _ in vals], color=TEXT2, fontsize=11.8) + ax2.set_xlim(-0.55, 5.00) 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) + color=PALE, fontsize=11.2, 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") diff --git a/docs/deck/make_figs_kernel.py b/docs/deck/make_figs_kernel.py index c2b9d209..73f04740 100644 --- a/docs/deck/make_figs_kernel.py +++ b/docs/deck/make_figs_kernel.py @@ -17,6 +17,9 @@ from matplotlib.colors import LinearSegmentedColormap from matplotlib.patches import Rectangle, FancyArrowPatch from pathlib import Path +sys.path.insert(0, str(Path(__file__).resolve().parent)) +import deckgate # noqa: E402 -- the projection floor, shared with make_figs.py + OUT = Path(__file__).resolve().parent.parent / "figures" OUT.mkdir(exist_ok=True) @@ -30,7 +33,7 @@ TEXT2 = "#A5B2C4" MUTED = "#6B7A90" plt.rcParams.update({ - "font.family": "DejaVu Sans", "font.size": 9, + "font.family": "DejaVu Sans", "font.size": 11.7, "text.color": PALE, "axes.labelcolor": TEXT2, "xtick.color": TEXT2, "ytick.color": TEXT2, "axes.edgecolor": RULE, "axes.facecolor": "none", @@ -43,11 +46,18 @@ CMAP = LinearSegmentedColormap.from_list( "deck", ["#080C12", "#10202F", ACCENT, AMBER, "#FFF3DC"]) +DPI = 220 + + def save(fig, name): - fig.savefig(OUT / f"{name}.png", dpi=220, transparent=False, + """Write the PNG, then hold it to the SAME projection floor make_figs.py + uses. These three figures were unchecked for a long time and were set at + 7-8 pt as a result -- roughly 8 pt on the screen, half the body size.""" + path = OUT / f"{name}.png" + fig.savefig(path, dpi=DPI, transparent=False, bbox_inches="tight", pad_inches=0.08) + deckgate.check(fig, path, name, DPI) plt.close(fig) - print("wrote", name) def bare(ax, keep=("left", "bottom")): @@ -77,11 +87,11 @@ def fig_frame(): for s in ax.spines.values(): s.set_color(RULE) ax.set_title("one frame, pedestal subtracted · 150×150 crop", - color=MUTED, fontsize=8, pad=7) + color=MUTED, fontsize=10.4, pad=7) cb = fig.colorbar(im, ax=ax, fraction=0.045, pad=0.02) cb.outline.set_edgecolor(RULE) - cb.ax.tick_params(labelsize=7, color=RULE) - cb.set_label("ADU above pedestal", color=MUTED, fontsize=7.5) + cb.ax.tick_params(labelsize=10.0, color=RULE) + cb.set_label("ADU above pedestal", color=MUTED, fontsize=9.8) # Pick a clean, well-isolated charge-sharing event: a local maximum of # moderate amplitude whose 3×3 core carries the charge and whose @@ -118,14 +128,19 @@ def fig_frame(): for dy in (-1, 0, 1): for dx in (-1, 0, 1): v = win[4 + dy, 4 + dx] + # A zoom cell is ~0.25 in wide and a value can be four digits, so + # these cannot clear the projection floor without colliding. They + # are texture: the argument is that the charge is SPREAD, and the + # 3x3 sum that proves it is set below the panel at full size. ax2.text(4 + dx, 4 + dy, f"{v:.0f}", ha="center", va="center", - color=BG if v > 500 else PALE, fontsize=7, + gid="texture", + color=BG if v > 500 else PALE, fontsize=6.4, fontweight="bold" if dx == 0 and dy == 0 else "normal", zorder=6) - ax2.set_title("9×9 zoom on one hit", color=MUTED, fontsize=8, pad=7) + ax2.set_title("9×9 zoom on one hit", color=MUTED, fontsize=10.4, pad=7) ax2.text(4, 9.2, f"3×3 sum = {win[3:6, 3:6].sum():.0f} ADU — one photon.\n" "The peak pixel holds only part of the charge.", - ha="center", va="top", color=TEXT2, fontsize=7.5) + ha="center", va="top", color=TEXT2, fontsize=9.8) save(fig, "fig_frame") @@ -133,11 +148,19 @@ def fig_frame(): def fig_tile(): B, r = 16, 1 # 16×16 block, 3×3 cluster → 1-px halo n = B + 2 * r - fig = plt.figure(figsize=(7.6, 3.1)) + # 6.9 in wide, not 7.6: the deck places this at 8.36 in, so a narrower + # canvas is a LARGER projected figure and larger projected type. The two + # panels sit 0.06 apart rather than 0.135, which is the gap the slide was + # paying for twice -- once inside the picture, once beside it. + fig = plt.figure(figsize=(6.9, 2.62)) - ax = fig.add_axes([0.0, 0.0, 0.44, 1.0]) + # The tile axes holds ONLY the tile. With set_aspect("equal") matplotlib + # shrinks the axes box to honour the aspect, so any data-space room reserved + # underneath for a legend gets squeezed and the legend is clipped. The key + # is drawn in figure coordinates instead, where nothing can take it back. + ax = fig.add_axes([0.015, 0.30, 0.44, 0.70]) ax.set_aspect("equal"); ax.axis("off") - ax.set_xlim(-0.6, n + 0.6); ax.set_ylim(-3.2, n + 1.3) + ax.set_xlim(-0.6, n + 0.6); ax.set_ylim(-0.6, n + 1.5) for i in range(n): for j in range(n): halo = i < r or j < r or i >= n - r or j >= n - r @@ -153,78 +176,88 @@ def fig_tile(): ax.add_patch(Rectangle((tj + r, n - 1 - (ti + r)), 0.92, 0.92, facecolor=AMBER, edgecolor="none")) ax.text(n / 2, n + 0.45, "shared-memory tile · 18 × 18", - ha="center", color=MUTED, fontsize=8) - for y, c, t in [(-1.05, AMBER, "the thread's own pixel"), - (-1.85, ACCENT, "its 3×3 neighbourhood"), - (-2.65, RULE, "halo — loaded, never centred on")]: - ax.add_patch(Rectangle((0, y), 0.7, 0.36, facecolor=c, edgecolor="none")) - ax.text(1.0, y + 0.18, t, va="center", color=TEXT2, fontsize=7.5) + ha="center", color=MUTED, fontsize=10.4) + for k, (c, t) in enumerate([(AMBER, "the thread's own pixel"), + (ACCENT, "its 3×3 neighbourhood"), + (RULE, "halo — loaded, never centred on")]): + fy = 0.20 - k * 0.095 + fig.patches.append(Rectangle((0.020, fy - 0.018), 0.022, 0.042, + transform=fig.transFigure, facecolor=c, + edgecolor="none", zorder=5)) + fig.text(0.055, fy, t, va="center", color=TEXT2, fontsize=9.8) # right: tile cost vs cluster size - ax2 = fig.add_axes([0.575, 0.20, 0.40, 0.62]) + ax2 = fig.add_axes([0.52, 0.20, 0.46, 0.62]) labels = ["3×3\n18×18", "5×5\n20×20", "7×7\n22×22", "9×9\n24×24"] kb = [(16 + 2 * (k // 2)) ** 2 * 4 / 1024 for k in (3, 5, 7, 9)] ax2.bar(np.arange(4), kb, width=0.55, color=ACCENT, zorder=3) for i, v in enumerate(kb): - ax2.text(i, v + 0.12, f"{v:.1f}", ha="center", color=PALE, fontsize=9, + ax2.text(i, v + 0.12, f"{v:.1f}", ha="center", color=PALE, fontsize=11.7, fontweight="bold") ax2.axhline(100, color=PALE, lw=1.2, ls="--") ax2.set_xticks(np.arange(4)); ax2.set_xticklabels(labels, color=TEXT2, - fontsize=8) - ax2.set_ylim(0, 3.4); ax2.set_yticks([]) + fontsize=10.4) + ax2.set_ylim(0, 4.4); ax2.set_yticks([]) bare(ax2, keep=("bottom",)) ax2.set_title("KB of shared memory per 16×16 block (float tile)", - color=MUTED, fontsize=8, pad=8) - ax2.text(3.55, 3.15, "100 KB available per SM on Ada\n" - "— shared memory is never the limit", - ha="right", va="top", color=PALE, fontsize=7.5) + color=MUTED, fontsize=10.4, pad=8) + ax2.text(3.55, 4.3, "100 KB available per SM on Ada\n" + "— shared memory is never the limit", + ha="right", va="top", color=PALE, fontsize=9.8) save(fig, "fig_tile") # ------------------------------------------------- 3. occupancy / registers def fig_occupancy(): - fig, (ax, ax2) = plt.subplots(1, 2, figsize=(11.2, 2.95), - gridspec_kw={"width_ratios": [1, 1.5]}) + # 1.7, not 1.5: the right panel carries three two-line tick labels and they + # were touching. Widening the panel is the fix; shrinking the labels below + # the projection floor is not. + fig, (ax, ax2) = plt.subplots(1, 2, figsize=(11.6, 2.95), + gridspec_kw={"width_ratios": [1, 1.7]}) # left — registers set the occupancy, per cluster size occ = [100.0, 33.3] ax.bar([0, 1], occ, width=0.5, color=[ACCENT, AMBER], zorder=3) for i, o in enumerate(occ): - ax.text(i, o + 3, f"{o:.0f}%", ha="center", color=PALE, fontsize=12, + ax.text(i, o + 3, f"{o:.0f}%", ha="center", color=PALE, fontsize=15.6, fontweight="bold") ax.set_xticks([0, 1]) - ax.set_xticklabels(["3×3 cluster\n38 regs/thread · 6 blocks/SM", - "9×9 cluster\n128 regs/thread · 2 blocks/SM"], - color=TEXT2, fontsize=8.5) + # Three short lines, not two long ones: the left panel is a third of the + # figure and the two labels were running into each other. + ax.set_xticklabels(["3×3 cluster\n38 regs/thread\n6 blocks/SM", + "9×9 cluster\n128 regs/thread\n2 blocks/SM"], + color=TEXT2, fontsize=10.4) ax.set_ylim(0, 122); ax.set_yticks([]) bare(ax, keep=("bottom",)) ax.set_title("achieved occupancy, 16×16 block · f32 build", color=MUTED, - fontsize=8.5, pad=8) + fontsize=11.1, pad=8) # right — block-size sweep, both cluster sizes o3 = [100.0, 100.0, 66.7] o9 = [33.3, 33.3, 0.0] halo3 = [56, 27, 13] - blocks = [f"{b}\nhalo +{h}% of the tile" - for b, h in zip(["8×8 · 64 threads", "16×16 · 256 threads", - "32×32 · 1024 threads"], halo3)] + blocks = [f"{b}\n{t}\nhalo +{h}%" + for b, t, h in zip(["8×8", "16×16", "32×32"], + ["64 threads", "256 threads", "1024 threads"], + halo3)] x = np.arange(3); w = 0.34 ax2.bar(x - w / 2, o3, width=w, color=ACCENT, zorder=3, label="3×3 cluster") ax2.bar(x + w / 2, o9, width=w, color=AMBER, zorder=3, label="9×9 cluster") for xi, (a, b) in enumerate(zip(o3, o9)): ax2.text(xi - w / 2, a + 3, f"{a:.0f}%", ha="center", color=TEXT2, - fontsize=8.5) + fontsize=11.1) ax2.text(xi + w / 2, b + 3, ("will not launch\n(registers)" if b == 0 else f"{b:.0f}%"), - ha="center", va="bottom", color=AMBER if b == 0 else TEXT2, - fontsize=8 if b == 0 else 8.5, + ha="left" if b == 0 else "center", va="bottom", + color=AMBER if b == 0 else TEXT2, + fontsize=10.4 if b == 0 else 11.1, fontweight="bold" if b == 0 else "normal") - ax2.set_xticks(x); ax2.set_xticklabels(blocks, color=TEXT2, fontsize=8.5) - ax2.set_ylim(0, 122); ax2.set_yticks([]) + ax2.set_xticks(x); ax2.set_xticklabels(blocks, color=TEXT2, fontsize=9.8) + ax2.set_ylim(0, 122); ax2.set_xlim(-0.6, 2.85); ax2.set_yticks([]) bare(ax2, keep=("bottom",)) - ax2.legend(frameon=False, fontsize=8.5, labelcolor=TEXT2, loc="upper right") + ax2.legend(frameon=False, fontsize=11.1, labelcolor=TEXT2, loc="upper right") ax2.set_title("occupancy vs block size (halo overhead quoted for 3×3)", - color=MUTED, fontsize=8.5, pad=8) + color=MUTED, fontsize=11.1, pad=8) fig.subplots_adjust(bottom=0.26) save(fig, "fig_occupancy") @@ -234,3 +267,4 @@ if __name__ == "__main__": fig_occupancy() fig_frame() print("done ->", OUT) + deckgate.report() diff --git a/docs/figures/fig_arc.png b/docs/figures/fig_arc.png index 85338339..9edbf7e5 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 5ebb94ba..668beb66 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_cancellation.png b/docs/figures/fig_cancellation.png index e00117a2..0b93b560 100644 Binary files a/docs/figures/fig_cancellation.png and b/docs/figures/fig_cancellation.png differ diff --git a/docs/figures/fig_f32_absolute.png b/docs/figures/fig_f32_absolute.png index 4a710079..a8f80efc 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 3fe07343..4439f808 100644 Binary files a/docs/figures/fig_f32_kernel.png and b/docs/figures/fig_f32_kernel.png differ diff --git a/docs/figures/fig_first_run.png b/docs/figures/fig_first_run.png index 0ad3017e..3109fb90 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_frame.png b/docs/figures/fig_frame.png index 94e4c46e..ddd814ad 100644 Binary files a/docs/figures/fig_frame.png and b/docs/figures/fig_frame.png differ diff --git a/docs/figures/fig_gpu_model.png b/docs/figures/fig_gpu_model.png index 8126e00b..2077eaa6 100644 Binary files a/docs/figures/fig_gpu_model.png and b/docs/figures/fig_gpu_model.png differ diff --git a/docs/figures/fig_graphs.png b/docs/figures/fig_graphs.png index c5a9adc5..3b5994c5 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 index ec2480c1..3ed85dac 100644 Binary files a/docs/figures/fig_measure.png and b/docs/figures/fig_measure.png differ diff --git a/docs/figures/fig_mismatch147.png b/docs/figures/fig_mismatch147.png index 50e526d5..394e4058 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 af330310..6c748369 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 90ebb805..e924474c 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 c5e14d4e..056a288a 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 2f1bdd17..ce77c108 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 c9068d70..cc99c097 100644 Binary files a/docs/figures/fig_overlap.png and b/docs/figures/fig_overlap.png differ diff --git a/docs/figures/fig_overlap_9x9.png b/docs/figures/fig_overlap_9x9.png index b6d847cd..6b544284 100644 Binary files a/docs/figures/fig_overlap_9x9.png and b/docs/figures/fig_overlap_9x9.png differ diff --git a/docs/figures/fig_pagefault.png b/docs/figures/fig_pagefault.png index 8a30cb35..adc01d34 100644 Binary files a/docs/figures/fig_pagefault.png and b/docs/figures/fig_pagefault.png differ diff --git a/docs/figures/fig_pedtiming.png b/docs/figures/fig_pedtiming.png index ebb178c3..12889900 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 df96c5d1..2d6489f8 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 4c8a1f64..6691ec2c 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 2a4c0c45..9eae4440 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 0dbd369a..2a6a62a7 100644 Binary files a/docs/figures/fig_streams.png and b/docs/figures/fig_streams.png differ diff --git a/docs/figures/fig_test3.png b/docs/figures/fig_test3.png index 6478a1eb..0c13b35b 100644 Binary files a/docs/figures/fig_test3.png and b/docs/figures/fig_test3.png differ diff --git a/docs/figures/fig_tile.png b/docs/figures/fig_tile.png index e073d302..2b465b8e 100644 Binary files a/docs/figures/fig_tile.png and b/docs/figures/fig_tile.png differ diff --git a/docs/figures/fig_varfloor.png b/docs/figures/fig_varfloor.png index b61b33ba..e44730c5 100644 Binary files a/docs/figures/fig_varfloor.png and b/docs/figures/fig_varfloor.png differ