CUDA Profiling and Performance Measurement The notes up to this point covered what makes CUDA kernels fast in theory: coalescing, tiling, occupancy, and asynchronous execution. This note covers how to measure whether a kernel is actually fast and how to find the bottleneck when it is not. Profiling is what separates guessing from optimization. Without it, you can apply every best practice and still be slower than a simpler implementation. ...
CUDA Unified Memory: One Pointer, Two Processors The previous note covered explicit CUDA streams and how to overlap host-device transfers with computation. This note covers Unified Memory, an alternative to explicit cudaMalloc / cudaMemcpy / cudaFree. It lets both CPU and GPU use a single pointer, at the cost of some performance complexity under the hood. Unified Memory is seductive: it removes the need to think about host and device pointers separately. But it does not remove the physical reality that CPU and GPU memory are separate. Understanding when it helps and when it hurts is essential. ...
CUDA Reduction: Parallel Sum, Max, and Other Tree Algorithms The previous note covered occupancy. This note covers reduction: taking a large array and producing a single value (a sum, max, min, product, or any associative binary operation). Reduction is one of the most important parallel algorithms on a GPU, and it teaches several key CUDA ideas at once: shared memory, thread cooperation, warp-level execution, and multi-kernel launches. This note builds the algorithm step by step, from a naive atomic version to a tree-based shared-memory reduction. ...
CUDA Streams and Host/Device Overlap The previous note covered reduction, the first parallel algorithm on the GPU side. This note moves to the host-device boundary: how to keep the GPU busy while the CPU prepares the next batch of data, and how to overlap data movement with computation. So far, every kernel launch and every cudaMemcpy has been synchronous. The host waits, the GPU works, the host waits again. CUDA streams let you break that pattern. ...
CUDA Occupancy: Filling the SM to Hide Latency The previous note showed how to use shared memory to fix uncoalesced access patterns. This note covers occupancy: how many warps can live on an SM at the same time, and why that number determines whether the GPU can hide memory latency. Occupancy is the first CUDA performance topic that is not about memory access at all. It is about keeping the warp schedulers busy. ...
CUDA Shared Memory Tiling The previous note showed that coalesced memory access is the key to memory-bound kernel performance. But some algorithms naturally want uncoalesced access patterns. The fix is shared-memory tiling: load data into fast on-chip SRAM in a coalesced way, rearrange it there, then write it back to global memory in a coalesced way. This note explains the tiling pattern, why __syncthreads() is essential, the bank-conflict problem, and how to choose tile sizes. ...
CUDA Memory Coalescing: Why Access Patterns Matter The previous note showed how to write and launch a CUDA kernel. This note answers the first performance question: why does the same kernel sometimes run 5× or 10× slower just because of how threads read memory? The answer is memory coalescing. Because GPU memory is optimized for bandwidth over latency, the hardware rewards access patterns where threads in a warp read or write contiguous addresses together. If they do not, the GPU wastes bandwidth and cycles fetching data that most threads ignore. ...
CUDA Programming Model: Writing and Launching Kernels The previous note covered the hierarchy: thread, warp, block, grid, and how those map to SMs. This note turns that hierarchy into actual code. By the end, you will have seen a complete CUDA program, understood every line, and know how a host program hands work to the GPU. CUDA is an extension of C/C++. Most of the code you write is ordinary C++. A small number of CUDA-specific pieces — kernel functions, the launch syntax, a few memory APIs, and built-in thread indices — turn a sequential program into a massively parallel one. ...
CUDA Thread Hierarchy: Grids, Blocks, Warps, and Threads The previous note mapped the hardware: SM, warp scheduler, register file, shared memory, L2, device memory. This note maps the software abstraction CUDA exposes on top of that hardware. The two maps fit together almost one-to-one, and once you see how, most of CUDA stops being arbitrary syntax and becomes named hardware concepts. CUDA organizes parallel work into four nested levels: Grid └── Block (threads in a block can share fast on-chip memory, and sync) └── Warp (32 threads, executed in lockstep on the hardware) └── Thread (your kernel code, from one thread's point of view) This note works from the bottom up: start with the thread (the thing you actually program), then warp, then block, then grid, then show how the whole tower maps onto the GPU die. ...
Anatomy of a GPU: Hardware Components The previous note covered why GPUs look the way they do — throughput over latency, thousands of simple cores instead of a few smart ones. This note zooms into the actual silicon: what physically sits on a GPU die, what each piece is called, and what job it does. The goal is to have concrete hardware nouns (SM, warp scheduler, register file, L2, …) in hand before those same nouns start showing up as CUDA concepts (threadIdx, __shared__, occupancy, …). ...