Skip to main content

2 docs tagged with "coalescing"

View all tags

Global Memory and Coalescing

Thread Indexing establishes that the fastest-varying array index should track threadIdx.x, so that adjacent threads in a warp touch adjacent addresses. This page is the mechanism that makes that rule matter: how the hardware actually turns a warp's 32 addresses into memory transactions, and why the difference between "adjacent" and "scattered" can be an 8x difference in delivered bandwidth for the exact same amount of useful data.

Memory Access Optimization

Global Memory and Coalescing establishes the 32-byte-sector model and works out the arithmetic for coalesced, strided, and AoS access patterns. This page turns that model into an ordered checklist for a memory-bound kernel: which lever to pull first, which to skip, and why the order matters more than any individual technique.