1. What Is a Bank?
CUDA shared memory is organized into memory banks.
Think of shared memory like:
1Shared Memory 2│ 3├── Bank 0 4├── Bank 1 5├── Bank 2 6├── Bank 3 7├── ... 8└── Bank N
The exact bank organization and behavior can vary by GPU architecture, but for current NVIDIA GPUs, a useful learning model is 32 banks.
Because a warp has:
132 threads
this creates an important relationship:
132 Threads 2 ↓ 332 Memory Banks
2. Basic Example
Suppose a warp accesses:
1Thread 0 → Bank 0 2Thread 1 → Bank 1 3Thread 2 → Bank 2 4Thread 3 → Bank 3 5... 6Thread 31 → Bank 31
This is a good pattern.
Conceptually:
1T0 ─────→ B0 2T1 ─────→ B1 3T2 ─────→ B2 4T3 ─────→ B3 5... 6T31 ─────→ B31
Different threads access different banks.
3. What Is a Bank Conflict?
A bank conflict occurs when multiple threads in a warp access different addresses that map to the same shared-memory bank.
For example:
1T0 ──┐ 2T1 ──┤ 3T2 ──┼──→ Bank 0 4T3 ──┘
Instead of being serviced efficiently in parallel, the accesses may need to be handled in multiple steps.
Conceptually:
1No conflict: 2 3T0 → B0 4T1 → B1 5T2 → B2 6T3 → B3 7 8 ↓ 9 10Parallel access 11 12 13Conflict: 14 15T0 → B0 16T1 → B0 17T2 → B0 18T3 → B0 19 20 ↓ 21 22Serialized access 23 ↓ 24Potential slowdown
4. Important Exception: Broadcast
There is an important special case.
If multiple threads access the same shared-memory location, NVIDIA GPUs can handle certain broadcast accesses efficiently.
For example:
1T0 ──┐ 2T1 ──┤ 3T2 ──┼──→ same address 4T3 ──┘
Don't simplify the rule to:
"Multiple threads accessing the same bank always means a bank conflict."
The important case is:
Different addresses mapping to the same bank.
5. How Banks Are Determined
For a simple learning model, suppose:
132 banks 24-byte word
Then for a 32-bit value:
1bank ≈ word_index % 32
For example:
1word 0 → bank 0 2word 1 → bank 1 3word 2 → bank 2 4... 5word 31 → bank 31 6word 32 → bank 0 7word 33 → bank 1
So:
10 → 0 21 → 1 32 → 2 4... 531 → 31 632 → 0
This modulo relationship is extremely useful for understanding bank conflicts.
6. Example: No Conflict
1__shared__ float data[32]; 2 3int tid = threadIdx.x; 4 5float x = data[tid];
For the first warp:
1T0 → data[0] 2T1 → data[1] 3T2 → data[2] 4... 5T31 → data[31]
Using:
1bank = index % 32
we get:
1T0 → Bank 0 2T1 → Bank 1 3T2 → Bank 2 4... 5T31 → Bank 31
Excellent.
7. Example: 2-Way Conflict
Consider:
1float x = data[tid * 2];
The first warp accesses:
1T0 → data[0] 2T1 → data[2] 3T2 → data[4] 4T3 → data[6] 5...
Banks:
1T0 → Bank 0 2T1 → Bank 2 3T2 → Bank 4 4... 5T16 → Bank 0 6T17 → Bank 2 7...
Notice:
1T0 → Bank 0 2T16 → Bank 0
Two different addresses map to the same bank.
This creates a 2-way bank conflict pattern.
8. 4-Way Conflict
Now:
1float x = data[tid * 4];
Then:
1T0 → Bank 0 2T1 → Bank 4 3T2 → Bank 8 4T3 → Bank 12 5...
After enough threads:
1T0 → Bank 0 2T8 → Bank 0 3T16 → Bank 0 4T24 → Bank 0
That's a 4-way conflict pattern.
9. 32-Way Conflict
Consider:
1float x = data[tid * 32];
Then:
1T0 → data[0] 2T1 → data[32] 3T2 → data[64] 4... 5T31 → data[992]
Since:
10 % 32 = 0 232 % 32 = 0 364 % 32 = 0 4...
all threads map to the same bank:
1T0 ──┐ 2T1 ──┤ 3T2 ──┤ 4... ├──→ Bank 0 5T31 ──┘
This is a severe conflict pattern.
10. Working CUDA Example
Let's create a simple kernel demonstrating shared-memory access.
1#include <stdio.h> 2#include <cuda_runtime.h> 3 4__global__ void bank_demo() 5{ 6 __shared__ float data[1024]; 7 8 int tid = threadIdx.x; 9 10 data[tid] = tid; 11 12 __syncthreads(); 13 14 float value = data[tid]; 15 16 printf( 17 "Thread %d: %.1f\n", 18 tid, 19 value 20 ); 21} 22 23int main() 24{ 25 bank_demo<<<1, 32>>>(); 26 27 cudaError_t err = 28 cudaGetLastError(); 29 30 if (err != cudaSuccess) 31 { 32 printf( 33 "Launch error: %s\n", 34 cudaGetErrorString(err) 35 ); 36 37 return 1; 38 } 39 40 err = cudaDeviceSynchronize(); 41 42 if (err != cudaSuccess) 43 { 44 printf( 45 "Execution error: %s\n", 46 cudaGetErrorString(err) 47 ); 48 49 return 1; 50 } 51 52 return 0; 53}
Compile:
1nvcc bank_demo.cu -o bank_demo
Run:
1./bank_demo
11. Important Point About This Example
This kernel:
1data[tid]
has a clean access pattern.
But:
1data[tid * 32]
would create a very different bank mapping.
You can experiment with:
1float value = data[tid];
then:
1float value = data[tid * 2];
then:
1float value = data[tid * 4];
then:
1float value = data[tid * 32];
This is a good practical experiment for understanding bank mapping.
12. The Classic Matrix Transpose Problem
Bank conflicts become especially important with matrix operations.
Consider:
1__shared__ float tile[32][32];
Suppose threads read:
1tile[threadIdx.x][threadIdx.y]
instead of:
1tile[threadIdx.y][threadIdx.x]
The access direction changes.
This can create a bank-conflict pattern depending on the warp's mapping.
That's why optimized matrix-transpose kernels often use:
132 × 32 tile
with padding.
13. Padding
One of the classic techniques for avoiding certain shared-memory bank conflicts is padding.
Instead of:
1__shared__ float tile[32][32];
use:
1__shared__ float tile[32][33];
Notice:
132
became:
133
Why?
Because the extra column changes the mapping of subsequent rows to banks.
14. How Padding Changes the Layout
Without padding:
132 × 32 2 3Row 0 → 32 elements 4Row 1 → 32 elements 5Row 2 → 32 elements 6...
Because 32 aligns with the number of banks, corresponding positions in different rows can map to the same bank.
With:
132 × 33
the row stride becomes:
133 elements
Conceptually:
1Row 0 → starts at 0 2Row 1 → starts at 33 3Row 2 → starts at 66 4Row 3 → starts at 99
Modulo 32:
10 21 32 43 5...
The bank mapping shifts between rows.
This can eliminate certain conflict patterns.
15. Padded Transpose
A common pattern is:
1#define TILE 32 2 3__global__ void transpose_padded( 4 const float* input, 5 float* output, 6 int N 7) 8{ 9 __shared__ 10 float tile[TILE][TILE + 1]; 11 12 int x = 13 blockIdx.x * TILE 14 + threadIdx.x; 15 16 int y = 17 blockIdx.y * TILE 18 + threadIdx.y; 19 20 if (x < N && y < N) 21 { 22 tile[threadIdx.y][threadIdx.x] = 23 input[y * N + x]; 24 } 25 26 __syncthreads(); 27 28 int tx = 29 blockIdx.y * TILE 30 + threadIdx.x; 31 32 int ty = 33 blockIdx.x * TILE 34 + threadIdx.y; 35 36 if (tx < N && ty < N) 37 { 38 output[ty * N + tx] = 39 tile[threadIdx.x][threadIdx.y]; 40 } 41}
The important line is:
1float tile[TILE][TILE + 1];
instead of:
1float tile[TILE][TILE];
16. Why +1?
Because:
132 columns
becomes:
133 columns
So:
1row stride = 33
rather than:
1row stride = 32
Since:
133 % 32 = 1
the starting bank shifts by one for each row.
Conceptually:
1Without padding: 2 3Row 0 → Bank 0 4Row 1 → Bank 0 5Row 2 → Bank 0 6Row 3 → Bank 0 7 8 9With padding: 10 11Row 0 → Bank 0 12Row 1 → Bank 1 13Row 2 → Bank 2 14Row 3 → Bank 3
This is the key idea.
17. Very Important: Padding Is Not Magic
Don't memorize:
132 → 33
as a universal rule.
The correct principle is:
Change the shared-memory layout so that threads that would otherwise access different addresses in the same bank are distributed across banks.
The padding amount depends on:
1data type 2tile shape 3access pattern 4bank organization 5GPU architecture
18. Bank Conflicts vs Global Memory Coalescing
These are different problems.
Global memory
You care about:
1Warp 2 ↓ 3Global memory addresses 4 ↓ 5Memory transactions 6 ↓ 7Coalescing
Shared memory
You care about:
1Warp 2 ↓ 3Shared-memory addresses 4 ↓ 5Memory banks 6 ↓ 7Bank conflicts
So remember:
1Global Memory 2→ Coalescing 3 4Shared Memory 5→ Bank Conflicts
19. Complete Mental Model
You now have:
1 WARP 2 │ 3 ┌───────┴────────┐ 4 ↓ ↓ 5 Global Memory Shared Memory 6 │ │ 7 ↓ ↓ 8 Coalescing Bank Mapping 9 │ │ 10 ↓ ↓ 11Memory Transactions Bank Conflicts
This distinction is fundamental.
20. How to Analyze a Bank Conflict
When you see:
1shared[index]
do this:
Step 1
Determine the index for each thread.
1T0 → index ? 2T1 → index ? 3T2 → index ? 4...
Step 2
Convert index to bank.
For a simple 32-bit-word model:
1bank = index % 32
Step 3
Look for duplicate banks.
Example:
1T0 → Bank 0 2T1 → Bank 1 3T2 → Bank 2 4T3 → Bank 3
Good.
But:
1T0 → Bank 0 2T1 → Bank 0 3T2 → Bank 0 4T3 → Bank 0
Potential conflict.
21. Practice Problem
Given:
1__shared__ float data[1024]; 2 3int tid = threadIdx.x; 4 5float x = 6 data[tid * 2];
For the first warp, calculate:
1Thread 0 → ? 2Thread 1 → ? 3Thread 2 → ? 4... 5Thread 31 → ?
Then calculate:
1Bank = 2index % 32
You should discover the repeating bank pattern.
Then try:
1data[tid * 4]
and:
1data[tid * 8]
and:
1data[tid * 32]
This exercise will make the concept much clearer than memorizing definitions.
22. Bank Conflict Formula
For the simplified 32-bank model:
1bank = 2(word_index) % 32
If the element size is 4 bytes:
1word_index = 2byte_address / 4
Therefore:
1bank = 2(byte_address / 4) % 32
This is a useful formula to remember.
23. What Happens During a Conflict?
Conceptually:
1No conflict 2 3T0 → B0 ─┐ 4T1 → B1 ─┤ 5T2 → B2 ─┤→ parallel access 6T3 → B3 ─┘
Conflict:
1T0 → B0 2T1 → B0 3T2 → B0 4T3 → B0
The hardware may need multiple servicing steps.
Therefore:
1Bank conflict 2 ↓ 3More serialized work 4 ↓ 5Lower shared-memory throughput 6 ↓ 7Potential kernel slowdown
The exact performance impact depends on the pattern and architecture.
24. Why AI Kernel Developers Care
This becomes very important in:
1GEMM 2Attention 3FlashAttention-style kernels 4Transpose 5Reduction 6Softmax 7LayerNorm 8Tensor tiling 9Convolution
For example:
1Global Memory 2 ↓ 3Coalesced load 4 ↓ 5Shared Memory 6 ↓ 7Bank-efficient access 8 ↓ 9Registers 10 ↓ 11Tensor Core / CUDA Core
A high-performance kernel needs all of these pieces working together.
25. Your Optimization Checklist
When writing a shared-memory kernel:
1□ What is my warp size? 2□ What does each thread access? 3□ What bank does each access map to? 4□ Are different addresses hitting the same bank? 5□ Is there a broadcast? 6□ Can padding fix the pattern? 7□ Is shared memory actually giving data reuse? 8□ Is register pressure becoming too high? 9□ Did performance improve after the change?
That last question is critical:
Always benchmark.
An optimization that looks theoretically better may not improve the complete kernel.
26. Phase 3 Progress
Your current roadmap becomes:
1Phase 3 — Advanced Memory Optimization 2 3├── Shared Memory Bank Conflicts ✅ 4├── Memory Transactions ← NEXT 5├── Cache Behavior 6├── Memory Bandwidth 7├── Pinned Memory 8├── Unified Memory 9└── Asynchronous Memory Transfers
The next topic should be Memory Transactions, because now you understand both sides:
1Global Memory 2 ↓ 3Memory Transactions + Coalescing 4 5Shared Memory 6 ↓ 7Banks + Bank Conflicts
After Memory Transactions, we'll connect that to cache behavior and memory bandwidth, which is where you'll start learning how to determine whether a real CUDA kernel is memory-bound or compute-bound.