code/day63-cuda-gdb/scale_debug.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 63: one wrong index, found from inside the kernel with cuda-gdb.
//
// scaleBuggy carries a deliberate bug: the store index is transposed,
// col * height + row where row * width + col belongs. Day 7 named this
// exact mistake and showed that it never leaves the buffer, so no error
// check and no sanitizer tool fires. scaleFixed is the same kernel with
// the index written once and used twice. The harness requires the buggy
// kernel to miss on exactly the count the transposition predicts and the
// fixed kernel to match everywhere, so the two variants gate each other.
//
// The debugger walk is scripted in commands.gdb, so the transcript the
// lesson quotes can be reproduced with one command.
//
// Build (release, for the timing rows):
// nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o scale_debug scale_debug.cu
// Build (device debug, for cuda-gdb):
// nvcc -std=c++17 -g -G -arch=sm_75 -o scale_debug_g scale_debug.cu
// Run: ./scale_debug
// Debug: cuda-gdb --batch -x commands.gdb ./scale_debug_g
//
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <cuda_runtime.h>
// The one error macro. This file is standalone, so it carries its own
// verbatim copy. `err_` has a trailing underscore so it cannot collide with a
// variable at the call site.
#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)
// 61 by 37 keeps the whole image small enough that a debugger transcript
// stays readable, and the two sides differ so the transposed index cannot
// hide behind a square, which is how day 7 said the habit survives.
constexpr size_t kWidth = 61;
constexpr size_t kHeight = 37;
constexpr size_t kElems = kWidth * kHeight; // 2,257
// 16 by 8 is the course's 2D shape for a 128-thread block: four whole
// warps, threadIdx.x fastest, so one warp is two rows of sixteen columns.
constexpr int kBlockDimX = 16;
constexpr int kBlockDimY = 8;
constexpr int kWarmupRuns = 3;
constexpr int kTimedRuns = 10;
constexpr float kRelTolerance = 1e-5f;
static_assert((kBlockDimX * kBlockDimY) % 32 == 0,
"block must be a whole number of warps");
// snippet: buggy-kernel
// out = 2 * in, one thread per pixel. The load index follows the course
// formula. The store index is the same formula retyped from memory, and
// retyped wrong.
__global__ void scaleBuggy(const float* in, float* out, size_t width,
size_t height) {
const size_t col =
blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
const size_t row =
blockIdx.y * static_cast<size_t>(blockDim.y) + threadIdx.y;
if (col < width && row < height) {
const size_t src = row * width + col;
// DELIBERATE BUG (day 63): transposed store index, the day 7
// classic. Every value it produces is still inside the buffer,
// so nothing faults; the value just lands in the wrong pixel.
// The fix is dst = src, as scaleFixed writes it.
const size_t dst = col * height + row;
out[dst] = 2.0f * in[src];
}
}
// end snippet
// snippet: fixed-kernel
// The same kernel with the index computed once and used on both sides,
// which is the cheapest way to make the two sides unable to disagree.
__global__ void scaleFixed(const float* in, float* out, size_t width,
size_t height) {
const size_t col =
blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
const size_t row =
blockIdx.y * static_cast<size_t>(blockDim.y) + threadIdx.y;
if (col < width && row < height) {
const size_t i = row * width + col;
out[i] = 2.0f * in[i];
}
}
// end snippet
// CPU reference. Plain loops, obviously correct, allocates nothing.
static void scaleCpu(const float* in, float* out, size_t width, size_t height) {
for (size_t row = 0; row < height; ++row) {
for (size_t col = 0; col < width; ++col) {
const size_t i = row * width + col;
out[i] = 2.0f * in[i];
}
}
}
// Returns the first index where got and want differ by more than the relative
// tolerance, or n if they agree everywhere.
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;
}
// How many elements the buggy kernel should get wrong. The transposed
// index agrees with the correct one only where col * height + row equals
// row * width + col, which rearranges to 36 * col == 60 * row, that is
// 3 * col == 5 * row. Counting those pairs here, instead of hard-coding
// 13, keeps the check honest if someone changes the image size.
// snippet: expected-mismatches
static size_t expectedBuggyMismatches() {
size_t fixedPoints = 0;
for (size_t row = 0; row < kHeight; ++row) {
for (size_t col = 0; col < kWidth; ++col) {
if (col * kHeight + row == row * kWidth + col) {
++fixedPoints;
}
}
}
return kElems - fixedPoints;
}
// end snippet
// Times a launch with CUDA events and returns the mean milliseconds per run.
//
// This is the one template and the one lambda allowed in module 1 to 3 code.
// Copy it verbatim; the alternative is six copies of the event boilerplate,
// which is how a warm-up goes missing from one of them.
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;
}
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);
#ifdef __CUDACC_DEBUG__
std::printf("build: device debug (-G)\n");
#else
std::printf("build: release (-O3)\n");
#endif
const dim3 block(kBlockDimX, kBlockDimY);
const dim3 blocks(
static_cast<unsigned int>((kWidth + kBlockDimX - 1) / kBlockDimX),
static_cast<unsigned int>((kHeight + kBlockDimY - 1) / kBlockDimY));
std::printf(
"image %zu x %zu, %zu pixels, %d x %d blocks in a %u x %u "
"grid\n",
kWidth, kHeight, kElems, kBlockDimX, kBlockDimY, blocks.x, blocks.y);
// in[i] = i, so every pixel holds its own row-major index and a value
// that lands in the wrong slot names the slot it came from.
const size_t bytes = kElems * sizeof(float);
std::vector<float> h_in(kElems);
std::vector<float> h_out(kElems);
std::vector<float> h_want(kElems);
for (size_t i = 0; i < kElems; ++i) {
h_in[i] = static_cast<float>(i);
}
scaleCpu(h_in.data(), h_want.data(), kWidth, kHeight);
float* d_in = nullptr;
float* 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));
// The buggy kernel first, because it is the one commands.gdb breaks in.
// Its gate is inverted: it must be wrong, and wrong by the exact count
// the transposition predicts, or the supplied bug is not the bug the
// lesson describes.
scaleBuggy<<<blocks, block>>>(d_in, d_out, kWidth, kHeight);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
size_t buggyMismatches = 0;
for (size_t i = 0; i < kElems; ++i) {
const float scale = (h_want[i] == 0.0f) ? 1.0f : std::fabs(h_want[i]);
if (std::fabs(h_out[i] - h_want[i]) > kRelTolerance * scale) {
++buggyMismatches;
}
}
const size_t expected = expectedBuggyMismatches();
std::printf(
"scaleBuggy: %zu of %zu pixels wrong (transposition "
"predicts %zu)\n",
buggyMismatches, kElems, expected);
// The fixed kernel overwrites the same output buffer.
scaleFixed<<<blocks, block>>>(d_in, d_out, kWidth, kHeight);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
const size_t bad =
firstMismatch(h_out.data(), h_want.data(), kElems, kRelTolerance);
// Timed so the two builds of this same file put a number on -G. The
// kernel is far under 50 us, so the mean over kTimedRuns launches is
// the launch cost plus the kernel; the ratio between builds is still
// the point, since everything but the flag is held equal.
const float fixedMs = timeKernel(
[&] { scaleFixed<<<blocks, block>>>(d_in, d_out, kWidth, kHeight); });
std::printf(
"scaleFixed: %.4f ms per launch, mean of %d after %d "
"warm-ups\n",
fixedMs, kTimedRuns, kWarmupRuns);
// Free before reporting, so the failure paths free too.
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_out));
int status = EXIT_SUCCESS;
if (buggyMismatches != expected) {
std::fprintf(stderr,
"scaleBuggy missed on %zu pixels, expected %zu: the "
"supplied bug is not behaving as documented\n",
buggyMismatches, expected);
status = EXIT_FAILURE;
}
if (bad != kElems) {
std::fprintf(stderr, "scaleFixed wrong at %zu: got %.9g, want %.9g\n",
bad, h_out[bad], h_want[bad]);
status = EXIT_FAILURE;
}
if (status == EXIT_SUCCESS) {
std::printf(
"scaleFixed matches the CPU reference on all %zu "
"pixels\n",
kElems);
}
return status;
}