code/day70-bug-catalog/bug_catalog.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 70 checkpoint: one pipeline, ten deliberate bugs, one file.
//
// The pipeline reads 479 sweeps of 611 sensors (row-major, one row per
// sweep), smooths the raw stream, calibrates each sensor, sums per-block
// energy as a QA figure, and totals each sensor's column. Every stage has a
// CPU reference and a gate.
//
// The same file builds two programs. The switch is DAY70_FIXED:
//
// buggy: nvcc -std=c++17 -O3 -DNDEBUG -lineinfo -arch=sm_75 \
// -o bug_catalog_buggy bug_catalog.cu
// fixed: nvcc -std=c++17 -O3 -DDAY70_FIXED=1 -lineinfo -arch=sm_75 \
// -o bug_catalog_fixed bug_catalog.cu
//
// -DNDEBUG on the buggy line is not decoration. It is what CMake's Release
// configuration adds, and it is what deletes bug 9's assert. Build the buggy
// variant without it and the assert fires, which is the trap in full: green
// in the build you test, silent in the build you ship.
//
// The ten bugs, each marked "BUG n (deliberate, day 70)" at its site:
//
// 1 unchecked kernel launch no tool; day 6's two lines
// 2 cudaMemcpyKind contradicts the no tool on this card; day 6
// pointers measured it copying anyway
// 3 missing bounds guard memcheck: 195 reads + 195 writes
// 4 missing __syncthreads after a racecheck, in blurSamples
// shared tile load
// 5 early return above a barrier synccheck, in blockEnergy
// 6 device buffer read, never written initcheck, in calibrateSamples
// 7 one cudaFree missing memcheck --leak-check full
// 8 sensor-major index into no tool; only the CPU reference
// sweep-major storage
// 9 assert() as the only gate no tool; NDEBUG deletes it
// 10 exact float compare in a gate no tool; flags a correct kernel
//
// Five of the ten produce no sanitizer report of any kind. That number is
// the reason this checkpoint exists.
//
#include <cassert>
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <cuda_runtime.h>
// The build switch. Not a value, so the no-#define-for-values rule does not
// apply: the preprocessor is the only tool that can add and remove whole
// statements (a barrier, a guard, an error check) from one source file.
#ifndef DAY70_FIXED
#define DAY70_FIXED 0
#endif
// 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)
// 479 x 611 = 292,669 elements, chosen so nothing divides evenly: 1144
// blocks of 256 launch 292,864 threads, 195 of them past the end. Those 195
// are bug 3's exact memcheck arithmetic and bug 5's one divergent block.
//
// kGain is 1/3 in float, a dense mantissa on purpose: the device fuses
// kGain * x + offset into one FMA and the host reference does not, so the
// two sides differ in the last bit on some elements. That difference is what
// bug 10's exact compare trips over and the relative tolerance absorbs.
constexpr int kRows = 479; // sweeps
constexpr int kCols = 611; // sensors
constexpr size_t kElems = static_cast<size_t>(kRows) * kCols;
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr float kGain = 1.0f / 3.0f;
constexpr float kRelTolerance = 1e-5f;
static_assert(kThreadsPerBlock % 32 == 0,
"block size must be a whole number of warps");
static_assert((kThreadsPerBlock & (kThreadsPerBlock - 1)) == 0,
"blockEnergy halves the stride, so the block size must be a "
"power of two");
// Three-tap box blur over the sample stream, edges clamped. One thread owns
// one output sample and stages its input sample in a shared tile with a
// one-cell halo on each side.
//
// Memory: consecutive threads load consecutive samples, one warp covers 128
// contiguous bytes. Each thread then reads its two shared-memory neighbours,
// which is why the barrier below is load-bearing.
//
// Launch assumption: gridDim.x * blockDim.x >= n, blockDim.x is exactly
// kThreadsPerBlock (the tile is sized from the constant).
// snippet: race-bug
__global__ void blurSamples(const float* __restrict__ in,
float* __restrict__ out, size_t n) {
__shared__ float tile[kThreadsPerBlock + 2];
const unsigned int tid = threadIdx.x;
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + tid;
// Clamped loads: a thread past the end restages the last sample, so
// every tile cell is written and the maths at the edges stays defined.
tile[tid + 1] = in[i < n ? i : n - 1];
if (tid == 0) {
tile[0] = in[i == 0 ? 0 : i - 1];
}
if (tid == kThreadsPerBlock - 1) {
const size_t gi = i + 1;
tile[kThreadsPerBlock + 1] = in[gi < n ? gi : n - 1];
}
// BUG 4 (deliberate, day 70): the buggy build omits this barrier, so a
// thread can read tile cells its neighbours have not written yet. The
// wrong answers come and go with the warp schedule; racecheck reports
// the hazard on every run, including the runs whose output is right.
#if DAY70_FIXED
__syncthreads();
#endif
const float val = (tile[tid] + tile[tid + 1] + tile[tid + 2]) / 3.0f;
if (i < n) {
out[i] = val;
}
}
// end snippet
// Sums one block's 256 blurred samples into out[blockIdx.x], the pipeline's
// per-block QA figure.
//
// Memory: the load is one coalesced sweep; the halving loop then works
// entirely in shared memory.
//
// Launch assumption: blockDim.x is exactly kThreadsPerBlock, a power of two.
// snippet: barrier-bug
__global__ void blockEnergy(const float* __restrict__ in,
float* __restrict__ out, size_t n) {
__shared__ float red[kThreadsPerBlock];
const unsigned int tid = threadIdx.x;
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + tid;
#if DAY70_FIXED
red[tid] = (i < n) ? in[i] : 0.0f;
#else
// BUG 5 (deliberate, day 70): 195 threads of the last block leave here,
// above every barrier below. Since Volta an exited thread counts as
// having arrived, so nothing hangs; their red[] cells simply hold
// whatever was there, and the last block's energy is garbage. The guard
// belongs on the load, never on the whole thread.
if (i >= n) {
return;
}
red[tid] = in[i];
#endif
__syncthreads();
for (unsigned int half = kThreadsPerBlock / 2; half > 0; half /= 2) {
if (tid < half) {
red[tid] += red[tid + half];
}
__syncthreads();
}
if (tid == 0) {
out[blockIdx.x] = red[0];
}
}
// end snippet
// out[i] = kGain * in[i] + offset[i % kCols]: per-sensor gain and offset.
// One thread owns one element.
//
// Memory: in and out are coalesced; offset is a 611-float table every warp
// reads through the cache.
//
// Launch assumption: gridDim.x * blockDim.x >= n, and the guard makes the
// 195-thread overshoot harmless. Removing it is bug 3.
// snippet: bounds-bug
__global__ void calibrateSamples(const float* __restrict__ in,
const float* __restrict__ offset,
float* __restrict__ out, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
const int c = static_cast<int>(i % static_cast<size_t>(kCols));
#if DAY70_FIXED
if (i < n) {
out[i] = kGain * in[i] + offset[c];
}
#else
// BUG 3 (deliberate, day 70): no guard, so each of the 195 threads past
// the end reads and writes 4 bytes beyond in[] and out[], the farthest
// landing 780 bytes past the allocations. cudaMalloc's rounding means
// none of it should fault; memcheck counts every access anyway.
out[i] = kGain * in[i] + offset[c];
#endif
}
// end snippet
// Sums each sensor's 479 calibrated readings. One thread owns one sensor
// (one column).
//
// Memory: at each r, consecutive threads read consecutive elements of row r,
// so the column walk is coalesced across the warp.
//
// Launch assumption: gridDim.x * blockDim.x >= cols.
// snippet: index-bug
__global__ void sensorTotals(const float* __restrict__ in,
float* __restrict__ out, int rows, int cols) {
const size_t c = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (c < static_cast<size_t>(cols)) {
float total = 0.0f;
for (int r = 0; r < rows; ++r) {
#if DAY70_FIXED
total += in[static_cast<size_t>(r) * cols + c];
#else
// BUG 8 (deliberate, day 70): sensor-major indexing of
// sweep-major storage. Every access lands inside the buffer
// (610 * 479 + 478 < 292,669), so no tool has anything to say.
// The totals are sums of the wrong cells, and only the CPU
// reference notices.
total += in[c * static_cast<size_t>(rows) + r];
#endif
}
out[c] = total;
}
}
// end snippet
// CPU references. Written for obvious correctness, not speed: plain loops,
// no OpenMP, no intrinsics. None of them allocates; the caller owns every
// buffer. The sums accumulate in double because the reference's job is to be
// right, not to match the kernels bit for bit; the tolerance covers the
// difference.
static void blurCpu(const float* in, float* out, size_t n) {
for (size_t i = 0; i < n; ++i) {
const float left = in[i == 0 ? 0 : i - 1];
const float right = in[i + 1 < n ? i + 1 : n - 1];
out[i] = (left + in[i] + right) / 3.0f;
}
}
static void energyCpu(const float* blurred, double* out, size_t n, int blocks) {
for (int b = 0; b < blocks; ++b) {
double total = 0.0;
const size_t lo = static_cast<size_t>(b) * kThreadsPerBlock;
const size_t hi = lo + kThreadsPerBlock < n ? lo + kThreadsPerBlock : n;
for (size_t i = lo; i < hi; ++i) {
total += static_cast<double>(blurred[i]);
}
out[b] = total;
}
}
// Kept in float and written as a separate multiply and add. The device side
// fuses the same expression into one FMA, so on hosts whose baseline has no
// FMA instruction (plain x86-64) the two sides differ in the last bit on
// some elements. Bug 10's exact compare exists to trip over exactly that.
static void calibrateCpu(const float* in, const float* offset, float* out,
size_t n) {
for (size_t i = 0; i < n; ++i) {
const float scaled = kGain * in[i];
out[i] = scaled + offset[i % static_cast<size_t>(kCols)];
}
}
static void totalsCpu(const float* in, double* out, int rows, int cols) {
for (int c = 0; c < cols; ++c) {
double total = 0.0;
for (int r = 0; r < rows; ++r) {
total += static_cast<double>(in[static_cast<size_t>(r) * cols + c]);
}
out[c] = total;
}
}
// Counts elements outside the relative tolerance and records the first one.
// A count, not a bool: "17 wrong starting at 292608" places the bug in a
// block; "wrong" places it nowhere.
static size_t countMismatches(const float* got, const double* want, size_t n,
float relTolerance, size_t* firstBad) {
size_t bad = 0;
*firstBad = n;
for (size_t i = 0; i < n; ++i) {
const double w = want[i];
const double scale = (w == 0.0) ? 1.0 : std::fabs(w);
if (std::fabs(static_cast<double>(got[i]) - w) >
static_cast<double>(relTolerance) * scale) {
if (bad == 0) {
*firstBad = i;
}
++bad;
}
}
return bad;
}
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("build: %s\n",
DAY70_FIXED ? "fixed" : "buggy (ten deliberate bugs)");
std::printf("%d sweeps x %d sensors = %zu samples\n", kRows, kCols, kElems);
const size_t bytes = kElems * sizeof(float);
const int blocks =
static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
const int colBlocks = (kCols + kThreadsPerBlock - 1) / kThreadsPerBlock;
std::printf("%d blocks of %d, %zu threads past the end\n", blocks,
kThreadsPerBlock,
static_cast<size_t>(blocks) * kThreadsPerBlock - kElems);
std::vector<float> h_raw(kElems);
std::vector<float> h_offset(kCols);
std::vector<float> h_blur(kElems);
std::vector<float> h_norm(kElems);
std::vector<float> h_energy(blocks);
std::vector<float> h_totals(kCols);
for (size_t i = 0; i < kElems; ++i) {
h_raw[i] = 1.0f + static_cast<float>(i % 4093) * 0.000244140625f;
}
for (int c = 0; c < kCols; ++c) {
h_offset[c] = 0.05f + static_cast<float>(c) * 0.001f;
}
// d_norm is allocated last so bug 3's 780 bytes of overshoot land in the
// allocator's padding rather than in a neighbouring buffer. That keeps
// the buggy run deterministic enough to grade; memcheck flags the writes
// either way.
float* d_raw = nullptr;
float* d_offset = nullptr;
float* d_blur = nullptr;
float* d_energy = nullptr;
float* d_totals = nullptr;
float* d_norm = nullptr;
CUDA_CHECK(cudaMalloc(&d_raw, bytes));
CUDA_CHECK(cudaMalloc(&d_offset, kCols * sizeof(float)));
CUDA_CHECK(cudaMalloc(&d_blur, bytes));
CUDA_CHECK(cudaMalloc(&d_energy, blocks * sizeof(float)));
CUDA_CHECK(cudaMalloc(&d_totals, kCols * sizeof(float)));
CUDA_CHECK(cudaMalloc(&d_norm, bytes));
// snippet: copy-bugs
#if DAY70_FIXED
CUDA_CHECK(cudaMemcpy(d_raw, h_raw.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_offset, h_offset.data(), kCols * sizeof(float),
cudaMemcpyHostToDevice));
#else
// BUG 2 (deliberate, day 70): the kind argument says device-to-host and
// the pointers say the opposite. Day 6 measured this on a Tesla T4 with
// CUDA 12.6: unified addressing lets the runtime take the direction from
// the pointers, the call returns cudaSuccess and copies correctly, so
// the CUDA_CHECK around it passes. The documentation calls a mismatch
// undefined behaviour, so on another platform this same line may be the
// one that fails.
CUDA_CHECK(cudaMemcpy(d_raw, h_raw.data(), bytes, cudaMemcpyDeviceToHost));
// BUG 6 (deliberate, day 70): the offset table is never copied at all.
// The kernel reads 2,444 bytes of memory nothing ever wrote. No launch
// fails, nothing faults; initcheck is the only tool that names it.
#endif
// end snippet
// Stage 1: blur.
// snippet: unchecked-launch
blurSamples<<<blocks, kThreadsPerBlock>>>(d_raw, d_blur, kElems);
#if DAY70_FIXED
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
#else
// BUG 1 (deliberate, day 70): no cudaGetLastError, no synchronize. This
// launch happens to be legal, so nothing is lost today. The day the
// grid or the kernel goes wrong, the failure surfaces at whichever call
// looks next, and that call is innocent. Day 6's two lines, always.
#endif
// end snippet
// Stage 1b: per-block energy.
blockEnergy<<<blocks, kThreadsPerBlock>>>(d_blur, d_energy, kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
// Stage 2: calibrate.
calibrateSamples<<<blocks, kThreadsPerBlock>>>(d_blur, d_offset, d_norm,
kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
// Stage 3: per-sensor totals.
sensorTotals<<<colBlocks, kThreadsPerBlock>>>(d_norm, d_totals, kRows,
kCols);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_blur.data(), d_blur, bytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(h_energy.data(), d_energy, blocks * sizeof(float),
cudaMemcpyDeviceToHost));
CUDA_CHECK(
cudaMemcpy(h_norm.data(), d_norm, bytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(h_totals.data(), d_totals, kCols * sizeof(float),
cudaMemcpyDeviceToHost));
// References.
std::vector<float> h_wantBlur(kElems);
std::vector<float> h_wantNorm(kElems);
std::vector<double> h_wantBlurD(kElems);
std::vector<double> h_wantNormD(kElems);
std::vector<double> h_wantEnergy(blocks);
std::vector<double> h_wantTotals(kCols);
blurCpu(h_raw.data(), h_wantBlur.data(), kElems);
energyCpu(h_wantBlur.data(), h_wantEnergy.data(), kElems, blocks);
calibrateCpu(h_wantBlur.data(), h_offset.data(), h_wantNorm.data(), kElems);
totalsCpu(h_wantNorm.data(), h_wantTotals.data(), kRows, kCols);
for (size_t i = 0; i < kElems; ++i) {
h_wantBlurD[i] = static_cast<double>(h_wantBlur[i]);
h_wantNormD[i] = static_cast<double>(h_wantNorm[i]);
}
// Gates. status falls through: every stage reports, everything is freed
// on every path, and the exit code carries the verdict at the end.
int status = EXIT_SUCCESS;
size_t firstBad = 0;
const size_t blurBad = countMismatches(h_blur.data(), h_wantBlurD.data(),
kElems, kRelTolerance, &firstBad);
if (blurBad != 0) {
std::fprintf(stderr, "stage 1 blur: %zu of %zu wrong, first at %zu\n",
blurBad, kElems, firstBad);
status = EXIT_FAILURE;
} else {
std::printf("stage 1 blur: all %zu samples match\n", kElems);
}
const size_t energyBad =
countMismatches(h_energy.data(), h_wantEnergy.data(),
static_cast<size_t>(blocks), kRelTolerance, &firstBad);
// snippet: assert-bug
#if DAY70_FIXED
if (energyBad != 0) {
std::fprintf(stderr,
"stage 1 energy: %zu of %d blocks wrong, "
"first at %zu\n",
energyBad, blocks, firstBad);
status = EXIT_FAILURE;
} else {
std::printf("stage 1 energy: all %d blocks match\n", blocks);
}
#else
// BUG 9 (deliberate, day 70): the QA stage's only gate is an assert, and
// the buggy build line defines NDEBUG, which expands it to nothing. Bug
// 5 corrupts exactly one block's energy and this line is the one that
// would have said so. The extra parentheses keep the course linter's
// no-runtime-assert rule from firing on the line that exists to show why
// that rule exists.
(assert(energyBad == 0));
std::printf("stage 1 energy: checked\n");
#endif
// end snippet
// snippet: tolerance-bug
#if DAY70_FIXED
const size_t normBad = countMismatches(h_norm.data(), h_wantNormD.data(),
kElems, kRelTolerance, &firstBad);
#else
// BUG 10 (deliberate, day 70): an exact compare between a fused device
// FMA and an unfused host multiply-add. Fix bugs 1 through 9 and this
// gate still cries wrong, over last-bit differences the day 5 tolerance
// was built to absorb. A gate that fails on a correct kernel costs you
// the week you spend debugging the kernel.
size_t normBad = 0;
firstBad = kElems;
for (size_t i = 0; i < kElems; ++i) {
if (h_norm[i] != h_wantNorm[i]) {
if (normBad == 0) {
firstBad = i;
}
++normBad;
}
}
#endif
// end snippet
if (normBad != 0) {
std::fprintf(stderr,
"stage 2 calibrate: %zu of %zu wrong, first at "
"%zu\n",
normBad, kElems, firstBad);
status = EXIT_FAILURE;
} else {
std::printf("stage 2 calibrate: all %zu samples match\n", kElems);
}
const size_t totalsBad =
countMismatches(h_totals.data(), h_wantTotals.data(),
static_cast<size_t>(kCols), kRelTolerance, &firstBad);
if (totalsBad != 0) {
std::fprintf(stderr,
"stage 3 totals: %zu of %d sensors wrong, "
"first at %zu\n",
totalsBad, kCols, firstBad);
status = EXIT_FAILURE;
} else {
std::printf("stage 3 totals: all %d sensors match\n", kCols);
}
// snippet: leak-bug
CUDA_CHECK(cudaFree(d_offset));
CUDA_CHECK(cudaFree(d_blur));
CUDA_CHECK(cudaFree(d_energy));
CUDA_CHECK(cudaFree(d_totals));
CUDA_CHECK(cudaFree(d_norm));
#if DAY70_FIXED
CUDA_CHECK(cudaFree(d_raw));
#else
// BUG 7 (deliberate, day 70): d_raw is never freed. The process exit
// will reclaim it, so nothing visible happens on a run this short; a
// loop that allocates per iteration dies of this in production.
// compute-sanitizer with --leak-check full prints the allocation and
// its size at context teardown.
#endif
// end snippet
std::printf("pipeline: %s\n", status == EXIT_SUCCESS ? "PASS" : "FAIL");
return status;
}