100 DAYS OF GPU Course preview
Day 15Module 2
in-technical-review

Shared memory bank conflicts explained

Here is the kernel a published CUDA course ships as its bank-conflict example, the one that is supposed to be the slow half of the comparison:

int conflictIndex = tx * 2 % TILE_WIDTH;  // TILE_WIDTH is 32
sharedData[ty][conflictIndex] = input[index];
__syncthreads();
output[index] = sharedData[ty][conflictIndex];

There is not one bank conflict in it. Threads tx and tx + 16 compute the same conflictIndex, so they write the same address, which is a race, and then read the same address, which the hardware broadcasts. Sixteen of the tile's 32 columns are never written.

The launch is one block, each thread does one shared write and one shared read, and the kernel is timed once. The result is launch overhead. (https://github.com/AdepojuJeremy/CUDA-120-DAYS--CHALLENGE/blob/main/daily-updates/day-12-Bank-Conflicts-in-Shared-Memory.md , checked 2026-08-30.)

By the end of this page you can look at any shared memory index and say how many ways it conflicts.

What a bank is, and what makes an access conflict

The most upvoted question on the subject, 131 votes, is not asking for the fix:

"I have been reading the programming guide for CUDA and OpenCL, and I cannot figure out what a bank conflict is. They just sort of dive into how to solve the problem without elaborating on the subject itself."

https://stackoverflow.com/questions/3841877/what-is-a-bank-conflict-doing-cuda-opencl-programming (checked 2026-08-30)

Start with the map. Shared memory is not one memory. NVIDIA: "To achieve high memory bandwidth for concurrent accesses, shared memory is divided into equally sized memory modules (banks) that can be accessed simultaneously." (https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html , section 10.2.3.1, checked 2026-08-30.)

The map from an address to a bank is one line of arithmetic, and the same page gives it: "On devices of compute capability 5.x or newer, each bank has a bandwidth of 32 bits every clock cycle, and successive 32-bit words are assigned to successive banks." So word 0 is in bank 0, word 31 in bank 31, word 32 back in bank 0.

The bank for word w is w mod 32. This model does not change between a T4 and a B200. Table 31 of the compute capabilities appendix has a "Number of shared memory banks" row, and it reads 32 for every capability this course targets (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html , checked 2026-08-29).

The next sentence is why it matters to you rather than to the hardware team: "The warp size is 32 threads and the number of banks is also 32, so bank conflicts can occur between any threads in the warp." One warp, 32 lanes, 32 banks, one request.

When two lanes want different words out of the same bank, the bank can only answer one of them per cycle: "If multiple addresses of a memory request map to the same memory bank, the accesses are serialized. The hardware splits a memory request that has bank conflicts into as many separate conflict-free requests as necessary, decreasing the effective bandwidth by a factor equal to the number of separate memory requests."

So the number to compute is the conflict degree: the largest number of distinct words any single bank is asked for. Degree 1 is one request. Degree 32 is 32 requests for what should have been one.

"Distinct" matters here, and the demo above got it wrong. Lanes reading the same address do not conflict: "The one exception here is when multiple threads in a warp address the same shared memory location, resulting in a broadcast. In this case, multiple broadcasts from different banks are coalesced into a single multicast from the requested shared memory locations to the threads."

Now put a stride on it. If lane L reads word L * s, it lands in bank L * s mod 32. Two lanes share a bank when (L - L') * s is a multiple of 32, which happens every 32 / gcd(s, 32) lanes.

Each occupied bank collects gcd(s, 32) lanes, and their words are all different. The degree is gcd(stride, 32).

Note. That widget is this page's only figure. A day with an interactive writes no separate diagram, so the widget carries its own static frame for a reader with no JavaScript, and column-32 is that frame.

Why stride 33 beats stride 32

The intuition you bring from day 11 is that a bigger stride is worse. It was true there and it was monotone: every step up the stride touched more sectors, then more 128-byte lines, then more DRAM pages, and the bandwidth fell the whole way.

Shared memory does not work like that. The model above has no line, sector, or cache. It has one map, w mod 32, and one question: do your 32 lanes cover 32 different remainders?

That makes the answer periodic instead of monotone. gcd(s, 32) is 32 at stride 32, and it is 1 at strides 31, 33, 63 and 65. Every odd stride is conflict free, however large.

Which is why + 1 is not a trick. A 32-wide tile puts row r at word 32r, so reading down a column is a stride of 32, the worst case there is. Make the row one word longer and row r starts at word 33r, the column read is a stride of 33, and gcd(33, 32) is 1.

Declaration Row r starts at Column read stride Degree
tile[32][32] word 32r 32 32
tile[32][33] word 33r 33 1

You are not aligning anything and you are not adding a gap for the hardware to skip. You are changing the row length from a multiple of the bank count to a number coprime with it. Padding is the cheap way to do that; day 73's tensor-core paths need the other way, an XOR swizzle, because they cannot afford the extra column.

Two measurements, because one would mislead

Full program in code/day15-bank-conflicts/bank_conflicts.cu. It runs two measurements, because either one on its own would teach the wrong thing.

The measurement follows four rules.

One kernel per measurement, one variable. The stride arrives as a runtime argument, so every row of the first table executes the same instructions in the same order and only the addresses move. A separate kernel per stride would let the compiler fold the constant and the rows would stop being comparable.

The accumulator is a receipt, not a result. Every word in the tile is 1.0f, so a thread that performed all 2,048 of its reads holds exactly 2,048, and the host checks all 81,920 outputs against that with no tolerance. A read the compiler deleted shows up as a wrong answer instead of as a fast row.

The model is checked at compile time, not asserted in prose. conflictDegree is constexpr, so the file will not build unless gcd(32, 32) is 32 and gcd(33, 32) is 1. A wrong value causes a build error.

Events, and a warm-up per kernel. Day 9 covers why a host clock around a launch measures the launch, and why the first launch of each kernel includes its own module load.

The probe is one kernel and one loop:

__global__ void probeSharedStride(const float* __restrict__ in,
                                  float* __restrict__ out, size_t n,
                                  int stride) {
    __shared__ float tile[kProbeWords];

    const unsigned int tid = threadIdx.x;
    for (unsigned int w = tid; w < kProbeWords; w += blockDim.x) {
        tile[w] = in[w];
    }
    __syncthreads();

    const unsigned int base = tid * static_cast<unsigned int>(stride);
    float acc = 0.0f;
    for (int k = 0; k < kProbeIters; ++k) {
        acc += tile[(base + static_cast<unsigned int>(k)) & kProbeMask];
    }

    const size_t t = blockIdx.x * static_cast<size_t>(blockDim.x) + tid;
    if (t < n) {
        out[t] = acc;
    }
}

Adding k moves every lane by the same amount, so the bank pattern is identical on every iteration and the mask cannot change it either, because 2048 is a multiple of 32.

The second measurement is the tiled transpose day 13 builds, which exists because the naive version of day 12 has to write a column. It is restated here over a matrix that is a whole number of tiles, so neither variant carries a bounds check and the two differ only in the width of one shared array.

It stores a tile row and loads a tile column with a barrier between the two, which is why day 14 came first. The store walks a row and is conflict free. The load walks a column, which is the 32-way case:

__global__ void transposeTiled(const float* __restrict__ in,
                               float* __restrict__ out, int n) {
    __shared__ float tile[kTileDim][kTileDim];

    const int tx = static_cast<int>(threadIdx.x);
    const int ty = static_cast<int>(threadIdx.y);
    const int col = static_cast<int>(blockIdx.x) * kTileDim + tx;
    const int row = static_cast<int>(blockIdx.y) * kTileDim + ty;

    for (int k = 0; k < kTileDim; k += kBlockRows) {
        tile[ty + k][tx] = in[static_cast<size_t>(row + k) * n + col];
    }
    __syncthreads();

    const int outCol = static_cast<int>(blockIdx.y) * kTileDim + tx;
    const int outRow = static_cast<int>(blockIdx.x) * kTileDim + ty;
    for (int k = 0; k < kTileDim; k += kBlockRows) {
        out[static_cast<size_t>(outRow + k) * n + outCol] = tile[tx][ty + k];
    }
}

The second variant is that kernel with one column added to the tile:

    __shared__ float tile[kTileDim][kTileDim + 1];

Both variants read and write global memory in coalesced 128-byte runs, both move the same 128 MiB, and the program prints that byte count so you can check it rather than trust it.

A straight copy over the same buffers runs first, on the same card in the same process, as the baseline. Day 13's tiled transpose beats the naive one but remains slower than the copy, and day 13 names the bank conflict as what is left. This page measures that claim, so both transposes are reported as a percentage of the copy and against each other.

Note. The padded kernel is the one that pays. It uses 128 more bytes of shared memory per block, and one extra integer add per shared access, because 33 is not a power of two and the compiler can fold 32 into a shift. If it still wins, it won against a handicap.

Results

Originally measured on a Tesla T4 with driver 595.84 and CUDA 12.6. The exact shipped program was re-verified with driver 580.173.02 and CUDA 13.0 (V13.0.88) on 2026-09-02; both transcripts remain in evidence. The conflict-degree sweep reproduced its ratios, but the transpose comparison drifted materially: padding changed runtime from 0.69x of unpadded under 12.6 to 0.96x under 13.0.

The original 12.6 capture follows.

GPU: Tesla T4 (compute capability 7.5)
Shared memory per block: 49152 bytes

Part 1: 2048 shared words per block, 2048 reads per thread
  stride     degree    time (ms)  vs stride 1
  ------     ------   ----------  -----------
       1          1        0.286         1.00
       2          2        0.563         1.97
       4          4        1.117         3.91
      32         32        8.886        31.10
      33          1        0.286         1.00

Part 2: 4096 x 4096 transpose, 128 MiB moved per launch
kernel                    time (ms)         GB/s    % of copy
----------------------   ----------   ----------   ----------
copy, no transpose            0.567        236.6        100.0
transpose tile[32][32]        0.994        135.1         57.1
transpose tile[32][33]        0.688        195.2         82.5

padded is 0.69x the plain kernel's time

All three rows above moved the same 128 MiB and all three read
and write global memory in 128-byte runs. The only difference
between the last two is the width of the shared tile.

The conflict degree is gcd(stride, 32) and the clock agrees to three significant figures. Stride 2 predicts a 2-way conflict and costs 1.97x. Stride 4 predicts 4-way and costs 3.91x.

Stride 32 predicts 32-way, every lane in the same bank, and costs 31.10x. Stride 33 is coprime with 32, predicts no conflict, and measures exactly the same as stride 1.

A 31x slowdown from shared memory, the fast memory, purely because of which lane reads which word.

In the original CUDA 12.6 run, one column of padding raised bandwidth by 45 percent. The unpadded tile[32][32] runs at 135.1 GB/s and tile[32][33] at 195.2, taking 0.69x the time. That single extra column costs 128 bytes per block and moves the kernel from 57 percent of copy bandwidth to 82.

The CUDA 13.0 re-run did not reproduce that magnitude: the unpadded and padded rows measured 0.706 and 0.680 ms, only a four percent time difference. The bank-index arithmetic still predicts the conflict and the synthetic sweep still measured the 32-way case at 31.24x, but this optimized transpose no longer supports a general 45 percent timing claim. Padding remains a targeted fix whose end-to-end benefit must be measured.

Run it yourself

A free Colab T4 or any card you own. The build line is in the repo's README:

nvcc -std=c++17 -O3 -arch=sm_75 -o bank_conflicts bank_conflicts.cu

There is no Compiler Explorer embed. Shared memory is not what rules it out: the tiles here are 8 KiB and 4 KiB against a 48 KiB per-block default on a T4. The transpose allocates two 64 MiB buffers and the correctness check transposes 16,777,216 elements on the CPU, which sits too close to Compiler Explorer's 20 second run cap to pin a lesson to.

Nsight Compute counts the conflict directly, with l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum, and the README gives the command. It needs root on a stock driver, so a hosted tier cannot run it.

Exercise

Add strides 3, 8 and 16 to the probe. Write down the degree you expect for each one before you build, from gcd(stride, 32), then run it and report each row's ratio to stride 1. Then say in one sentence why stride 3 does not sit between stride 2 and stride 4.

Time: 25 to 40 minutes. Submit: the three predicted degrees, the three measured ratios, and the sentence.

Check: the harness re-runs the accumulator check on your rows, so a stride that got fast by skipping reads fails instead of winning, and it prints the thread index and the value it found. It then checks the ordering the degrees predict: your stride 3 row must be closer to stride 1 than to stride 4, and your stride 16 row must be slower than your stride 8 row. It reports every ratio whether it passes or not, because the ratios are the answer and the pass is only the gate.

Hint 1

Take stride 3 and write out, for lanes 0 through 31, which bank each one lands in. How many different banks appear in that list, and how many lanes end up sharing one?

Hint 2

The degree is gcd(stride, 32), so check what the stride shares with 32. Which of 3, 8 and 16 share a factor with 32? A stride does not have to be small to be conflict free, it has to be odd.

Solution

Degrees 1, 8 and 16. Stride 3 sends lane L to bank 3L mod 32, and as L runs from 0 to 31 that covers all 32 banks exactly once, so it is conflict free and its ratio should sit next to stride 1's. Stride 8 puts four lanes in each of eight banks; stride 16 puts sixteen lanes in each of two.

Stride 3 does not sit between stride 2 and stride 4 because the degree is not a function of how large the stride is. It is a divisibility question, and 3 shares nothing with 32. What you plotted is not a curve at all: gcd(s, 32) only ever takes the values 1, 2, 4, 8, 16 and 32, and it drops back to 1 at every odd stride.

Keep this rule: shared memory conflict degree depends on the factors a stride shares with 32.

Pitfalls

Your slow kernel has every lane on the same address. That is a broadcast, and the hardware serves it in one request, so the two variants measure the same thing. Before calling an access a conflict, check that the lanes' word indices are distinct.

If two lanes compute the same index, you have written a race, not a conflict. That is the bug in the kernel at the top of this page.

You time a kernel too small to measure. One block, one shared read per thread, and one timed launch make shared memory a rounding error beside launch overhead and module load. Give every kernel a warm-up and enough work for the clock to measure the shared-memory access.

Day 9 explains why.

You add + 1 everywhere and lose blocks per SM. Padding spends shared memory, and on a kernel already near the per-block limit one extra column can cost you a resident block and more than the conflict did. The tiles on this page are far from that limit, so it does not bite here; check your occupancy before assuming the same on a kernel with a big tile.

You pad the allocation and keep the old width in the index. A flat __shared__ float tile[32 * 33]; indexed as tile[row * 32 + col] compiles, runs, reads the wrong element and is still 32-way conflicted, because the index never learned about the extra column. The padded width belongs in one place: let a two-dimensional declaration do the multiply, or name the width once and use that name in every index.

You expect the padding to show up in a kernel that is not shared memory bound. The degree tells you how many requests one instruction becomes. If that instruction is a small part of a kernel waiting on global memory, fixing it changes little.

Measure the kernel before you claim a win.

You profile as a normal user and get nothing. ncu reports ERR_NVGPUCTRPERM, in full: The user running <tool_name/application_name> does not have permission to access NVIDIA GPU Performance Counters or the Hardware Event System on the target device (https://developer.nvidia.com/nvidia-development-tools-solutions-err_nvgpuctrperm-permission-issue-performance-counters , checked 2026-08-30).

On a machine where you have root, sudo ncu works because the stock driver default RmProfilingAdminOnly is 1. Colab and Kaggle give you neither.

Go deeper

Next

Day 16 puts a tile in a loop and reads one of its operands down a column on every iteration, so the arithmetic on this page decides the shape of that tile before you write it. Day 44 changes the rules by loading 16 bytes per lane, which the hardware splits into phases with their own conflict behaviour, and day 73 replaces padding with an XOR swizzle for the tensor-core paths, where the extra column is one the tile cannot spare.