code/day61-sanitizer/memory_bugs.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 61: three planted memory bugs that pass every test.
//
// Two small kernels, each in a correct and a deliberately broken variant:
// an adjacent difference (broken variant reads one element past the input)
// and a radius-1 windowed sum staged through shared memory (broken variant
// writes one cell past its tile). A third bug is a cudaMalloc with no
// cudaFree. The bare run checks every element of every kernel against a
// CPU reference and exits 0, because none of the three bugs reaches any
// byte the harness looks at. compute-sanitizer is the only thing here
// that can tell the broken variants from the fixed ones, which is the
// lesson.
//
// The switch: kRunBuggyPass below. true (the supplied state) runs the
// fixed pass and then the buggy pass plus the leak. false runs only the
// fixed pass and frees everything, and the same binary comes back clean
// under memcheck and --leak-check full.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o memory_bugs
// memory_bugs.cu
// Run: ./memory_bugs
// Tool: compute-sanitizer commands are in README.md.
//
#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)
// 6151 elements is 24 full blocks of 256 plus 7, so every guard runs, and
// the last element sits at thread (6,0,0) of block (24,0,0). Those are the
// coordinates the memcheck report of bug 1 must name, derived here rather
// than read back from the tool.
constexpr size_t kElems = 6151;
constexpr int kThreadsPerBlock = 256; // 8 warps
// The tile holds the block's 256 elements plus one halo cell each side.
constexpr int kTileCells = kThreadsPerBlock + 2;
// The switch between the supplied program and the clean one. Supplied
// state is true: fixed pass, then buggy pass, then the leak.
constexpr bool kRunBuggyPass = true;
static_assert(kThreadsPerBlock % 32 == 0,
"block size must be a whole number of warps");
// Fixed adjacent difference: out[i] = in[i+1] - in[i], and out[n-1] = 0.
// One thread owns one element.
//
// Memory: consecutive threads read and write consecutive elements, fully
// coalesced; each thread also reads its right neighbour, which the warp
// to its right just loaded, so the overlap is served by cache.
//
// Launch assumption: gridDim.x * blockDim.x >= n.
__global__ void diffNext(const unsigned int* __restrict__ in,
unsigned int* __restrict__ out, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
out[i] = (i + 1 < n) ? in[i + 1] - in[i] : 0u;
}
}
// snippet: buggy-read
// The same kernel with the guard moved after the load.
//
// Deliberate bug 1 (day 61): the load of in[i + 1] is unconditional, so
// the last thread (i == n - 1, thread (6,0,0) of block (24,0,0) at these
// sizes) reads in[n], four bytes past the end of a 24,604-byte
// allocation. The bare run cannot see it: cudaMalloc hands out more than
// it promises, the hardware only faults at the edge of a mapping, and
// the select below throws the loaded value away for exactly that thread.
// The store through `tap` is why the compiler cannot quietly guard the
// load for us: the value is always used, just never checked.
__global__ void diffNextPastEnd(const unsigned int* __restrict__ in,
unsigned int* __restrict__ out,
unsigned int* __restrict__ tap, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
const unsigned int next = in[i + 1]; // reads in[n] when i == n - 1
tap[i] = next;
out[i] = (i + 1 < n) ? next - in[i] : 0u;
}
}
// end snippet
// Fixed windowed sum: out[i] = in[i-1] + in[i] + in[i+1], with zeros past
// either end of the array. One thread owns one output element.
//
// Memory: the staging loop and the output store are both consecutive
// threads on consecutive addresses, fully coalesced. tile[c] holds
// in[blockStart + c - 1]; threads 0 and 1 wrap around once to load the
// two cells past their own.
//
// Launch assumption: blockDim.x == kThreadsPerBlock, because the tile is
// sized from the constant.
__global__ void windowSum(const unsigned int* __restrict__ in,
unsigned int* __restrict__ out, size_t n) {
__shared__ unsigned int tile[kTileCells];
const unsigned int tid = threadIdx.x;
const size_t blockStart = blockIdx.x * static_cast<size_t>(blockDim.x);
const size_t i = blockStart + tid;
for (int c = static_cast<int>(tid); c < kTileCells; c += kThreadsPerBlock) {
const size_t g = blockStart + static_cast<size_t>(c);
tile[c] = (g >= 1 && g <= n) ? in[g - 1] : 0u;
}
__syncthreads();
if (i < n) {
out[i] = tile[tid] + tile[tid + 1] + tile[tid + 2];
}
}
// The same kernel with <= where < belongs in the staging loop.
//
// Deliberate bug 2 (day 61): thread (2,0,0) of every block runs the loop
// once more at c == kTileCells and writes tile[258], one cell past the
// 258-cell tile. Every cell the sum reads is still written correctly, and
// the stray cell lands in the block's own shared-memory padding: shared
// memory is allocated to blocks in 256-byte units on compute capability
// 7.x (cudaOccSMemAllocationGranularity in the toolkit's
// cuda_occupancy.h, checked 2026-09-01), so a 1,032-byte tile owns 1,280
// bytes. The global load stays guarded, so this is one shared write per
// block and nothing else.
__global__ void windowSumOverrun(const unsigned int* __restrict__ in,
unsigned int* __restrict__ out, size_t n) {
__shared__ unsigned int tile[kTileCells];
const unsigned int tid = threadIdx.x;
const size_t blockStart = blockIdx.x * static_cast<size_t>(blockDim.x);
const size_t i = blockStart + tid;
// snippet: buggy-halo
for (int c = static_cast<int>(tid); c <= kTileCells;
c += kThreadsPerBlock) { // <= runs one cell too far
const size_t g = blockStart + static_cast<size_t>(c);
tile[c] = (g >= 1 && g <= n) ? in[g - 1] : 0u;
}
// end snippet
__syncthreads();
if (i < n) {
out[i] = tile[tid] + tile[tid + 1] + tile[tid + 2];
}
}
// CPU references. Written for obvious correctness; they never allocate.
static void diffNextCpu(const unsigned int* in, unsigned int* out, size_t n) {
for (size_t i = 0; i < n; ++i) {
out[i] = (i + 1 < n) ? in[i + 1] - in[i] : 0u;
}
}
static void windowSumCpu(const unsigned int* in, unsigned int* out, size_t n) {
for (size_t i = 0; i < n; ++i) {
const unsigned int left = (i >= 1) ? in[i - 1] : 0u;
const unsigned int right = (i + 1 < n) ? in[i + 1] : 0u;
out[i] = left + in[i] + right;
}
}
// Exact compare: everything here is unsigned integer arithmetic, so the
// two sides must match bit for bit and there is no tolerance. Returns the
// first differing index, or n.
static size_t firstMismatch(const unsigned int* got, const unsigned int* want,
size_t n) {
for (size_t i = 0; i < n; ++i) {
if (got[i] != want[i]) {
return i;
}
}
return n;
}
// Runs one already-launched kernel's result through the reference. On a
// mismatch it reports and flips status, then falls through: cleanup at
// the bottom of main must run on every path, or a --leak-check transcript
// reports every live buffer instead of the one real leak.
static void checkResult(const char* name, const unsigned int* got,
const unsigned int* want, size_t n, int* status) {
const size_t bad = firstMismatch(got, want, n);
if (bad == n) {
std::printf("%-18s all %zu elements match\n", name, n);
} else {
std::fprintf(stderr, "%-18s wrong at %zu: got %u, want %u\n", name, bad,
got[bad], want[bad]);
*status = EXIT_FAILURE;
}
}
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);
const size_t bytes = kElems * sizeof(unsigned int);
const int blocks =
static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
std::printf("n = %zu, %d blocks of %d, %zu bytes per buffer\n", kElems,
blocks, kThreadsPerBlock, bytes);
std::vector<unsigned int> h_in(kElems);
std::vector<unsigned int> h_out(kElems);
std::vector<unsigned int> h_want(kElems);
for (size_t i = 0; i < kElems; ++i) {
h_in[i] = static_cast<unsigned int>((i % 97) * 3 + (i % 13));
}
unsigned int* d_in = nullptr;
unsigned int* d_out = nullptr;
CUDA_CHECK(cudaMalloc(&d_in, bytes));
CUDA_CHECK(cudaMalloc(&d_out, bytes));
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));
int status = EXIT_SUCCESS;
// Fixed pass. Both correct kernels run and are checked first, so a
// sanitizer transcript of this binary shows two clean kernels before
// any report, and the report's kernel names carry the diagnosis.
diffNextCpu(h_in.data(), h_want.data(), kElems);
diffNext<<<blocks, kThreadsPerBlock>>>(d_in, d_out, kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
checkResult("diffNext", h_out.data(), h_want.data(), kElems, &status);
windowSumCpu(h_in.data(), h_want.data(), kElems);
windowSum<<<blocks, kThreadsPerBlock>>>(d_in, d_out, kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
checkResult("windowSum", h_out.data(), h_want.data(), kElems, &status);
if (kRunBuggyPass) {
// snippet: leak
// Deliberate bug 3 (day 61): a scratch buffer somebody added while
// debugging bug 1 and never removed. It is written by every thread
// of diffNextPastEnd, read by nobody, and freed nowhere on any
// path, so --leak-check full must report exactly one leaked
// allocation of 24,604 bytes.
unsigned int* d_tap = nullptr;
CUDA_CHECK(cudaMalloc(&d_tap, bytes));
// end snippet
diffNextCpu(h_in.data(), h_want.data(), kElems);
diffNextPastEnd<<<blocks, kThreadsPerBlock>>>(d_in, d_out, d_tap,
kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
checkResult("diffNextPastEnd", h_out.data(), h_want.data(), kElems,
&status);
windowSumCpu(h_in.data(), h_want.data(), kElems);
windowSumOverrun<<<blocks, kThreadsPerBlock>>>(d_in, d_out, kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
checkResult("windowSumOverrun", h_out.data(), h_want.data(), kElems,
&status);
if (status == EXIT_SUCCESS) {
std::printf(
"every check passed. two of these kernels and one\n"
"allocation are still wrong; see README.md for the\n"
"compute-sanitizer runs that prove it.\n");
}
}
// One cleanup block, reached on every path. d_tap is deliberately
// absent from it; that absence is bug 3.
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_out));
return status;
}