code/day64-printf/printf_assert.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 64: printf, assert and when printf lies.
//
// One source file, four builds:
//
// default the full demonstration: the printf FIFO shrunk
// as far as the driver allows, a buffering demo,
// an overflow demo, then a device-side assert
// that kills the context
// -DDAY64_LOSE_OUTPUT=1 the trap: the last kernel's lines are dropped
// because the program exits without a sync
// -DNDEBUG the assert is deleted by the preprocessor and
// the poisoned element sails through, which is why
// no gate in this course is an assert
// -DDAY64_POISON=0 the fixed variant: same kernel, clean data,
// every status below is cudaSuccess
//
// What it demonstrates: device printf writes argument records into a
// fixed-size circular buffer that is flushed at sync points, not into a
// terminal. Output can be overwritten (a full FIFO drops the oldest
// records), reordered (the scheduler owns the order) or lost outright
// (exit without a flush point). And a fired device assert aborts the
// context, after which every CUDA call in the process reports the assert.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o printf_assert printf_assert.cu
// Run: ./printf_assert
//
#include <cassert>
#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.
#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)
// Build switches, set from the nvcc line. A compile-time switch that CI
// flips with -D cannot be a constexpr, so these two defaults are the only
// #define values in the file.
#ifndef DAY64_LOSE_OUTPUT
#define DAY64_LOSE_OUTPUT 0
#endif
#ifndef DAY64_POISON
#define DAY64_POISON 1
#endif
// The smallest FIFO worth asking for. The driver is free to round the
// request up to its own granularity and does: on the verification node
// this 4 KiB ask comes back as 262,144 bytes. Part 2 therefore derives
// its line count from what was granted. The FIFO stores each record's
// arguments and metadata rather than the formatted text, so how many
// lines survive is a driver property; that the survivors are the newest
// is documented and is the point.
constexpr size_t kFifoBytes = 4096;
constexpr int kProbeThreads = 32; // one warp
constexpr int kAssertThreads = 32;
constexpr size_t kPoisonedIndex = 5;
static_assert(kProbeThreads % 32 == 0,
"the buffering demo reasons about whole warps");
// Each thread prints one line naming itself. No memory is touched: the
// record goes into the device-side printf FIFO and sits there until a
// flush point, which is the whole demonstration.
//
// The 32 lines of one warp arrive in whatever order the runtime drains
// them; day 1 measured the same non-promise across blocks.
//
// Launch assumption: one block of kProbeThreads threads.
__global__ void printPerThread(void) {
printf("device: thread %2u has printed\n", threadIdx.x);
}
// One thread prints `lines` numbered lines. A single thread's loop is the
// one production order the scheduler cannot change, so when the circular
// FIFO wraps, "the oldest records are overwritten" becomes checkable: the
// survivors must be one contiguous run ending at the last line.
//
// Launch assumption: <<<1, 1>>>. More threads would scramble the
// production order the overflow argument depends on.
// snippet: overflow-kernel
__global__ void printPastFifo(int lines) {
for (int k = 0; k < lines; ++k) {
printf("fifo line %04d of %d\n", k, lines);
}
}
// end snippet
// Asserts that in[i] holds the value the host promised to put there. The
// host poisons one element on purpose, so exactly one thread's assert
// fires, and the message names the block and thread that failed. That
// message is the tool: printf shows what you asked it to show, assert
// points at the thread where the data went wrong, then takes the whole
// context with it.
//
// Memory: consecutive threads read consecutive floats, one coalesced
// 128-byte load for the warp. Nothing is written.
//
// Launch assumption: gridDim.x * blockDim.x >= n.
// snippet: assert-kernel
__global__ void assertEachElement(const float* __restrict__ in, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
const float want = static_cast<float>(i);
// NDEBUG deletes the next line; the -DNDEBUG build proves it.
assert(in[i] == want); // deliberate device-side assert (day 64)
}
}
// end snippet
static void report(const char* what, cudaError_t status) {
std::printf(" %-30s %-24s %s\n", what, cudaGetErrorName(status),
cudaGetErrorString(status));
}
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);
// Part 0: the FIFO is a size you can read and set. Shrinking it is
// what makes the overflow in part 2 reproducible. The set has to
// happen before any kernel that calls printf has launched; after one,
// cudaDeviceSetLimit returns cudaErrorInvalidValue.
// snippet: fifo-shrink
size_t fifoDefault = 0;
CUDA_CHECK(cudaDeviceGetLimit(&fifoDefault, cudaLimitPrintfFifoSize));
std::printf("printf FIFO default: %zu bytes\n", fifoDefault);
CUDA_CHECK(cudaDeviceSetLimit(cudaLimitPrintfFifoSize, kFifoBytes));
size_t fifoNow = 0;
CUDA_CHECK(cudaDeviceGetLimit(&fifoNow, cudaLimitPrintfFifoSize));
std::printf("printf FIFO now: %zu bytes\n", fifoNow);
// end snippet
// The driver rounds this limit up to its own granularity, so asking
// for 4 KiB does not mean getting 4 KiB. Measured on a Tesla T4 with
// driver 595.84: the request above lands at 262,144 bytes. The gate
// is therefore "the driver honoured a request at all", not equality,
// and part 2 sizes itself from what was actually granted.
if (fifoNow == 0 || fifoNow > fifoDefault) {
std::fprintf(stderr, "FIFO request refused: %zu bytes\n", fifoNow);
return EXIT_FAILURE;
}
// Part 1: the lines exist before the sync, in the buffer. The fflush
// calls pin the host markers in place so the transcript shows where
// the device lines land even when stdout is a file.
std::printf("part 1: %d device lines are about to be buffered\n",
kProbeThreads);
printPerThread<<<1, kProbeThreads>>>();
CUDA_CHECK(cudaGetLastError());
std::printf("host: launch returned, no sync yet\n");
std::fflush(stdout);
CUDA_CHECK(cudaDeviceSynchronize());
std::fflush(stdout);
std::printf("host: sync returned, the buffer is flushed\n");
// Part 2: overflow whatever the driver granted. A record is the
// format pointer plus the arguments plus metadata, always more than
// eight bytes, so one line per eight granted bytes is guaranteed to
// wrap the buffer. Count the survivors with `grep -c '^fifo line'`
// on the transcript; the first surviving number is how far the wrap
// reached back.
const int fifoLines = static_cast<int>(fifoNow / 8);
std::printf("part 2: one thread prints %d lines into %zu bytes\n",
fifoLines, fifoNow);
std::fflush(stdout);
printPastFifo<<<1, 1>>>(fifoLines);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
std::fflush(stdout);
std::printf("host: overflow demo flushed\n");
#if DAY64_LOSE_OUTPUT
// snippet: lost-output
// deliberate bug (day 64): launch a printing kernel, then exit with
// no sync. The 12.6 guide, section 7.35.2: "the buffer is not flushed
// automatically when the program exits." These 32 lines are dropped.
// The fix is the cudaDeviceSynchronize() part 1 runs on this same
// kernel: this transcript holds 32 `device: thread` lines, not 64.
std::printf("part 3: launching, then exiting with no sync\n");
std::fflush(stdout);
printPerThread<<<1, kProbeThreads>>>();
// deliberately unchecked (day 64): even cudaGetLastError() is not a
// flush point, but the point lands harder with nothing at all here.
return EXIT_SUCCESS;
// end snippet
#else
// Part 3: one poisoned element, one assert. Fill 32 floats with their
// own indices, then overwrite one on purpose.
const size_t n = static_cast<size_t>(kAssertThreads);
std::vector<float> h_in(n);
for (size_t i = 0; i < n; ++i) {
h_in[i] = static_cast<float>(i);
}
#if DAY64_POISON
// deliberate bug (day 64): element 5 holds the wrong value, so thread
// [5,0,0]'s assert fires and the message names the culprit. Build
// with -DDAY64_POISON=0 for the fixed variant.
h_in[kPoisonedIndex] = -1.0f;
#endif
float* d_in = nullptr;
CUDA_CHECK(cudaMalloc(&d_in, n * sizeof(float)));
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), n * sizeof(float),
cudaMemcpyHostToDevice));
// What this build expects from every call after the launch. Poisoned
// and assert compiled in: cudaErrorAssert everywhere. Fixed data or
// -DNDEBUG: cudaSuccess everywhere.
#if DAY64_POISON && !defined(NDEBUG)
const cudaError_t want = cudaErrorAssert;
#else
const cudaError_t want = cudaSuccess;
#endif
#ifdef NDEBUG
const int ndebug = 1;
#else
const int ndebug = 0;
#endif
std::printf("part 3: assert kernel (poison=%d, ndebug=%d, expect %s)\n",
DAY64_POISON, ndebug, cudaGetErrorName(want));
std::fflush(stdout);
assertEachElement<<<1, kAssertThreads>>>(d_in, n);
// Statuses from here down are captured and gated by hand rather than
// wrapped in CUDA_CHECK, because on the poisoned build the expected
// value is not cudaSuccess and the macro would exit on the exact
// result this program exists to show.
const cudaError_t atLaunch = cudaGetLastError();
const cudaError_t atSync = cudaDeviceSynchronize();
std::fflush(stdout);
report("cudaGetLastError at launch", atLaunch);
report("cudaDeviceSynchronize", atSync);
int status = EXIT_SUCCESS;
if (atLaunch != cudaSuccess) {
std::fprintf(stderr, "the launch itself was refused: %s\n",
cudaGetErrorName(atLaunch));
status = EXIT_FAILURE;
}
if (atSync != want) {
std::fprintf(stderr, "sync returned %s, this build expected %s\n",
cudaGetErrorName(atSync), cudaGetErrorName(want));
status = EXIT_FAILURE;
}
// After a fired assert the context is in day 6's sticky class, so the
// frees below are attempted on every path and whether they can still
// do their job is itself the result. On the fixed and NDEBUG builds
// they must succeed.
float* d_probe = nullptr;
const cudaError_t atMalloc = cudaMalloc(&d_probe, sizeof(float));
report("a fresh cudaMalloc", atMalloc);
cudaError_t atProbeFree = atMalloc;
if (atMalloc == cudaSuccess) {
atProbeFree = cudaFree(d_probe);
}
const cudaError_t atFree = cudaFree(d_in);
report("cudaFree(d_in)", atFree);
if (atMalloc != want || atProbeFree != want || atFree != want) {
std::fprintf(stderr, "post-assert statuses disagree with %s\n",
cudaGetErrorName(want));
status = EXIT_FAILURE;
}
// Reported, never gated: NVIDIA's two documents disagree about what
// comes next. The 12.6 guide's assertion section says no more
// commands reach the device "until cudaDeviceReset() is called to
// reinitialize the device"; the runtime reference's cudaErrorAssert
// entry says "the process must be terminated and relaunched". This
// run is the tiebreaker for this card. cudaDeviceReset() is banned in
// lesson code as cleanup; this call is the subject, not cleanup.
if (want == cudaErrorAssert) {
const cudaError_t atReset = cudaDeviceReset(); // deliberate, day 64
report("cudaDeviceReset", atReset);
float* d_after = nullptr;
const cudaError_t afterReset = cudaMalloc(&d_after, sizeof(float));
report("cudaMalloc after reset", afterReset);
if (afterReset == cudaSuccess) {
const cudaError_t freeAfter = cudaFree(d_after);
report("cudaFree after reset", freeAfter);
std::printf(
"reset revived the process: the guide's sentence "
"held\n");
} else {
std::printf(
"reset did not revive it: the reference's sentence "
"held\n");
}
}
if (status == EXIT_SUCCESS) {
std::printf("every status matched what this build expected\n");
}
return status;
#endif
}