Bank Conflict

A bank conflict occurs when two or more threads within the same GPU warp attempt to access different data words stored in the same memory bank of shared memory during a single instruction cycle. This forces the hardware to serialize the accesses, turning a parallel operation into a sequential one. The root cause is the interleaved bank addressing scheme used to maximize bandwidth, where consecutive 32-bit words are assigned to consecutive banks in a round-robin fashion (e.g., 32 banks for 32-bit words).

The primary impact is a drastic reduction in effective memory bandwidth and increased latency. A single conflict can cause a 2x slowdown; a worst-case scenario where all 32 threads in a warp access the same bank serializes all accesses, resulting in a 32x slowdown. This stalls the warp scheduler, reducing instruction-level parallelism (ILP) and overall GPU occupancy. The performance penalty is multiplicative, making it a critical optimization target for kernels heavy on shared memory operations like matrix transposes or reductions.

Conflicts arise from specific, predictable stride patterns in thread memory access:

• Strided Access with a stride that is a multiple of the bank count: The most common cause. For 32 banks, a stride of 32, 64, 96, etc., will cause all threads to map to the same bank.
• Random or Irregular Access: If thread IDs map unpredictably to addresses, collisions are likely.
• Diagonal Access in 2D Arrays: Accessing elements like array[threadIdx.x][threadIdx.x] often leads to bank conflicts.
• Broadcast Access (a special case): When all threads read the same address, it is optimized by hardware into a broadcast and does not cause a conflict.

The GPU hardware resolves conflicts by splitting the simultaneous access request into multiple, sequential transaction cycles. The number of cycles equals the number of threads that collided on a single bank (the degree of conflict). Modern GPUs (Compute Capability 2.0+) have 32-bit and 64-bit bank modes. In 64-bit mode, successive 64-bit words are assigned to successive banks, effectively halving the bank count to 16 for 64-bit accesses, which changes the conflict dynamics for double precision data.

Developers use several strategies to avoid or eliminate bank conflicts:

• Memory Padding: Adding an extra column (a "pad") to shared memory arrays changes the base address stride, shifting threads into different banks. For a 2D array tile[32][32], declaring it as tile[32][33] can eliminate stride-32 conflicts.
• Data Layout Transformation: Reorganizing data in shared memory (e.g., using a different indexing scheme) to ensure coalesced or conflict-free access patterns.
• Using Alternate Algorithms: Choosing parallel algorithms designed for conflict-free shared memory access, such as certain parallel reduction or scan implementations that use sequential addressing.
• Bank Size Awareness: Writing code that is aware of the bank width (32-bit vs. 64-bit) for the target architecture.

It is crucial to distinguish bank conflicts from memory coalescing. They are complementary optimizations for different memory hierarchies:

• Bank Conflicts: Occur in shared memory (on-chip, software-managed cache). Optimization focuses on access stride relative to the 32-bank structure.
• Coalescing: Concerns global memory (off-chip DRAM). Optimization focuses on having consecutive threads access consecutive 128-byte aligned segments to minimize transactions. A kernel can be perfectly coalesced for global memory access yet suffer severe bank conflicts in shared memory, and vice-versa. Both must be optimized for peak performance.

An optimal global memory access pattern where consecutive threads in a warp access consecutive memory addresses. This allows the GPU's memory controller to combine these requests into a single, wide transaction, maximizing effective bandwidth. It is the global memory equivalent of avoiding bank conflicts in shared memory. Key patterns include:

• Thread i accesses element base_address + i.
• Accesses are aligned to 32-byte or 128-byte boundaries.
• Non-coalesced access scatters requests across memory, causing severe performance penalties.

The order and structure in which threads within a GPU warp read from or write to memory. The pattern dictates performance. For shared memory, the critical factor is the mapping of thread indices to memory banks. For global memory, the stride and alignment determine coalescing. Engineers analyze and redesign algorithms (e.g., using padding, matrix transposition, or Z-order curves) to transform suboptimal patterns like strided or random access into efficient, contiguous, and bank-conflict-free patterns.

A fundamental unit of execution and scheduling on an NVIDIA GPU, consisting of 32 threads that execute the same instruction in lockstep (Single Instruction, Multiple Threads - SIMT). Bank conflicts are defined at the warp level because all 32 threads execute simultaneously. Memory requests from threads within a warp are issued together, making their collective access pattern the determinant for conflicts. Understanding warp execution is essential for predicting and avoiding serialization in shared memory.

A hardware subdivision within shared memory. On modern NVIDIA GPUs, shared memory is typically partitioned into 32 equally sized banks (matching the warp size). Successive 4-byte words (32-bit) are assigned to successive banks in a round-robin fashion. When two or more threads in a warp request data from different addresses within the same bank, a conflict occurs, and the accesses are processed sequentially. Techniques like address padding or memory layout transformation are used to avoid bank conflicts.

The organization of memory subsystems in a GPU into multiple levels with different capacities, latencies, and bandwidths. Bank conflicts specifically affect the shared memory/L1 cache level. The hierarchy typically includes:

• Registers: Fastest, per-thread storage.
• Shared Memory / L1 Cache: Low-latency, block-level (banked).
• L2 Cache: Larger, shared across all SMs.
• Global Memory (HBM/GDDR): High-bandwidth, high-latency DRAM. Optimization involves keeping data as high in this hierarchy as possible and ensuring access patterns are efficient for each level.