Coalesced vs. Uncoalesced Memory Access

A Tale of Two Kernels: The Copy Experiment

To isolate the effect of memory access patterns, we will use one of the simplest possible operations: a data copy. Since this operation does minimal computation, its performance is almost entirely bound by memory bandwidth. This makes it a perfect laboratory for studying memory access efficiency.

We will write two kernels:

1. copyCoalesced: Each thread copies one element from an input array to an output array using a direct, one-to-one mapping. This represents the ideal access pattern.
2. copyUncoalesced: Each thread copies an element, but the access to the input array will be strided, deliberately violating the principle of coalescing.

The following resource provides a complete, self-contained C++/CUDA program that implements this experiment.

Analyzing the Access Patterns

Let's dissect the two kernels from the reading to solidify why one is coalesced and the other is not. Both kernels use the standard calculation for a global thread ID:

int index = blockIdx.x * blockDim.x + threadIdx.x;

The Coalesced Kernel

__global__ void copyDataCoalesced(float *in, float *out, int n) {
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < n) {
        out[index] = in[index];
    }
}

Here, thread t within a warp accesses in[base + t], and thread t+1 accesses in[base + t+1]. Their memory requests are for adjacent 4-byte floats. The hardware's Load/Store Unit sees that the 32 threads in the warp are requesting a perfectly contiguous 128-byte block of memory. It can service all 32 requests with a single 128-byte memory transaction (or four 32-byte transactions), achieving 100% bus efficiency.

The Uncoalesced Kernel

__global__ void copyDataNonCoalesced(float *in, float *out, int n) {
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < n) {
        out[index] = in[(index * 2) % n];
    }
}

The write to out[index] is still perfectly coalesced, but the read from in is not. Let's look at the addresses requested by two consecutive threads, t and t+1:

• Thread t requests address corresponding to in[(base + t) * 2].
• Thread t+1 requests address corresponding to in[(base + t+1) * 2] = in[(base * 2) + 2t + 2].

The memory locations they request are separated by 2 * sizeof(float) = 8 bytes. This is a strided access. To satisfy the requests for the whole warp, the hardware must issue more, smaller memory transactions. For each 32-byte sector it fetches, only a fraction of the data is actually used by the threads in the warp. This wasted bandwidth is the direct cause of the performance degradation.

This simple strided copy is a proxy for a very common real-world problem: accessing columns of a matrix stored in row-major format.

Measuring Performance with Nsight Compute

To quantify the performance difference, we will use NVIDIA Nsight Compute (ncu), the command-line profiler. This tool executes your program, collects detailed performance data from the GPU's hardware counters for each kernel launch, and presents a report. Given your goal of becoming a kernel engineer, proficiency with ncu is non-negotiable.

After compiling the code from the reading into an executable (e.g., benchmark), you would run the profiler like this:

ncu --metrics dram__throughput.avg.pct_of_peak_sustained_elapsed,l1tex__throughput.avg.pct_of_peak_sustained_elapsed,gpu__time_duration.sum ./benchmark

Let's break down the requested metrics:

• gpu__time_duration.sum: The total time the kernel spent running, in microseconds. This is our primary measure of performance.
• dram__throughput.avg.pct_of_peak_sustained_elapsed: This tells us what percentage of the GPU's theoretical peak global memory (DRAM) bandwidth was actually utilized.
• l1tex__throughput.avg.pct_of_peak_sustained_elapsed: The percentage of the L1 cache bandwidth utilized. In the uncoalesced case, more requests are generated, but they are inefficient, leading to lower L1 throughput as seen in the resource's results.

When you run this, ncu will output a detailed report for each kernel. The copyDataCoalesced kernel will show a shorter duration and higher memory throughput percentages compared to copyDataNonCoalesced. You have now empirically proven the performance cost of uncoalesced memory access.

From Simple Copy to Matrix Multiplication

The simple copy experiment is a building block for understanding more complex kernels. The principles are identical. Let's revisit the naive matrix multiplication kernel. The core of the operation is:
tmp += A[row * K + i] * B[i * N + col];

Consider a thread block where consecutive threadIdx.x values correspond to changing col.

• Accessing A: All threads in a warp have the same row but different col. They will all access A[row * K + i] for a given i. This is a broadcast, which is efficient.
• Accessing B: Threads in a warp have different col values. Thread t accesses B[i * N + col_t] and thread t+1 accesses B[i * N + col_{t+1}]. Since col_t and col_{t+1} are consecutive, the access to matrix B is coalesced.

Now, what if we mapped consecutive threadIdx.x to changing row instead? This is the central idea behind the performance difference in the videos from our last lesson. The following article details exactly how to make this change in code.

The change seems subtle—just a re-indexing of threads—but as you now understand from our copy experiment, the impact on how the hardware fetches memory is profound. This is the essence of GPU kernel optimization: aligning your algorithm with the realities of the hardware.

Conclusion

In this lesson, you have taken a crucial step from theory to practice. You now have a concrete, measurable understanding of memory coalescing.

Key Takeaways:

• Practice Validates Theory: We implemented both a coalesced and an uncoalesced kernel to directly observe the performance difference predicted by theory.
• Profiling is Essential: You were introduced to NVIDIA's Nsight Compute (ncu) as the primary tool for measuring kernel performance and identifying bottlenecks like poor memory access.
• Strided Access is the Enemy: The most common form of uncoalesced access is strided access, where consecutive threads access memory with a fixed gap between them. This wastes memory bandwidth and is a primary target for optimization.
• Indexing is Key: Achieving coalesced access is a matter of correctly mapping thread indices to memory indices so that consecutive threads access contiguous data.

We have successfully improved how efficiently we use the available global memory bandwidth. However, global memory is slow, even at 100% efficiency. The next level of optimization is to avoid accessing it whenever possible.

In our next lesson, you will learn about shared memory. This is a small, programmer-managed cache on each Streaming Multiprocessor that is orders of magnitude faster than global memory. You will learn how to use it to explicitly cache data, enabling data reuse within a thread block and dramatically reducing the need for costly global memory traffic.