100 DAYS OF GPU Course preview
← Glossary
CUDA glossaryMemory
CC 7.5

What is a shared memory bank on a GPU?

One of 32 independent slices of shared memory, each serving 4 bytes per cycle, addressed by (byteAddress / 4) % 32.

Ask a warp for 32 floats out of shared memory and the hardware does not go to one place for them. The memory is cut into 32 modules that answer at the same time, and successive 32-bit words go to successive modules: word 0 to bank 0, word 31 to bank 31, word 32 back round to bank 0. NVIDIA's best-practices guide states both halves of that, first that shared memory "is divided into equally sized memory modules (banks) that can be accessed simultaneously", then that "each bank has a bandwidth of 32 bits every clock cycle, and successive 32-bit words are assigned to successive banks". So the bank holding word w is w mod 32, and the shape you declared your array with never enters the calculation.

Neither constant is arbitrary. There are as many banks as there are lanes in a warp, and a bank returns exactly the 4 bytes a lane normally wants, so the ceiling on a shared read is one word per lane per cycle and the hardware reaches it when the 32 lanes sit on 32 different banks. Table 31 of the compute-capability appendix carries a "Number of shared memory banks" row, and it reads 32 for every architecture this course touches, so the arithmetic you learn on a free T4 still holds on an H100.

The move that makes this usable is to stop thinking in rows and columns. __shared__ float tile[32][32] is 1024 words laid out in a line, so tile[r][c] is word 32r + c, whose bank is c mod 32. Every element of column c is in bank c, all 32 of them. That is why walking a row is the cheapest thing you can do to shared memory and walking a column of a 32-wide tile is the most expensive. Widen the declaration to tile[32][33] and tile[r][c] becomes word 33r + c in bank (r + c) mod 32, so the same column walk now visits all 32 banks, because each row starts one bank further along than the last.

One case looks like a collision and is not. When several lanes read the same address, the guide says the accesses "are coalesced into a single multicast", so a bank only ever serializes when the words asked of it are different. Banks turn access patterns into a question about remainders, never about distance, which is what the whole of bank conflicts rests on.

Measured

On a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75, day 15 read a 2048-word shared tile 2048 times per thread and varied one runtime argument: the word stride between neighbouring lanes.

Lane L reads word Banks the 32 lanes cover Time vs the best row
L 32 0.286 ms 1.00
32L 1 8.886 ms 31.10
33L 32 0.286 ms 1.00

Same kernel, same instruction count, same 2048 reads per thread, and a factor of 31 between the first row and the second. The third row is the first row's time to three decimal places, which is the bank map showing through the clock: 33 words on from a bank is one bank on from it, so the lanes fan out again. The 4-byte width is what makes the middle row cost 32 requests. One bank can put one word on the wire per cycle, and 32 lanes wanted 32 different words out of it.

Diagram

bank-conflicts, preset bank-map: a 32 by 32 grid of tile words shaded by w mod 32, with one row and one column picked out.

Alt text: "Words of a 32-wide shared tile shaded by bank. A row read touches all 32 banks once each; a column read touches one bank 32 times."

Related terms

Where you meet this

Sources

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.