code/day06-error-checking/error_checking.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 6: what a CUDA call is trying to tell you, and when.
//
// This program makes two mistakes on purpose and reports which call saw each
// one. Nothing here is timed. The file is almost nothing but synchronisation
// points, and a sync changes what a timed region measures, so day 9 starts
// the timing and this day does not.
//
// Four parts, in this order:
// 1. a call that fails, read with cudaPeekAtLastError and cudaGetLastError
// 2. the same context afterwards, still working, because that error is not
// sticky
// 3. an out-of-bounds write inside a kernel, which surfaces at the next
// call that looks rather than at the launch
// 4. what clearing a sticky error does, which is nothing
//
// Part 3 destroys the context on purpose, so it runs last and nothing after
// it can succeed. That is the demonstration, not a bug.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o error_checking error_checking.cu
// Run: ./error_checking
//
// 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, byte for byte the copy day 5 handed over. This file is
// standalone, the way a Compiler Explorer buffer is, so it carries its own
// copy rather than including one. `err_` carries a trailing underscore so it
// cannot collide with a variable at the call site, the do/while makes the
// macro one statement so it survives a braceless `if`, and `#call` turns the
// call into its own source text so the message names what failed.
// snippet: check-macro
#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)
// end snippet
// 611 is day 5's size, kept so the bounds check runs on every launch: at 256
// threads per block that is 3 blocks, 768 threads and 157 with no element.
// kRelTolerance is not zero because a GPU may fuse a multiply and an add
// where the host compiler does not; on these inputs both sides are exact.
constexpr size_t kElems = 611;
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr float kRelTolerance = 1e-5f;
// One gibibyte past the end of a 2,444 byte buffer. The overshoot is this
// large deliberately: cudaMalloc rounds an allocation up and the pages after
// a small buffer are often already mapped, so a short overrun can write into
// memory the driver owns and never fault at all. A gibibyte cannot.
constexpr size_t kOvershoot = 1ull << 28;
static_assert(kThreadsPerBlock % 32 == 0,
"block size must be a whole number of warps");
static_assert(kOvershoot * sizeof(float) >= (1ull << 30),
"the overshoot has to be far larger than the slack cudaMalloc "
"rounds an allocation up to, or the write lands in mapped "
"memory, nothing faults, and parts 3 and 4 prove nothing");
// out[i] = in[i] + 1. One thread owns one element.
//
// Memory: consecutive threads take consecutive elements, so one warp's 32
// addresses cover 128 contiguous bytes.
//
// Launch assumption: gridDim.x * blockDim.x >= n. Part 2 uses this kernel to
// show the context still works, so it has to be correct.
__global__ void addOne(const float* in, float* out, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
out[i] = in[i] + 1.0f;
}
}
// Writes one float `overshoot` elements past the buffer it was handed. Every
// live thread does it, so the fault does not depend on which thread the
// scheduler runs first.
//
// Memory: the warp's 32 addresses are still contiguous with each other, and
// every one of them is outside the allocation.
//
// Launch assumption: none. This kernel is a fault generator and it is wrong
// at every size. Its writes are never read back.
__global__ void writePastEnd(float* out, size_t n, size_t overshoot) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
out[i + overshoot] = 1.0f;
}
}
// CPU reference. Written for obvious correctness, not speed: plain loop, no
// intrinsics. It never allocates; the caller owns every buffer.
static void addOneCpu(const float* in, float* out, size_t n) {
for (size_t i = 0; i < n; ++i) {
out[i] = in[i] + 1.0f;
}
}
// 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 512" names the block, "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;
}
// Prints one status the way a reader searches for it: the enum name from
// cudaGetErrorName beside the message from cudaGetErrorString. The macro
// above prints only the message, because the message is what people paste
// into a search box; the name is what the runtime API reference is indexed by.
static void report(const char* label, cudaError_t err) {
std::printf(" %-32s %-30s %s\n", label, cudaGetErrorName(err),
cudaGetErrorString(err));
}
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\n", prop.name, prop.major,
prop.minor);
const size_t bytes = kElems * sizeof(float);
const int blocks =
static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
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);
}
float* d_in = nullptr;
float* d_out = nullptr;
CUDA_CHECK(cudaMalloc(&d_in, bytes));
CUDA_CHECK(cudaMalloc(&d_out, bytes));
// Part 1. A copy whose direction argument disagrees with its pointers.
// The runtime documents this case as undefined behaviour, so the code it
// returns is not promised. What the part demonstrates is the distance:
// whatever comes back, it comes back from this call and not a later one.
std::printf("part 1: one failed call, and who can still see it\n");
// snippet: peek-versus-get
// A null destination. This fails synchronously with
// cudaErrorInvalidValue, the runtime rejects it before anything is
// queued, and the context is untouched, which is what part 2 needs.
//
// The obvious choice, passing the wrong cudaMemcpyKind, does not work.
// Measured on a Tesla T4 with CUDA 12.6: a host-to-device copy tagged
// cudaMemcpyDeviceToHost returns cudaSuccess and copies correctly,
// because unified addressing lets the runtime read the real direction
// off the pointers and the kind argument is a hint it can overrule.
// That is worth knowing on its own: the direction flag will not catch
// your mistake for you.
const cudaError_t badCopy =
cudaMemcpy(nullptr, h_in.data(), bytes, cudaMemcpyHostToDevice);
const cudaError_t peek1 = cudaPeekAtLastError();
const cudaError_t peek2 = cudaPeekAtLastError();
const cudaError_t taken = cudaGetLastError();
const cudaError_t peek3 = cudaPeekAtLastError();
// end snippet
report("cudaMemcpy returned", badCopy);
report("cudaPeekAtLastError", peek1);
report("cudaPeekAtLastError, again", peek2);
report("cudaGetLastError", taken);
report("cudaPeekAtLastError, after get", peek3);
// Failures below set `wrong` and fall through to one cleanup block at the
// end. An early `return` here would skip the cudaFree calls, which on the
// page that teaches error handling would be its own small joke.
int wrong = 0;
if (badCopy == cudaSuccess) {
std::fprintf(stderr,
"the copy to a null destination succeeded, so part 1 "
"shows nothing on this device\n");
++wrong;
}
if (peek1 != badCopy || peek2 != badCopy || taken != badCopy) {
std::fprintf(stderr,
"the stored status did not match what the call "
"returned\n");
++wrong;
}
if (peek3 != cudaSuccess) {
std::fprintf(stderr, "cudaGetLastError did not clear the status\n");
++wrong;
}
// Part 2. Same context, same buffers, one correct copy and one correct
// launch. It works, which is what non-sticky means: the failed call above
// damaged nothing.
std::printf("\npart 2: the same context after a non-sticky error\n");
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));
addOne<<<blocks, kThreadsPerBlock>>>(d_in, d_out, kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
addOneCpu(h_in.data(), h_want.data(), kElems);
const size_t bad =
firstMismatch(h_out.data(), h_want.data(), kElems, kRelTolerance);
if (bad != kElems) {
std::fprintf(stderr, "addOne wrong at %zu: got %.9g, want %.9g\n", bad,
h_out[bad], h_want[bad]);
++wrong;
}
std::printf(" all %zu elements match, so the context survived it\n",
kElems);
// Part 3. An out-of-bounds write inside a kernel. The launch is legal, so
// the launch is accepted; the fault happens later, on the device, and the
// host hears about it from whichever runtime call looks next.
//
// The two lines day 5 puts after every launch are deliberately missing
// here. This is what their absence costs.
std::printf("\npart 3: an error that arrives late\n");
// snippet: late-error
writePastEnd<<<blocks, kThreadsPerBlock>>>(d_out, kElems, kOvershoot);
const cudaError_t atLaunch = cudaPeekAtLastError();
report("peek, straight after the launch", atLaunch);
// Nothing is wrong with this copy. It reads a buffer that exists, into a
// vector that exists, with the direction its pointers agree on. It is
// just the next runtime call, which is the whole reason a CUDA error can
// name a line that did nothing.
const cudaError_t atNextCall =
cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost);
// end snippet
report("the next cudaMemcpy returned", atNextCall);
if (atNextCall == cudaSuccess) {
std::fprintf(stderr,
"the out-of-bounds write did not fault here, so parts 3 "
"and 4 show nothing. Re-run under compute-sanitizer "
"--tool memcheck\n");
++wrong;
}
// Part 4. Sticky. cudaGetLastError clears the runtime's stored status and
// the next call sets it straight back, because the context itself is the
// thing that is broken. No call in this process can recover it.
std::printf("\npart 4: what clearing a sticky error buys you\n");
const cudaError_t cleared = cudaGetLastError();
const cudaError_t afterClear = cudaPeekAtLastError();
float* d_probe = nullptr;
const cudaError_t freshCall = cudaMalloc(&d_probe, sizeof(float));
report("cudaGetLastError returned", cleared);
report("peek, right after clearing it", afterClear);
report("a fresh cudaMalloc returned", freshCall);
if (freshCall == cudaSuccess) {
std::fprintf(stderr,
"an unrelated allocation succeeded after the fault, so "
"the error was not sticky on this device\n");
++wrong;
}
// Every cudaMalloc in this file has a cudaFree, and after part 3 not one
// of them can work. They are called anyway and their status printed,
// because "even cudaFree returns the same code" is the plainest statement
// of what sticky means. CUDA_CHECK is deliberately not used on these two:
// it would exit with a failure code on a program that is behaving exactly
// as designed. d_probe is still null, so there is nothing to free.
report("cudaFree(d_in) returned", cudaFree(d_in));
report("cudaFree(d_out) returned", cudaFree(d_out));
if (d_probe != nullptr) {
// Only reachable when part 4 did not behave as described, which is
// already recorded in `wrong`. Freeing it keeps the file's claim that
// every allocation here has a matching free literally true.
report("cudaFree(d_probe) returned", cudaFree(d_probe));
}
std::printf("\nthe context is gone. Only a new process gets CUDA back.\n");
// One exit, after the frees. Every check above records into `wrong` and
// falls through to here, so a failing run frees exactly what a passing run
// frees. On the page that teaches error handling, leaking on the error
// path would be its own small joke.
if (wrong > 0) {
std::fprintf(stderr,
"\n%d check(s) failed: this device did not behave the way "
"the lesson describes\n",
wrong);
return EXIT_FAILURE;
}
return EXIT_SUCCESS;
}