Skip to main content
Back to path

Welcome to a Grasp lesson

Create a free account to keep studying and create your own study path.

Lesson 33

Understanding and Avoiding Shared Memory Bank Conflicts

Recommended resources

Introduction

In our last lesson, we achieved a significant performance boost by implementing a tiled matrix transpose. We used shared memory as a high-speed scratchpad to transform inefficient, uncoalesced global memory writes into efficient, coalesced ones. However, we ended with a question: why is the tiled transpose still not as fast as a simple memory copy? The answer lies in a subtle but critical aspect of the GPU architecture that we will explore today: shared memory bank conflicts.

This lesson will deconstruct the internal architecture of shared memory. You will learn what banks are, how memory addresses are mapped to them, and how certain access patterns can lead to conflicts that serialize execution and degrade performance. Most importantly, we will revisit our matrix transpose kernel, diagnose the exact cause of its bank conflicts, and implement a simple yet powerful technique—memory padding—to eliminate them. This is a foundational concept in high-level CUDA optimization, moving beyond correctness to true performance engineering.

The Architecture of Shared Memory: Parallelism via Banks

To understand bank conflicts, we first need to refine our mental model of shared memory. It is not a single, monolithic block of RAM. Instead, to achieve its extremely high bandwidth, shared memory is divided into multiple independent modules called banks. On all modern NVIDIA GPUs, shared memory is organized into 32 banks.

This number is not arbitrary; it's designed to perfectly match the number of threads in a warp (32). The ideal scenario is that each of the 32 threads in a warp accesses a different bank simultaneously. When this happens, all 32 memory requests can be serviced in parallel in a single clock cycle, achieving maximum throughput.

So, how does the GPU decide which bank holds which memory address? The memory is interleaved across the banks word by word. A "word" in this context is a 4-byte (32-bit) chunk of data, the size of a standard float or int.

  • The first word (bytes 0-3) of your shared memory array is in Bank 0.
  • The second word (bytes 4-7) is in Bank 1.
  • ...
  • The 32nd word (bytes 124-127) is in Bank 31.
  • The 33rd word (bytes 128-131) wraps around and is back in Bank 0.

This organization is visualized in the diagram below.

Diagram of shared memory organization. Successive 4-byte words are mapped to successive banks in a round-robin fashion. Word 0, Word 32, and Word 64 all reside in Bank 0.

The following video provides an excellent introduction to this banked architecture and the concept of warps.

NVIDIA CUDA Tutorial 9: Bank Conflicts

This tutorial by Creel on YouTube provides a clear, practical explanation of bank conflicts.

Watch the first two sections of the video. The first part, Warps and Shared Memory, will briefly review how threads are grouped into warps and introduce the idea of bank conflicts. The second part, Shared Memory Organization, explains the physical layout of shared memory into 32 banks and how words are mapped to them. This is the core mechanism we need to understand.

The mapping rule can be expressed with a simple formula. For a given byte address A within the shared memory block, the bank it belongs to is calculated as:

where A / 4 gives the word index. This formalizes the interleaved, or "strided," layout of data across the banks.

Defining Bank Conflicts

Now that we understand the banked architecture, we can define a bank conflict. A bank conflict occurs when two or more threads within the same warp attempt to access different memory addresses that reside in the same bank.

Since a bank can only service one request per cycle, a conflict forces the memory accesses to be serialized. For example, in a 2-way bank conflict, where two threads hit the same bank, the second request must wait for the first to complete. This effectively halves the shared memory bandwidth for that instruction. In a 32-way conflict, all 32 threads try to access different addresses in the same bank, forcing the warp to execute completely serially for that access—a 32x slowdown!

It's important to note two key exceptions that do not cause conflicts:

  1. Broadcast: If all threads in a warp read from the exact same address (e.g., a configuration parameter), the hardware is smart enough to perform a single read and broadcast the value to all threads. This is very fast.
  2. Multicast: If multiple threads (but not necessarily all) read from the exact same address, the value is read once and "multicast" to the requesting threads. This is also fast.

Therefore, a performance-degrading bank conflict occurs only when multiple threads access different addresses within the same bank in the same instruction.

NVIDIA CUDA Tutorial 9: Bank Conflicts

Let's continue with the Creel video to solidify this definition.

Watch the section from What is a Bank Conflict? where the presenter clearly distinguishes between a conflict, a broadcast, and a multicast.

Unmasking the Conflict in Our Tiled Transpose

Let's connect this theory directly to the transposeCoalesced kernel from our previous lesson.

The problem arises during the second phase of the algorithm, where threads read from the shared memory tile to write to global memory:

// Recall: TILE_DIM was 32
__shared__ float tile[TILE_DIM][TILE_DIM];

// ... load tile, then __syncthreads() ...

// This is where the conflict occurs!
odata[...] = tile[threadIdx.x][threadIdx.y + j];

Consider a warp of 32 threads. These threads all execute in lockstep and will share the same threadIdx.y and blockIdx values but have threadIdx.x values from 0 to 31.

