CUDA专家手册：GPU编程权威指南 (高性能计算技术丛书) (（美）Nicholas Wilt) (z-library.sk, 1lib.sk, z-lib.sk)

【One-Line Pitch】 A deep, architecture-first guide to CUDA for developers who already write GPU code and want to understand why it runs fast or slow—covering hardware internals, memory models, concurrency, and advanced optimization patterns with concrete, measurable examples. 【Book Arc】 - **Opening (~0%–11%)**: Sets the stage with GPU history and the shift to heterogeneous CPU+GPU computing, then lays out the book’s three-part structure: hardware/software architecture, programming model, and advanced case studies. Establishes why bottleneck analysis must precede optimization. - **Early (~11%–33%)**: Dives into hardware foundations—virtual memory, page tables, TLB behavior, and the host interface—then explains the CUDA runtime vs. driver API split, context management (including the floating-context trick for multi-GPU), and introduces streams/events as the coarse-grained concurrency mechanism. Also covers host memory types (pageable vs. pinned) and why pinned memory enables async copies. - **Middle (~33%–56%)**: Moves into the programming model proper: global/static/shared memory allocation, cache behavior across SM 2.x and 3.x, warp-level memory access patterns, and the cost of CPU/GPU synchronization. Includes practical measurement experiments (e.g., null kernel launch overhead) and a first look at reduction kernels—the book’s recurring example. - **Late (~56%–78%)**: Explores the SM instruction set (SASS), special function units (SFU) and their precision trade-offs, then covers multi-GPU programming via peer-to-peer addressing and copy, plus texture and surface memory in depth—including pitch, normalized coordinates, and the new-in-Fermi surface read/write intrinsics. - **Ending (~78%–100%)**: Focuses on advanced algorithm patterns: warp-level reductions using shuffles and atomics, parallel scan (prefix sum) with circuit diagrams and recursive implementations, and a full case study on normalized cross-correlation for image matching—showing how to combine texture cache, shared memory, and tiling for real performance gains. 【Key Takeaways】 - **Bottleneck analysis before optimization** (Opening): Profile to find the actual hotspot (GPU kernel vs. data transfer) before touching code; optimizing the wrong part wastes effort. Use timing code or profilers to quantify first. - **Pinned memory is the gateway to async transfers** (Early): Pageable memory can’t be accessed by the GPU directly; CUDA uses staging buffers. Pinned (page-locked) memory enables faster copies, async operations, and even direct kernel access via mapped memory—critical for hiding transfer latency. - **Streams and events are the coarse-grained concurrency toolkit** (Early): They let you overlap CPU/GPU work, DMA copies with kernel execution, and even run multiple kernels or GPUs in parallel. Operations within a stream are serial; concurrency comes from using multiple streams. - **Context management is a hidden performance lever** (Early): The driver API’s push/pop context stack lets a single thread drive multiple GPUs or libraries create “floating” contexts without knowing their caller. On WDDM (Windows Vista+), popping a context with pending commands triggers a kernel thunk—avoid it for speed. - **Cache behavior varies by architecture—know your SM** (Middle): SM 2.x caches global memory in L1; SM 3.x doesn’t (use `__ldg()` or `const __restrict__` to route through texture cache). Shared memory vs. L1 cache split is configurable on SM 3.x. These details change which optimization works. - **Synchronization is expensive—measure it** (Middle): Full CPU/GPU syncs (e.g., synchronous memcpy) can cost ~8 microseconds per call on real hardware. Async operations and careful stream usage hide driver overhead; even in CUDA 1.0, kernel launches were async. - **Warp-level primitives beat shared memory for reductions** (Late): Shuffle instructions, atomics, and memory fences let you implement single-pass reductions without multiple kernel launches or shared memory bank conflicts. This is the modern pattern for high-performance reductions. - **Texture and surface memory are specialized tools, not defaults** (Late): Textures offer hardware filtering, normalized coordinates, and cache-friendly access for read-only data; surfaces (SM 2.x+) add write access to CUDA arrays. Use them when the access pattern fits—e.g., image processing with overfetch and boundary handling. 【Reading Tips】 - **Skim the historical intro (Ch. 1) and jump to Ch. 2–3 for architecture**: The early GPU history is interesting but not actionable. The hardware diagrams and memory-model explanations in Ch. 2–3 are the foundation for everything later—read them carefully. - **Deep-read Ch. 5 (memory) and Ch. 6 (streams/events)**: These are the most practically useful chapters for everyday CUDA coding. The pinned-memory discussion and the concurrency experiments (e.g., `concurrencyMemcpyKernel.cu`) directly translate to real performance wins. - **Treat Ch. 8 (SM instruction set) as a reference, not a cover-to-cover read**: The SASS-level details and SFU precision tables are valuable when you’re optimizing hot loops, but you don’t need them upfront. Bookmark the tables and come back when profiling points to arithmetic bottlenecks. - **Work through the reduction and scan examples (Ch. 12–13) with code in hand**: These are the book’s core algorithmic case studies. The progression from simple to optimized versions (atomics, shuffles, zero-padding, templates) is the best way to internalize the optimization mindset. - **Expect some dated material**: The book references CUDA 5.0, Fermi/Kepler architectures, and AWS EC2 instances from that era. Modern GPUs (Volta+) have changed some details (e.g., independent thread scheduling, different cache hierarchies), so cross-check with current NVIDIA docs for the latest hardware. 【Coverage Limits】 This guide synthesizes the provided excerpts, which cover roughly the first 15 chapters and the glossary. The book’s later chapters (beyond Ch. 15) and any appendices are not represented here.