Memory Coalescing for Global Memory Access

1. The "Why" Behind Coalescing

Imagine you need to retrieve 32 different books from a large library. It would be far more efficient to grab them all from a single shelf in one trip rather than running to 32 different aisles. The GPU's memory system operates on a similar principle.

The GPU executes threads in groups of 32 called warps. When a warp needs to read data from global memory, the hardware doesn't fetch data one byte at a time for each thread. Instead, it performs large memory transactions to fetch entire segments of memory (e.g., 32 or 128 bytes at a time).

Memory coalescing occurs when the 32 threads in a warp access consecutive locations in global memory. The GPU can then service all 32 requests with one or very few memory transactions. Conversely, if threads access scattered, non-consecutive memory locations (a "strided" or random access pattern), the GPU is forced to issue many separate transactions, wasting bandwidth and time.

To see the dramatic impact this can have, let's start with a short video that demonstrates a 4.5x performance difference between two kernels that only differ in their memory access pattern.

The official CUDA Best Practices Guide provides the technical definition and diagrams for this concept. It's a critical reference for understanding coalescing.

As the guide explains, for modern GPUs (Compute Capability 6.0+), the rule is simple: the hardware will issue the minimum number of 32-byte transactions required to fulfill the requests of all threads in the warp. Your goal is to make your memory access patterns fit neatly within these transactions.

2. Identifying Coalesced and Uncoalesced Access Patterns

To apply this knowledge, you need to be able to spot these patterns in your code. Given your background in C++, you know that multi-dimensional arrays are typically stored in memory in row-major order. This means that elements in the same row are contiguous in memory.

This has direct implications for CUDA kernels. Consider a kernel processing a 2D matrix float* A of width W:

Coalesced Access:
Threads access consecutive elements along a row. This is the ideal pattern.

int tid = blockIdx.x * blockDim.x + threadIdx.x;
float value = A[tid]; // Perfect coalescing!

When processing a 2D matrix, this translates to threads with consecutive threadIdx.x values accessing consecutive columns in the same row.

Uncoalesced (Strided) Access:
Threads access elements separated by a fixed stride. This is a common anti-pattern, especially when accessing columns of a matrix.

int tid = blockIdx.x * blockDim.x + threadIdx.x;
int stride = W; // Width of the matrix
float value = A[tid * stride]; // Uncoalesced access!

Each thread jumps stride elements past the previous one. If stride is large, each thread's access will fall into a different memory transaction, leading to maximum inefficiency.

The matrix multiplication example from the first video (LINK) is a perfect illustration. Let's analyze it more closely.

3. Strategies for Achieving Coalescing

When you identify an uncoalesced access pattern, you generally have two main strategies to fix it: change the memory access pattern in your kernel, or change the data layout in memory.

Strategy 1: Change the Access Pattern

Sometimes, a simple change to how threads are assigned to data can solve the problem. In the matrix multiplication example, swapping how thread indices (x, y) map to the output matrix C can change which input matrix (A or B) is accessed with a stride.

Strategy 2: Change the Data Layout

This is often a more robust solution. Instead of forcing your kernel to work with an inefficient layout, you change the layout itself.

Matrix Transposition
If your algorithm requires column-wise access to a matrix, a very effective strategy is to transpose the matrix in global memory beforehand. After transposition, the columns become rows, and accessing them becomes a perfectly coalesced operation.

Structure of Arrays (SoA) vs. Array of Structures (AoS)

This is a classic data layout problem that will be familiar from your C++ experience. Suppose you have data for particles, each with a position (x, y, z) and mass (m).

• Array of Structures (AoS): struct Particle { float x, y, z, m; }; Particle particles[N];

  • Memory layout: x0, y0, z0, m0, x1, y1, z1, m1, ...
  • If a warp of threads tries to read only the x component of all particles, the access is strided and uncoalesced.

• Structure of Arrays (SoA): struct Particles { float x[N], y[N], z[N], m[N]; };

  • Memory layout: x0, x1, ..., xN, y0, y1, ..., yN, ...
  • Now, if a warp reads the x components, the access is perfectly coalesced.

For CUDA, the SoA pattern is almost always preferred for achieving good memory performance.

Strategy 3 (Preview): Use Shared Memory

A powerful technique, which we will study in depth in the next lesson, is to use the on-chip shared memory as a user-managed cache. The pattern is:

1. A block of threads collectively loads a "tile" of data from global memory into shared memory. This load is carefully arranged to be coalesced.
2. The threads then perform their computations using the data from the extremely fast shared memory, accessing it in any pattern (e.g., strided) without penalty.

This combines the best of both worlds: efficient, high-bandwidth access from global memory and flexible, low-latency access on-chip.

4. Advanced Topic: Vectorized Memory Access

We can take coalescing a step further. Instead of having each thread load a single 4-byte float, we can have each thread load a larger chunk, like a 16-byte float4. This is called vectorized memory access. It allows the GPU to use wider load instructions (e.g., 128-bit loads instead of 32-bit loads), further increasing the memory throughput.

This requires your data to be properly aligned in memory. Since you have a C++ background, you'll appreciate the nuance of how this is communicated to the compiler.

This is an advanced but powerful technique for squeezing maximum bandwidth out of the memory system.

Conclusion

You've just learned one of the most important concepts in CUDA optimization. Uncoalesced memory access is one of the most common performance killers in beginner CUDA code, and now you know how to diagnose and fix it.

Key Takeaways:

• Coalescing is Key: Global memory is accessed in wide transactions. Coalesced access occurs when threads in a warp request consecutive memory locations, maximizing bandwidth.
• Row-Major Dominates: Accessing data along a row is fast; accessing it down a column is slow (strided and uncoalesced).
• Identify and Fix: Watch for strided access patterns in your kernels, especially with multi-dimensional data.
• Solution Toolkit:
  • Change the access pattern in the kernel.
  • Change the data layout (transpose matrices, use SoA instead of AoS).
  • (Preview) Use shared memory to stage data.
• Vectorize for More Speed: Use vector types like float4 to enable wider, more efficient memory instructions.

In our next lesson, we will dive deep into using shared memory to reduce global memory traffic. You've already seen it previewed here as a way to fix uncoalesced access; we will now explore its full potential as a programmable cache to enable complex, high-performance computations.