Avoiding Bank Conflicts
Why padding can speed up shared access.
Avoiding Bank Conflicts is a free CUDA Academy lesson on CoddyKit — lesson 3 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
Shared Memory Has Banks
Shared memory is split into 32 banks, one for each thread in a warp. Spreading accesses across them lets all 32 threads read at full speed.
How Addresses Map
Consecutive 4-byte words land in consecutive banks, wrapping around after 32. So word 0 is bank 0, word 32 is bank 0 again.
What a Bank Conflict Is
A bank conflict happens when two threads in a warp hit the same bank but different words. The hardware must serialize those accesses.
Serialization Costs Speed
An N-way conflict turns one fast access into N slow ones. A 32-way conflict can make shared memory as slow as if it were single-file. 🐢
The Happy Case: Stride One
When each thread reads tile[threadIdx.x], every thread hits a distinct bank. That is conflict-free and runs at full bandwidth.
float v = tile[threadIdx.x];The Trap: Stride of 32
Reading with a stride of 32 sends every thread to the same bank. That is the worst case, a full 32-way conflict.
float v = tile[threadIdx.x * 32];Even Strides Hurt Too
Any even stride that shares a factor with 32 causes partial conflicts. A stride of 2, for example, gives a 2-way conflict across the warp.
The Broadcast Exception
Good news: if all threads read the same address, the hardware broadcasts it in one shot. Same word is fine, only same bank with different words conflicts.
Padding to the Rescue
For 2D tiles, add one extra column with +1 padding. This shifts each row so column accesses no longer all land in one bank.
__shared__ float tile[32][33];Why +1 Works
The extra column makes the row length coprime with 32. Now stepping down a column visits a different bank each time, killing the conflict.
Profile, Do Not Guess
Bank conflicts are invisible in source code. Use Nsight Compute to measure shared-memory conflicts before spending effort fixing them.
Quick Check
Let us check what makes shared access fast.
Recap
You learned that shared memory has 32 banks, that same-bank different-word access serializes, and that +1 padding fixes column conflicts. Next: dynamic shared memory. 🎯
Frequently asked questions
Is the “Avoiding Bank Conflicts” lesson free?
Yes — the full text of “Avoiding Bank Conflicts” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “Avoiding Bank Conflicts”?
Why padding can speed up shared access. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 3 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Avoiding Bank Conflicts” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- Declaring __shared__ Arrays
- Synchronizing with __syncthreads
- Avoiding Bank Conflicts
- Dynamic Shared Memory