Optimizing Parallel Reduction with Shared Memory and Sequential Addressing

From Divergence to Bank Conflicts

The primary villain of our previous kernel was branch divergence. To eliminate it, we must change how threads select their work. Instead of using a divergent if condition to decide if a thread works, we can have more threads participate but compute a different index to work on.

Let's examine an improved version of the interleaved addressing loop. This corresponds to the "REDUCE-1" step in the optimization guide we've been following.

The new loop logic looks like this:

for(unsigned int s = 1; s < blockDim.x; s *= 2) {
    int index = 2 * s * tid;
    if (index + s < blockDim.x) { // Non-divergent bounds check
        sdata[index] += sdata[index + s];
    }
    __syncthreads();
}

This is a significant improvement. The expensive modulo (%) operator is gone. More importantly, the if condition is now a simple bounds check. In the early stages of the loop, most threads within a warp will pass this check, following the same execution path and avoiding divergence.

However, we've traded one problem for another. This new strided access pattern, index = 2 * s * tid, creates a new, more subtle bottleneck: shared memory bank conflicts.

Understanding Shared Memory Banks

To grasp why this new pattern is still suboptimal, we need to look deeper at the hardware. Shared memory is not a single, monolithic block of RAM. It is physically partitioned into a number of smaller, independent memory modules called banks. On modern NVIDIA GPUs, there are typically 32 banks.

Successive 4-byte words in shared memory are mapped to successive banks.

• Word 0 maps to Bank 0.
• Word 1 maps to Bank 1.
• ...
• Word 31 maps to Bank 31.
• Word 32 maps back to Bank 0.
• In general, word[i] maps to bank[i % 32].

A bank can only service one memory request from a warp per clock cycle. If multiple threads in a warp try to access different addresses that fall within the same bank, those requests are serialized. This serialization is a bank conflict, and it negates the speed advantage of shared memory.

The following video provides an excellent and detailed explanation of warps, the bank structure, and what constitutes a bank conflict.

Now, let's analyze our reduce1 kernel's access pattern: sdata[2 * s * tid].
Consider the first iteration where s=1. The access index is 2 * tid.

• threadIdx.x = 0 accesses sdata[0], which is in bank 0.
• threadIdx.x = 1 accesses sdata[2], which is in bank 2.
• threadIdx.x = 15 accesses sdata[30], which is in bank 30.
• threadIdx.x = 16 accesses sdata[32]. Since 32 % 32 = 0, this is bank 0.

Here is the conflict: thread 0 and thread 16 both need to access bank 0 simultaneously. The hardware must serialize these two requests, making the operation take twice as long. This is a "2-way" bank conflict. Since this pattern repeats across all even-numbered banks, the entire memory access for the warp is slowed down.

The diagram below, based on the NVIDIA slides, illustrates this strided, conflict-prone access pattern.

The Solution: Sequential Addressing

To eliminate bank conflicts, we must ensure that threads in a warp access different banks. The most effective way to do this is to have them access consecutive memory locations. This is called sequential addressing.

Instead of a strided access where all threads are active, we revert to a model where only a subset of threads are active. But this time, we do it smartly.
The new logic is: if (tid < s) { sdata[tid] += sdata[tid + s]; }.

Let's analyze this pattern:

• The loop now counts down the stride: for(unsigned int s = blockDim.x/2; s > 0; s >>= 1).
• In each step, the first s threads are active (if (tid < s)).
• Each active thread tid accesses sdata[tid] and sdata[tid + s].

Since the active threads (0, 1, 2, ..., s-1) have consecutive IDs, their access to sdata[tid] is perfectly sequential. Each thread in the first warp (threads 0-31) accesses a unique bank (banks 0-31). This is conflict-free. Similarly, their access to sdata[tid + s] is also sequential and thus conflict-free. This simple change completely resolves the bank conflict issue.

The resources below explain this superior approach, which corresponds to "REDUCE-2" in the Medium article.

The difference in the access pattern is striking.

Now, let's see the full implementation in the Medium article.

By adopting sequential addressing, we have now solved both of the major problems from our initial implementation:

1. Warp Divergence: The if (tid < s) condition is not divergent for most of the reduction. For a large s, all 32 threads in the first several warps satisfy the condition and execute identically. Divergence only becomes a factor in the very last few steps when s < 32.
2. Bank Conflicts: The memory access pattern within the loop is now fully sequential and conflict-free.

This is a much more robust and performant algorithm. It respects the underlying hardware architecture, which is the essence of high-performance computing.

Conclusion

In this lesson, we took a critical step forward in optimizing our parallel reduction. We moved from diagnosing problems to implementing effective solutions, a core loop in performance engineering.

Key Takeaways:

• Fixing Divergence: We replaced the divergent if (tid % (2*s) == 0) condition with a non-divergent indexing scheme (index = 2*s*tid), which improved warp efficiency.
• Shared Memory Bank Conflicts: We learned that shared memory is divided into banks and that simultaneous access to the same bank by different threads in a warp leads to serialization. Our divergence fix, unfortunately, introduced a strided access pattern that caused severe bank conflicts.
• Sequential Addressing: The solution to bank conflicts is sequential addressing. By using a reversed loop and the condition if (tid < s), we ensure that active threads are contiguous and access consecutive memory locations, making the operation conflict-free.
• The Code: The core of our optimized reduction loop is now for(s=N/2; s>0; s>>=1) if(tid < s) sdata[tid] += sdata[tid+s];.

We've made tremendous progress. Our kernel is now significantly faster. However, if you look closely at the sequential addressing loop, you might notice a new, albeit smaller, inefficiency: if (tid < s). In the very first iteration, when s = blockDim.x / 2, half of the threads in the block are sitting idle. In our next lesson, we will begin tackling these more subtle inefficiencies, starting with a powerful technique to put idle threads to work and further optimize the data flow from global memory using warp-level primitives.