When these threads execute tile[threadIdx.x][threadIdx.y + j], they are accessing a column of the tile array.

  • Thread 0 reads tile[0][y]. Let's say this is word W. This maps to Bank W % 32.
  • Thread 1 reads tile[1][y]. In a row-major layout, this address is 32 words away from tile[0][y]. Its word index is W + 32. This maps to Bank (W + 32) % 32.
  • Thread 2 reads tile[2][y]. Its word index is W + 64. This maps to Bank (W + 64) % 32.

Do you see the pattern? Since our stride between rows is 32, all these accesses map to the same bank!

(W + k * 32) % 32 is always equal to W % 32.

We have inadvertently created a 32-way bank conflict. Every thread in the warp tries to access a different address in the same bank, causing the reads from shared memory to be fully serialized. This is the primary reason our tiled transpose still lags in performance.

This "power of two" stride problem is one of the most common pitfalls in CUDA programming. A more formal analysis explains why odd strides are safe while even strides are risky.

A Comprehensive Analysis of CUDA Shared Memory, Bank ... - Uplatz

This detailed article from Uplatz provides a rigorous, architectural view of bank conflicts.

Please read the section titled Deconstruction of Access Patterns. Focus on the distinction between odd and even strides. The key insight is that any stride S that is coprime with 32 (i.e., odd) will simply permute the bank accesses, resulting in no conflicts. The subsection on "The 'Power of Two' Trap in 2D Tiling" directly describes the exact problem we've just diagnosed in our transpose kernel.

The Solution: Memory Padding

Now that we have identified the problem—a column-access stride of 32—the solution becomes clear: we must change the stride. The simplest way to do this is with padding.

Instead of declaring our shared memory tile with a width of 32, we add one dummy element to each row, making the width 33.

Original Code:
__shared__ float tile[TILE_DIM][TILE_DIM]; (i.e., [32][32])

Padded Code:
__shared__ float tile[TILE_DIM][TILE_DIM + 1]; (i.e., [32][33])

With this single change, the stride between consecutive rows is now 33 words. Since 33 is odd (and thus coprime with 32), when our warp accesses a column, each thread will access a unique bank.

  • Thread 0 accesses a word W in Bank W % 32.
  • Thread 1 accesses word W + 33 in Bank (W + 33) % 32 = (W + 1) % 32.
  • Thread 2 accesses word W + 66 in Bank (W + 66) % 32 = (W + 2) % 32.
  • ... and so on.

The bank conflict is completely eliminated. The only cost is a small amount of extra shared memory.

Let's watch a practical demonstration of how access patterns and padding affect performance.

NVIDIA CUDA Tutorial 9: Bank Conflicts

We'll return to the Creel video for a hands-on look at these patterns.

First, watch the visual examples from Request Patterns. This section beautifully illustrates how different strides (like *2, *3, *12) create conflicts, and how padding (*13 instead of *12) resolves them. Then, watch the coding demo from Timing the Patterns. Here, the presenter measures the raw clock cycles for different access patterns, including the use of padding on a struct to remove conflicts. This provides concrete evidence of the performance impact we've been discussing.

Diagnosing Conflicts with Nsight Compute

In a real-world scenario, you wouldn't just guess that you have bank conflicts. You would use a profiler to diagnose them. The primary tool for this is NVIDIA Nsight Compute. When you profile your kernel, it provides a detailed memory workload analysis.

An Nsight Compute memory chart view. The alert at the top, "Shared Load Bank Conflicts," is a clear indicator of a problem. Profiler metrics confirm that a large number of shared memory requests resulted in conflicts.

The profiler gives you specific metrics to look for:

  • l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum: This is a raw count of bank conflicts that occurred during shared memory loads. A non-zero value is your first clue.
  • L1 Wavefronts Shared Excessive: This is an even better metric. An "ideal" instruction should generate one "wavefront" (a request to the memory unit). If bank conflicts cause serialization, the hardware must issue excessive wavefronts. This metric directly quantifies the number of re-issued requests, giving you a measure of the performance penalty.

By inspecting these values, you can confirm whether bank conflicts are a bottleneck and whether your optimizations (like padding) have been successful.

Conclusion

Today, we've pulled back the curtain on shared memory to reveal its inner workings. You now understand that its immense bandwidth comes from a parallel, banked architecture that requires careful management of memory access patterns.

Key Takeaways:

  • Shared memory is physically divided into 32 banks, allowing for 32 parallel accesses per cycle.
  • A bank conflict occurs when multiple threads in a warp access different 4-byte words in the same bank, forcing the accesses to be serialized and reducing bandwidth.
  • Accessing data with a stride that is a multiple of 32 (like accessing columns in a 32-wide 2D array) is a classic cause of severe bank conflicts. This was the hidden bottleneck in our tiled matrix transpose.
  • The most common solution is padding the shared memory array (e.g., changing a width of 32 to 33) to make the access stride an odd number, which guarantees conflict-free access.
  • Profiling tools like Nsight Compute are essential for identifying and quantifying bank conflicts in real kernels.

You have now mastered the core principles of shared memory optimization. In our next lesson, we will broaden our focus to other specialized memory spaces, such as constant memory and the read-only data cache, which provide additional tools for optimizing data access in your CUDA kernels.