Intuition The one core ideaShared memory is split into 32 parallel lanes called banks , and all 32 threads of a warp can read at once only if each thread lands in a different bank . Everything in this topic is one question repeated: "given how my threads index the array, do any two of them land in the same bank?"
Before you can answer that question you need to own a small pile of words and symbols. The parent note throws around warp , bank , word , stride , mod, gcd, and [r][c] addressing as if you already know them. This page builds every one of them from nothing, in an order where each piece rests on the previous.
A thread is one worker running your program. On a GPU you launch thousands of them, and they all run the same line of code but on different data — thread number 5 handles element 5, thread number 6 handles element 6, and so on.
The number that tells a thread "who am I" is written threadIdx.x. If you have 32 threads, threadIdx.x runs 0 , 1 , 2 , … , 31 i
A warp is a fixed bundle of 32 threads that the hardware runs in lockstep — they all execute the same instruction at the same instant. When those 32 threads hit a shared-memory read, all 32 memory requests are fired simultaneously . That simultaneity is the whole reason banks matter.
Intuition Why 32 keeps appearingThere are 32 threads in a warp and 32 banks in shared memory. This is not a coincidence — the hardware is built so that a perfectly-behaved warp maps its 32 threads onto its 32 banks one-to-one. Every "conflict" is a failure of that one-to-one match.
Definition Byte and byte-addressA byte is 8 bits — the smallest chunk of memory that has its own numbered slot. That number is the byte-address : byte 0, byte 1, byte 2, … Think of memory as an infinitely long ruler where every millimetre is numbered.
Definition Word (a 4-byte word)A float (and an int) takes 4 bytes . Four consecutive bytes glued together and treated as one value is called a 32-bit word (32 bits = 4 bytes). So byte-addresses 0 , 1 , 2 , 3 inside word 0 4 , 5 , 6 , 7 1
To turn a byte-address into a word-number you divide by 4 and throw away the remainder:
word_index = ⌊ 4 byte_address ⌋
Those ⌊ ⌋ floor symbol — "round down to a whole number." We need it because addresses 0, 1, 2, 3 must all collapse to word 0.
Intuition Why bother with words at all?The banks are 4 bytes wide, so the bank hardware doesn't care about individual bytes — it only cares which 4-byte word you asked for. Converting to word-index first is the step that lets the next formula ignore the messy byte details.
Definition Modulo, writtenmod or mod a mod n remainder left over after you divide a n 7 mod 3 = 1 7 = 2 × 3 + 1 0 , 1 , … , n − 1
The picture is a clock face . A clock does mod 12 mod 32
mod 32 and not something else?
With only 32 physical banks, word number 32 has nowhere new to go — the round-robin must wrap. mod 32 is exactly the arithmetic of "keep dealing cards into 32 piles, one card per pile, starting over at pile 0 each lap." Any tool that captures wrap-around would do; mod is the standard name for it.
s Stride is the gap, measured in words , between the element thread i i + 1 shared[i] the gap is 1 word (stride 1). In shared[i*2] the gap is 2 words (stride 2). In matrix[i][col] with 32 columns, moving one row jumps 32 words (stride 32).
So thread i i ⋅ s
bank ( i ) = ( i ⋅ s ) mod 32.
This single line is the engine of the whole topic. Everything the parent note computes is this line with a specific s
Worked example Feel the stride
Stride 1: banks are 0 , 1 , 2 , … , 31 one-to-one , no conflict.
Stride 2: banks are 0 , 2 , 4 , … , 30 , 0 , 2 , … and thread 16. Collision.
Stride 32: every thread's bank is ( i ⋅ 32 ) mod 32 = 0
Definition Greatest common divisor,g cd( a , b ) g cd( a , b ) biggest whole number that divides both a b g cd( 2 , 32 ) = 2 g cd( 3 , 32 ) = 1 g cd( 16 , 32 ) = 16 g cd( s , 32 ) = 1 s 32 coprime (they share no factor but 1).
right tool
We want the smallest step k k ( k ⋅ s ) mod 32 = 0 k ⋅ s shared factors between s 32
The payoff formula the parent uses:
conflict_degree = g cd( s , 32 )
Coprime stride (g cd= 1 i ↦ ( i ⋅ s ) mod 32 all 32 banks — a permutation, perfectly conflict-free. (This is why the odd-stride .x field access is safe.)
g cd= d > 1 32/ d d d
Definition Row-major layoutA 2-D array matrix[R][C] is really a 1-D strip in memory. Row-major means row 0 is laid out first, then row 1, etc. Element [r][c] sits at word:
word = r ⋅ C + c C c → c + 1 r → r + 1 C
That is why reading a column (matrix[i][col], varying the row i C C = 32 C = 33 33 mod 32 = 1 in bank-space ).
Stride s and bank of thread i
gcd predicts conflict degree
Row-major r times C plus c
Bank conflicts in shared memory
Banks live inside the wider 6.2.07-Shared-memory-architecture picture.
The same "different lanes at once" idea in global memory is 6.2.08-Memory-coalescing .
The column-walk conflict is the star villain of 6.3.01-Matrix-transpose-optimization .
How much shared memory you can afford ties into 6.2.10-Occupancy-and-resource-limits .
Threads acting in lockstep also underlies 6.2.11-Warp-divergence .
Return to the parent: the Hinglish version of the topic .
Cover the right side and answer out loud before revealing.
A thread is one worker running the same code on its own data, identified by threadIdx.x.
A warp is a bundle of 32 threads that execute in lockstep and fire their memory requests simultaneously.
A byte-address vs a word a byte is one numbered slot; a 4-byte word is four bytes glued into one float/int.
Why divide a byte-address by 4 to convert it to a word-index, since banks are 4 bytes wide and only care about words.
a mod n the remainder after dividing a n 0 … n − 1
The master bank formula bank_index = word_index mod 32
What stride s the gap in words between the elements read by thread i i + 1 i ( i ⋅ s ) mod 32
g cd( a , b ) the largest whole number dividing both a b g cd( s , 32 )
Why coprime strides are safe if g cd( s , 32 ) = 1 i ↦ ( i ⋅ s ) mod 32
Word address of matrix[r][c] (row-major, C r ⋅ C + c C
Why padding C 33 mod 32 = 1
