code/day15-bank-conflicts/bank_conflicts.cuThis is the source used by the lesson and its recorded evidence. Compile commands and expected output live in the directory README.
// SPDX-License-Identifier: MIT
//
// Day 15: shared memory bank conflicts, measured.
//
// Two measurements, because either one alone would mislead.
//
// Part 1 isolates the conflict. One kernel reads shared memory kProbeIters
// times per thread at a stride the host varies, so the only thing that
// changes between rows is which bank each lane lands in. The instruction mix
// is identical: stride arrives as a runtime argument, so every row runs the
// same code.
//
// Part 2 is what a conflict is worth inside a kernel you would actually
// write: the tiled transpose, once with a 32-wide tile and once with a
// 33-wide one. Those two kernels differ in one character. A straight copy
// over the same buffers runs first as the ceiling, measured on the same card
// in the same process, so the two transposes are reported as a fraction of
// what this card can stream rather than against each other alone.
//
// Every row within a part moves the same bytes, which is the discipline day
// 11 exists to teach. Part 1's rows each perform kProbeIters shared reads per
// thread; part 2's rows each move 2 * kN * kN floats through global memory.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o bank_conflicts bank_conflicts.cu
// Run: ./bank_conflicts
//
// Verified 2026-08-30 on a Tesla T4 (compute capability 7.5), driver
// 595.84, CUDA 12.6 (V12.6.85). Transcript: evidence/run-2026-08-30.txt
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <cuda_runtime.h>
// The one error macro. This file is standalone, the way a Compiler Explorer
// embed is, so it carries its own verbatim copy. `err_` carries a trailing
// underscore so it cannot collide with a variable at the call site, and the
// do/while makes the macro one statement so it survives a braceless `if`.
#define CUDA_CHECK(call) \
do { \
cudaError_t err_ = (call); \
if (err_ != cudaSuccess) { \
std::fprintf(stderr, "CUDA error %s:%d: %s: %s\n", __FILE__, \
__LINE__, #call, cudaGetErrorString(err_)); \
std::exit(EXIT_FAILURE); \
} \
} while (0)
// 32 banks, 4 bytes each, on every compute capability this course targets.
// Table 31 of the compute capabilities appendix gives 32 for all of them.
constexpr int kBanks = 32;
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr int kWarmupRuns = 3;
constexpr int kTimedRuns = 10;
constexpr float kRelTolerance = 1e-5f;
// Part 1. 2048 words is 8 KiB of shared memory per block, which leaves the
// block-count limit rather than the shared-memory limit deciding occupancy on
// every card in the course's matrix. kProbeIters is large enough that the
// shared reads, not the launch, are what the timer sees.
constexpr unsigned int kProbeWords = 2048;
constexpr unsigned int kProbeMask = kProbeWords - 1;
constexpr int kProbeIters = 2048;
constexpr int kProbeBlocks = 320; // 8 per SM on the T4's 40 SMs
constexpr size_t kProbeOutElems =
static_cast<size_t>(kProbeBlocks) * kThreadsPerBlock;
// Part 2. A 4096 x 4096 matrix is 64 MiB per buffer. kBlockRows is 8, so one
// thread moves four elements and the block is 32 x 8, which is 256 threads.
constexpr int kTileDim = 32;
constexpr int kBlockRows = 8;
constexpr int kN = 4096;
constexpr size_t kMatrixElems = static_cast<size_t>(kN) * kN;
// The conflict degree for a warp whose lane L reads shared word L * stride.
// Lanes L and L' share a bank when (L - L') * stride is a multiple of 32, so
// the colliding lanes are the ones 32 / gcd(stride, 32) apart, there are
// gcd(stride, 32) of them per bank, and their words differ. Euclid, written
// constexpr so the static_asserts below can call it.
constexpr int conflictDegree(int stride) {
int a = stride % kBanks;
int b = kBanks;
while (a != 0) {
const int t = b % a;
b = a;
a = t;
}
return b;
}
static_assert(kThreadsPerBlock % kBanks == 0,
"block size must be a whole number of warps");
static_assert((kProbeWords & (kProbeWords - 1)) == 0,
"the probe masks its index, so the word count must be a power "
"of two");
static_assert(kProbeWords % kBanks == 0,
"masking must not change which bank an index lands in, which "
"needs the word count to be a multiple of the bank count");
static_assert(kTileDim == kBanks,
"the transpose reads a tile column, and the column is a 32-way "
"conflict only because the row length equals the bank count");
static_assert(kTileDim * kBlockRows == kThreadsPerBlock,
"the transpose block is kTileDim by kBlockRows");
static_assert(kTileDim % kBlockRows == 0,
"each thread moves kTileDim / kBlockRows elements");
static_assert(kN % kTileDim == 0,
"the matrix side must be a whole number of tiles, which is why "
"neither transpose kernel carries a bounds check");
// The whole lesson, checked by the compiler. A stride of 32 words puts all 32
// lanes in one bank; a stride of 33 puts them in 33 consecutive banks, which
// wraps to all 32 of them. Adding one column is not a trick, it is the only
// change that makes gcd(width, 32) equal 1.
static_assert(conflictDegree(1) == 1, "stride 1 is conflict free");
static_assert(conflictDegree(2) == 2, "stride 2 is a two-way conflict");
static_assert(conflictDegree(4) == 4, "stride 4 is a four-way conflict");
static_assert(conflictDegree(kTileDim) == kBanks,
"an unpadded 32-wide tile column is a 32-way conflict");
static_assert(conflictDegree(kTileDim + 1) == 1,
"one column of padding makes the same column access conflict "
"free");
// Reads kProbeIters shared words per thread and accumulates them.
//
// One thread: copies its share of `in` into the tile, waits, then walks the
// tile kProbeIters times.
//
// One warp: lane L reads word (L * stride + k) mod kProbeWords, so the 32
// lanes are `stride` words apart and lane L lands in bank (L * stride) mod
// 32. Adding k shifts every lane by the same amount, so the bank pattern is
// the same on every iteration and the degree is conflictDegree(stride).
//
// Launch assumption: kProbeWords is a power of two and a multiple of 32, so
// the mask cannot change which bank an index lands in, and the grid covers
// exactly n threads.
//
// What this measures and what it does not: the accumulator is the reason the
// reads cannot be deleted, not a number anyone wants. Every element of `in`
// is 1.0f, so the sum is kProbeIters exactly and main() checks that, which
// also proves every read happened.
// snippet: probe
__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;
}
}
// end snippet
// The ceiling both transposes are measured against: the same bytes, the same
// launch shape, no transpose and no shared memory. One thread moves
// kTileDim / kBlockRows elements straight through.
//
// One warp: 32 lanes covering 128 contiguous bytes on the read and the same
// on the write, which is the best this access pattern can do. Day 13 times
// the same row, which is what lets the two pages be compared.
__global__ void copyRowMajor(const float* __restrict__ in,
float* __restrict__ out, int n) {
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) {
const size_t elem = static_cast<size_t>(row + k) * n + col;
out[elem] = in[elem];
}
}
// Transposes an n x n matrix through a 32 x 32 shared tile. One thread moves
// kTileDim / kBlockRows elements.
//
// One warp: blockDim.x is 32, so a warp is one row of the block and both
// global accesses cover 128 contiguous bytes. That is why the tile is here at
// all, and day 13 is where it came from.
//
// The shared store walks a tile row, so the 32 lanes land in 32 different
// banks. The shared load walks a tile column: lane L reads word L * 32 + c,
// every lane lands in bank c mod 32, and the 32 words are distinct, so the
// read serializes into 32 requests.
//
// Launch assumption: block is (kTileDim, kBlockRows) and n is a multiple of
// kTileDim, both fixed by static_assert above, which is why there is no
// bounds check under the barrier.
// snippet: transpose-plain
__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];
}
}
// end snippet
// The same kernel with one character changed: the tile is 33 words wide.
//
// Row r now starts at word 33r instead of 32r, so the column read at lane L
// is word L * 33 + c, whose bank is (L + c) mod 32. Every lane gets its own
// bank and the read takes one request instead of 32.
//
// The padding costs 128 bytes of shared memory per block and one integer add
// per shared access, because 33 is not a power of two and the compiler can
// fold 32 into a shift. Both costs are paid by the faster kernel, which is
// the direction that cannot flatter the result.
__global__ void transposeTiledPadded(const float* __restrict__ in,
float* __restrict__ out, int n) {
// snippet: padded-tile
__shared__ float tile[kTileDim][kTileDim + 1];
// end snippet
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];
}
}
// CPU reference. Written for obvious correctness, not speed: plain loops, no
// blocking, no OpenMP. It never allocates; the caller owns every buffer.
static void transposeCpu(const float* in, float* out, int n) {
for (int r = 0; r < n; ++r) {
for (int c = 0; c < n; ++c) {
out[static_cast<size_t>(c) * n + r] =
in[static_cast<size_t>(r) * n + c];
}
}
}
// Returns the first index where got and want differ by more than the relative
// tolerance, or n if they agree everywhere. Returning the index rather than a
// bool is the point: "wrong at 4096" names the tile, "wrong" does not.
static size_t firstMismatch(const float* got, const float* want, size_t n,
float relTolerance) {
for (size_t i = 0; i < n; ++i) {
const float scale = (want[i] == 0.0f) ? 1.0f : std::fabs(want[i]);
if (std::fabs(got[i] - want[i]) > relTolerance * scale) {
return i;
}
}
return n;
}
// Times a launch with CUDA events and returns the mean milliseconds per run.
// A host clock around a launch measures the launch, because launches are
// asynchronous. Day 9 takes that apart. Copy this helper verbatim.
template <typename LaunchFn>
static float timeKernel(LaunchFn launch) {
cudaEvent_t start, stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
// Warm up this kernel, not just the first kernel in the program. Lazy
// module loading has been the default since CUDA 12.2 on Linux, so the
// first launch of each kernel pays its own load.
for (int i = 0; i < kWarmupRuns; ++i) {
launch();
}
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaEventRecord(start));
for (int i = 0; i < kTimedRuns; ++i) {
launch();
}
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
CUDA_CHECK(cudaGetLastError());
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
CUDA_CHECK(cudaEventDestroy(start));
CUDA_CHECK(cudaEventDestroy(stop));
return ms / kTimedRuns;
}
// Both transpose kernels move every element in and every element out, so the
// figure comes from one constant rather than from each kernel, where the two
// could drift apart.
static double transposeGBs(float ms) {
const double bytes =
2.0 * static_cast<double>(kMatrixElems) * sizeof(float);
return bytes / (static_cast<double>(ms) * 1.0e-3) / 1.0e9;
}
int main() {
const int device = 0;
CUDA_CHECK(cudaSetDevice(device));
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
std::printf("GPU: %s (compute capability %d.%d)\n", prop.name, prop.major,
prop.minor);
std::printf("Shared memory per block: %zu bytes\n", prop.sharedMemPerBlock);
// Every failure below records itself and falls through to the one cleanup
// block at the bottom, so no path can return with device memory still
// allocated.
int failures = 0;
float* d_probeIn = nullptr;
float* d_probeOut = nullptr;
float* d_in = nullptr;
float* d_out = nullptr;
const size_t probeInBytes = kProbeWords * sizeof(float);
const size_t probeOutBytes = kProbeOutElems * sizeof(float);
const size_t matrixBytes = kMatrixElems * sizeof(float);
CUDA_CHECK(cudaMalloc(&d_probeIn, probeInBytes));
CUDA_CHECK(cudaMalloc(&d_probeOut, probeOutBytes));
CUDA_CHECK(cudaMalloc(&d_in, matrixBytes));
CUDA_CHECK(cudaMalloc(&d_out, matrixBytes));
// Every probe word is 1.0f, so a thread that performed all kProbeIters
// reads holds exactly kProbeIters, and 2048 is far below the 2^24 where
// float stops counting integers exactly.
std::vector<float> h_probeIn(kProbeWords, 1.0f);
std::vector<float> h_probeOut(kProbeOutElems);
CUDA_CHECK(cudaMemcpy(d_probeIn, h_probeIn.data(), probeInBytes,
cudaMemcpyHostToDevice));
// kMatrixElems is 2^24 exactly, and every integer below 2^24 is exact in
// float, so element i holds a value nothing else in the matrix holds. An
// index bug in either kernel therefore cannot hide behind a repeated
// value.
std::vector<float> h_in(kMatrixElems);
std::vector<float> h_out(kMatrixElems);
std::vector<float> h_want(kMatrixElems);
for (size_t i = 0; i < kMatrixElems; ++i) {
h_in[i] = static_cast<float>(i);
}
CUDA_CHECK(
cudaMemcpy(d_in, h_in.data(), matrixBytes, cudaMemcpyHostToDevice));
transposeCpu(h_in.data(), h_want.data(), kN);
const int strides[] = {1, 2, 4, kTileDim, kTileDim + 1};
const int kNumStrides = sizeof(strides) / sizeof(strides[0]);
float strideMs[kNumStrides];
std::printf("\nPart 1: %u shared words per block, %d reads per thread\n",
kProbeWords, kProbeIters);
std::printf("%8s %10s %12s %12s\n", "stride", "degree", "time (ms)",
"vs stride 1");
std::printf("%8s %10s %12s %12s\n", "------", "------", "----------",
"-----------");
int rows = 0;
for (int s = 0; s < kNumStrides; ++s) {
const int stride = strides[s];
// One untimed launch first, checked, so a wrong answer is reported
// before a number that came from it reaches the table.
probeSharedStride<<<kProbeBlocks, kThreadsPerBlock>>>(
d_probeIn, d_probeOut, kProbeOutElems, stride);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_probeOut.data(), d_probeOut, probeOutBytes,
cudaMemcpyDeviceToHost));
for (size_t t = 0; t < kProbeOutElems; ++t) {
if (h_probeOut[t] != static_cast<float>(kProbeIters)) {
std::fprintf(stderr,
"probe stride %d wrong at thread %zu: got %.9g, "
"want %d\n",
stride, t, h_probeOut[t], kProbeIters);
++failures;
break;
}
}
strideMs[s] = timeKernel([&] {
probeSharedStride<<<kProbeBlocks, kThreadsPerBlock>>>(
d_probeIn, d_probeOut, kProbeOutElems, stride);
});
++rows;
std::printf("%8d %10d %12.3f %12.2f\n", stride, conflictDegree(stride),
strideMs[s], strideMs[s] / strideMs[0]);
}
const dim3 tileBlock(kTileDim, kBlockRows);
const dim3 tileGrid(kN / kTileDim, kN / kTileDim);
std::printf("\nPart 2: %d x %d transpose, %.0f MiB moved per launch\n", kN,
kN, 2.0 * static_cast<double>(matrixBytes) / (1024.0 * 1024.0));
std::printf("%-22s %12s %12s %12s\n", "kernel", "time (ms)", "GB/s",
"% of copy");
std::printf("%-22s %12s %12s %12s\n", "----------------------",
"----------", "----------", "----------");
// The copy is the ceiling, so it is measured first and on the same card in
// the same process. Its answer is its input, so it checks against h_in.
copyRowMajor<<<tileGrid, tileBlock>>>(d_in, d_out, kN);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_out.data(), d_out, matrixBytes, cudaMemcpyDeviceToHost));
size_t bad =
firstMismatch(h_out.data(), h_in.data(), kMatrixElems, kRelTolerance);
if (bad != kMatrixElems) {
std::fprintf(stderr, "copyRowMajor wrong at %zu: got %.9g, want %.9g\n",
bad, h_out[bad], h_in[bad]);
++failures;
}
const float copyMs = timeKernel(
[&] { copyRowMajor<<<tileGrid, tileBlock>>>(d_in, d_out, kN); });
++rows;
std::printf("%-22s %12.3f %12.1f %12.1f\n", "copy, no transpose", copyMs,
transposeGBs(copyMs), 100.0);
transposeTiled<<<tileGrid, tileBlock>>>(d_in, d_out, kN);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_out.data(), d_out, matrixBytes, cudaMemcpyDeviceToHost));
bad =
firstMismatch(h_out.data(), h_want.data(), kMatrixElems, kRelTolerance);
if (bad != kMatrixElems) {
std::fprintf(stderr,
"transposeTiled wrong at %zu: got %.9g, want %.9g\n", bad,
h_out[bad], h_want[bad]);
++failures;
}
const float plainMs = timeKernel(
[&] { transposeTiled<<<tileGrid, tileBlock>>>(d_in, d_out, kN); });
++rows;
std::printf("%-22s %12.3f %12.1f %12.1f\n", "transpose tile[32][32]",
plainMs, transposeGBs(plainMs), 100.0 * copyMs / plainMs);
transposeTiledPadded<<<tileGrid, tileBlock>>>(d_in, d_out, kN);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_out.data(), d_out, matrixBytes, cudaMemcpyDeviceToHost));
bad =
firstMismatch(h_out.data(), h_want.data(), kMatrixElems, kRelTolerance);
if (bad != kMatrixElems) {
std::fprintf(stderr,
"transposeTiledPadded wrong at %zu: got %.9g, want %.9g\n",
bad, h_out[bad], h_want[bad]);
++failures;
}
const float paddedMs = timeKernel([&] {
transposeTiledPadded<<<tileGrid, tileBlock>>>(d_in, d_out, kN);
});
++rows;
std::printf("%-22s %12.3f %12.1f %12.1f\n", "transpose tile[32][33]",
paddedMs, transposeGBs(paddedMs), 100.0 * copyMs / paddedMs);
std::printf("\npadded is %.2fx the plain kernel's time\n",
paddedMs / plainMs);
std::printf(
"\nAll three rows above moved the same %.0f MiB and all three read\n"
"and write global memory in 128-byte runs. The only difference\n"
"between the last two is the width of the shared tile.\n",
2.0 * static_cast<double>(matrixBytes) / (1024.0 * 1024.0));
// The row count is checked so the lesson's tables cannot drift from what
// the program prints. A real branch, not an assert: CI builds Release,
// Release defines NDEBUG, and NDEBUG deletes assert(), so the check would
// be missing from exactly the build that matters.
const int kExpectedRows = kNumStrides + 3;
if (rows != kExpectedRows) {
std::fprintf(stderr,
"printed %d rows, expected %d; the lesson's tables and "
"this program disagree\n",
rows, kExpectedRows);
++failures;
}
CUDA_CHECK(cudaFree(d_probeIn));
CUDA_CHECK(cudaFree(d_probeOut));
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_out));
if (failures != 0) {
std::fprintf(stderr, "%d check(s) failed\n", failures);
return EXIT_FAILURE;
}
return EXIT_SUCCESS;
}