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 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. ...