GPU Computing · Cluster
Kernel Programming
Launch geometry, synchronization, and patterns.
- 01 Thread Hierarchy and Global Index Arithmetic in a Kernel Launch How a grid of blocks and threads is flattened into buffer indices, where the arithmetic overflows, and what the launch dimensions deliberately do not guarantee.
- 02 Tail Handling and Bounds Checks in Partial Tiles Why the last block of a launch is the one that returns early, how predicated access differs from an early return, and what goes wrong when barriers and zero-length buffers meet an oversized grid.
- 03 Grid-Stride Loops and Launch Size Independence Why the grid-stride loop survives shrunken launches and changing problem sizes, what it costs against a one-element-per-thread kernel, and when it is the wrong shape.
- 04 Block Barriers and Divergent Synchronization What __syncthreads() actually orders within a block, why a barrier reached by only some threads is undefined, and how warp-level primitives replace it where blocks are too coarse.
- 05 Tiled Stencils and Shared Memory Halos How an overlapping tile plus a barrier turns repeated global reads into repeated shared-memory reads, how halo overhead scales with tile size, and what register blocking adds.
- 06 Tree Reductions and Floating-Point Summation Order Why a pairwise tree sum is not the same number as a linear sum, how much error the tree removes, and how shuffle-based and shared-memory reduction stages compose.
- 07 Single-Pass Prefix Scan with Decoupled Look-Back How a scan produces every prefix in one launch without a chain of kernels, what each block publishes and waits for, and why the algorithm stays linear in total work.
- 08 Atomics and Contention in Counters and Aggregates Why a single global atomic counter serialises a kernel, how per-warp and per-block aggregation cuts that traffic, and which failure modes survive the optimisation.
- 09 Streams, Events, and Concurrent Kernel Execution What a stream actually orders, why two kernels can overlap or fail to, and how events turn a stream of launches into a dependency graph.
- 10 Grid-Wide Barriers and Persistent Work Queues How a barrier that spans every block is built from atomics and counters, what residency it requires, and how a persistent grid replaces the launch queue with a device-side one.