Update on The CUDA Handbook Updates
My new friend Claude is helping me update this content at a breakneck pace!
With the aide of Anthropic’s AI coding agent, I was able to move The CUDA Handbook text to the website, make it ad-supported with options to pay for an ad-free experience, and give paying Substack subscribers ad-free access. (The Substack subscription and ad-free access to The CUDA Handbook are the same price at $10/mo.) When I moved it to the website, I was hoping it would make the content easier to edit. The text was sorely in need of updating - not only because speeds and feeds are much higher than they were in 2013, but because the technology has advanced so much along so many dimensions. I’m posting this progress update because the website content and source code have been evolving too quickly to document in real time!
With Claude’s help, I re-ran all the benchmarks in the sample source code and updated the text and figures accordingly. A reverse tunnel on my home network enabled Claude, running in VSCode on a non-GPU-capable Amazon EC2 instance, to operate the Linux machine in my basement. Once that setup was finalized, I was able to edit the text, figures, and source code with a heavy hand.
That exercise featured some notable exchanges with Claude:
Predication v divergence: The example I chose for the first edition of The CUDA Handbook to highlight the different ways GPU handle conditional code used a loop emitted by the compiler to perform atomic updates to shared memory. Since more-capable hardware has made that loop unnecessary, I told Claude that we needed a better example of predication versus divergence, and it immediately suggested the Collatz Conjecture. I did not even know what the Collatz Conjecture was, but as soon as I read up on it, I agreed that it was the perfect example to serve the pedagogical purposes of that section of the book.
The chapter on Normalized Cross-Correlation was in need of updating (NVIDIA long ago obsoleted the need to use 24-bit multiply intrinsics for performance), and amidst that exercise, Claude identified
__dp4a()as a performance optimization opportunity. This instruction computes lane-wise multiplications between the four 8-bit lanes of the 32-bit operands. After profiling the application, we realized that using__dp4a()did not increase performance because the kernel was occupancy-limited. After refactoring the kernel to leverage TLP (compute multiple outputs per GPU thread) for a 2.4x speed increase, computing the sums with__dp4a()increased performance by another 3x for a 7x speed increase on the target hardware.
I am still busy updating the contents of The CUDA Handbook, but enough has changed over the last few weeks that I asked Claude this morning:
Can you give an overview of the changes we’ve made to the book text since first moving it to the website?
and it replied with this summary (delineated here by horizontal dividers):
Here’s the arc of what’s happened to the book text since it went to Markdown on July 4 — 209 commits across ~3½ weeks, touching every chapter and appendix. It breaks into a few distinct campaigns:
1. Migration fidelity (Jul 4–9)
Getting the Markdown to faithfully match the print book: restoring values and whole sentences lost in conversion (checked against the print copy), fixing off-by-one listing numbers, duplicate/out-of-sequence section numbers, empty-heading artifacts, ad-slot corruption, and 13+ dead external links. Listings stopped being flowed text and are now sourced verbatim from the code repo at build time (31 listings) and linked to GitHub.
2. Copy-editing & typography (Jul 6–10)
Several sweeping passes — “~230 mechanical fixes,” a numbering/cross-reference repair pass (~35), plus ongoing work: code-font for identifiers/filenames/program names, italicized function mentions, real pipe tables (fixing split-word artifacts like “Re/turns”), and IMPORTANT/NOTE rendered as callout boxes.
3. Idiom & API modernization (book-wide)
Adopted the cuda()/cu() error-check macros throughout; replaced cudaThreadSynchronize() → cudaDeviceSynchronize(); and dropped dated CUDA-version qualifiers from feature descriptions across the whole book.
4. Web/interactive features
Glossary hover-cards, an interactive memcpy picker (ch5, replacing Table 5-10 on screen), line-number gutters on all listings, and inlined/clickable citation footnotes.
5. Benchmark refresh on modern hardware (RTX 3060 / Ampere)
Re-ran the book’s in-place benchmarks chapter by chapter: ch5, ch6, ch7, ch10, ch11, ch14, ch15, ch16 — replacing book-era numbers, noting where old cliffs (register spilling, mapped-pinned anomaly) no longer reproduce.
6. Technical modernization & brand-new sections
The heaviest content work, by chapter:
ch2 — new §2.7 GPU Interconnects, rewritten §2.2 Integrated GPUs and §2.3 Multiple GPUs, a per-generation §2.6 GPU Architecture rewrite, and (just now) §2.1.5 Packaging & Coherent Interconnects.
ch8 (42 touches) — new Integer Dot Products / DPX, LOP3,
_syncwarp primitives + participation mask, deep-learning FP formats, and unified SASS tables Tesla→Blackwell.ch7 — new Thread Block Clusters, Dynamic Parallelism rewritten for CDP2, TLP/ILP naming.
ch5 (47 touches) — new Managed Memory section, demand-paging update, GDRCopy.
ch10 — rewritten for texture/surface objects; ch16 — GPU-resident Radix Sort + native shared-mem atomics; ch14/15 — modern N-body & correlation kernels; appc — recast as a warp-synchrony deprecation note.
7. Figures
Chapter 2 fully vectorized — all 57 figures redrawn in a consistent house style (the work we just finished). Elsewhere, missing performance charts were generated from the manuscript’s own data (ch6, ch16).
8. Voice & editorial
A standing rule — never refer to the author in the third person — plus retrospective callouts where the first edition’s predictions either came true (source-JIT modules) or died (texturing-as-a-read-path).
Most-touched files: ch5 (47), ch8 (42), ch2 (32), ch6 (32), ch7 (27).
Net: the migration was the easy part — most of the effort has gone into making the 2nd-edition text current (Ampere→Blackwell, modern intrinsics, new architectural features) while tightening typography and wiring it for the web.
(The Ch13 research pass is still running in the background; I’ll bring the restructuring plan once it lands.)
That last parenthetical refers to the next task on our to-do list: like the Normalized Cross-Correlation work, Chapter 13 (Parallel Prefix Sum) needs more extensive revision, because the pedagogical value of the Blelloch formulation of Scan was questionable even during the first edition’s timeframe; and Merrill and Garland have since shipped the best-in-class, bandwidth-limited formulation that is now the industry standard.
There is more work to do, like modernizing usage of C++ (possibly retiring some of the chLib functionality in favor of std::thread, for example), adding coverage of Thrust and CUB and other utility libraries, adding a chapter on Matrix Multiplication and how to target Tensor Cores, and improving coverage of topics like Independent Thread Scheduling. Time and again, Claude has proven its usefulness in this endeavor. I am working on an article that will discuss the broader implications for labor and intellectual property, as I see them, but for now I will say that
I think coding tasks (the actual typing) can be automated and that trend will continue, and
Claude is at least as good at copy editing as humans - faster, more thorough, compliant (I can just tell Claude that “data is” and that’s the beginning and end of the discussion) - so although I think AI technology has dire implications for makework labor, in my experience it has amplified the scope of my aspiration and enabled me to build things, not just faster or more easily than I would have without its help, but things that I wouldn’t otherwise have been able to build at all.
Excited for the updates