code/day65-deadlocks/deadlocks.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 65: deadlocks and hangs. Three kernels that never finish, run under a
// watchdog so this program always does.
//
// Case 1: a __syncthreads() only half the block can reach, because the other
// half spins on a flag that is written after the barrier.
// Case 2: a homemade grid-wide barrier launched with one block more than the
// card can hold resident. The resident blocks spin for an arrival
// that needs an SM slot they will never give up.
// Case 3: a loop whose exit test is `x != 1.0f` while x steps by 0.1f, which
// lands above 1.0f without ever equaling it.
//
// Every case runs twice, fixed then buggy, each in its own forked child
// process. The parent makes no CUDA call anywhere: a process that has
// initialized a CUDA context must not fork children that use the GPU, and a
// stuck kernel cannot be cancelled from its own process by any smaller tool
// than tearing the whole context down. SIGKILL on a child plus process exit
// is what reliably returns the device.
//
// The three buggy kernels are DELIBERATE BUGS, labelled at the line that
// carries each one, with the fixed variant beside every bug in this same
// file so the harness can diff behaviour rather than describe it.
//
// Nothing here is timed. The subject is programs for which "how long" has
// no answer.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o deadlocks deadlocks.cu
// Run: ./deadlocks (the full watchdogged table)
// DAY65_CASE=1 ./deadlocks (hang on purpose, in the foreground, for
// nvidia-smi and cuda-gdb; cases 2 and 3
// likewise; kill it yourself)
//
// Linux only: the watchdog is fork(), waitpid() and SIGKILL.
//
#include <cerrno>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <ctime>
#include <vector>
#include <signal.h>
#include <sys/types.h>
#include <sys/wait.h>
#include <unistd.h>
#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)
constexpr int kThreadsPerBlock = 256; // 8 warps
// Producers in the hand-off case; the other 128 threads consume. A multiple
// of 32, so the split in case 1 lands on a warp boundary and no warp has
// lanes on both sides of it.
constexpr unsigned int kHalf = 128;
constexpr size_t kLoopElems = 65536;
// Additions of float 0.1f before the running total passes 1.0f. Float 0.1f
// is slightly above one tenth, so the tenth addition lands above 1.0f.
constexpr unsigned int kExpectedSteps = 10;
// Long enough that a completing case is never killed by accident (each case
// does a few small copies and one launch), short enough that six runs stay
// under a minute even when three of them spin flat out until killed.
constexpr int kWatchdogSeconds = 5;
static_assert(kThreadsPerBlock % 32 == 0,
"block size must be a whole number of warps");
static_assert(kHalf * 2 == kThreadsPerBlock,
"the hand-off splits one block into producers and consumers");
// DELIBERATE BUG (day 65, case 1): the barrier sits inside `tid < kHalf`, so
// at most 128 of the block's 256 threads can ever arrive, while the other
// 128 spin on `flag`, which is written two lines below the barrier. A
// circular wait: nobody moves, and the launch never returns.
//
// One producer thread copies one float into the tile; one consumer thread is
// meant to copy it back out. A warp's 32 addresses are 32 consecutive floats
// on both sides. Launch assumption: one block of exactly kThreadsPerBlock.
// snippet: handoff-blocked
__global__ void handoffBlocked(const float* __restrict__ in,
float* __restrict__ out,
unsigned int* __restrict__ flag) {
__shared__ float tile[kHalf];
const unsigned int tid = threadIdx.x;
if (tid < kHalf) {
tile[tid] = in[tid];
__syncthreads(); // deliberate bug: 128 arrivals, 256 expected
if (tid == 0) {
atomicExch(flag, 1u); // the consumers are waiting on this line
}
} else {
while (atomicAdd(flag, 0u) == 0u) {
}
out[tid - kHalf] = tile[tid - kHalf];
}
}
// end snippet
// The fix: one barrier that every thread reaches. The producers' stores are
// then ordered before the consumers' loads by the barrier itself, and no
// flag is needed at all. Same memory pattern and launch assumption as above.
__global__ void handoffOrdered(const float* __restrict__ in,
float* __restrict__ out) {
__shared__ float tile[kHalf];
const unsigned int tid = threadIdx.x;
if (tid < kHalf) {
tile[tid] = in[tid];
}
__syncthreads();
if (tid >= kHalf) {
out[tid - kHalf] = tile[tid - kHalf];
}
}
// A homemade grid-wide barrier: thread 0 of each block counts its block in
// with a device-scope atomic (day 27's ticket, minus the fence, because
// nothing is published here except the count itself), then every thread
// spins until all gridDim.x blocks have arrived.
//
// This is correct only while every block in the grid is resident at once.
// Blocks are never preempted, so a spinning block holds its SM slot until it
// exits, and a block the scheduler cannot place waits forever behind it. The
// launch size decides which program this is, and main() owns that choice.
//
// One thread reads and writes single counters; there is no per-element
// memory pattern. Launch assumption: gridDim.x is at most the card's
// resident-block capacity at this block size, or this kernel never returns.
__global__ void gridlockBarrier(unsigned int* __restrict__ arrived,
unsigned int* __restrict__ seen) {
if (threadIdx.x == 0) {
atomicAdd(arrived, 1u);
}
// volatile, so the spin re-reads memory instead of a stale register.
volatile unsigned int* watch = arrived;
while (*watch < gridDim.x) {
}
if (threadIdx.x == 0) {
seen[blockIdx.x] = *watch;
}
}
// Steps a running total upward by in[i] until it reaches 1.0f and records
// how many additions that took. in[] holds 0.1f everywhere, so the fixed
// kernel should report kExpectedSteps for every element.
//
// One thread owns one element; a warp reads 32 consecutive floats and
// writes 32 consecutive counts. Launch assumption: grid covers n.
//
// DELIBERATE BUG (day 65, case 3): the exit test is `!=`. Ten additions of
// float 0.1f land just above 1.0f, never on it. And once x grows so large
// that x + 0.1f rounds back to x, the loop stops making progress at all,
// so the condition is stuck at true forever.
__global__ void stepsToOneStuck(const float* __restrict__ in,
unsigned int* __restrict__ steps, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
float x = 0.0f;
unsigned int count = 0u;
// snippet: loop-guard
while (x != 1.0f) { // deliberate bug: the test must be <, not !=
x += in[i];
++count;
}
// end snippet
steps[i] = count;
}
}
// The fix: a convergence test that is an inequality. `x < 1.0f` goes false
// the moment the total reaches or passes 1.0f, whether or not any float on
// the way happens to equal it. Also records the final x, so the host can
// show that the total passed 1.0f without ever equaling it.
__global__ void stepsToOneConverging(const float* __restrict__ in,
unsigned int* __restrict__ steps,
float* __restrict__ finalX, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
float x = 0.0f;
unsigned int count = 0u;
while (x < 1.0f) {
x += in[i];
++count;
}
steps[i] = count;
finalX[i] = x;
}
}
// Everything below runs in a forked child (or in the foreground under
// DAY65_CASE), never in the watchdog parent.
static int probeDevice(bool buggy) {
static_cast<void>(buggy);
cudaDeviceProp prop;
int kernelExecTimeout = 0;
CUDA_CHECK(cudaSetDevice(0));
CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
CUDA_CHECK(cudaDeviceGetAttribute(&kernelExecTimeout,
cudaDevAttrKernelExecTimeout, 0));
// The execution-timeout attribute is the driver's own watchdog, present on
// cards that drive a display. Where it prints "off", nothing but this
// program's fork-and-kill parent will ever stop a stuck kernel.
std::printf("GPU: %s (compute capability %d.%d, %d SMs)\n", prop.name,
prop.major, prop.minor, prop.multiProcessorCount);
std::printf("driver run time limit on kernels: %s\n",
kernelExecTimeout != 0 ? "on" : "off");
std::printf("watchdog: %d s per case, SIGKILL on expiry\n\n",
kWatchdogSeconds);
return EXIT_SUCCESS;
}
// Case 1. The buggy launch never returns; the child hangs inside the
// cudaDeviceSynchronize() below and the parent kills it. The fixed launch
// must produce out[c] == in[c] for all 128 hand-off elements.
static int runCaseBarrier(bool buggy) {
std::vector<float> h_in(kHalf);
std::vector<float> h_out(kHalf, 0.0f);
for (unsigned int c = 0; c < kHalf; ++c) {
h_in[c] = static_cast<float>(c + 1);
}
const size_t bytes = kHalf * sizeof(float);
float* d_in = nullptr;
float* d_out = nullptr;
unsigned int* d_flag = nullptr;
CUDA_CHECK(cudaMalloc(&d_in, bytes));
CUDA_CHECK(cudaMalloc(&d_out, bytes));
CUDA_CHECK(cudaMalloc(&d_flag, sizeof(unsigned int)));
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemset(d_out, 0, bytes));
CUDA_CHECK(cudaMemset(d_flag, 0, sizeof(unsigned int)));
if (buggy) {
handoffBlocked<<<1, kThreadsPerBlock>>>(d_in, d_out, d_flag);
} else {
handoffOrdered<<<1, kThreadsPerBlock>>>(d_in, d_out);
}
CUDA_CHECK(cudaGetLastError());
// The launch above returned immediately either way; this line is where
// a hung kernel becomes a hung program.
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_out));
CUDA_CHECK(cudaFree(d_flag));
int status = EXIT_SUCCESS;
for (unsigned int c = 0; c < kHalf; ++c) {
if (h_out[c] != h_in[c]) {
std::fprintf(stderr, "hand-off wrong at %u: got %.9g, want %.9g\n",
c, h_out[c], h_in[c]);
status = EXIT_FAILURE;
break;
}
}
return status;
}
// Case 2. Same kernel both times; only the grid size changes. The capacity
// is asked from the occupancy API rather than assumed, so the same binary
// wedges any card by exactly one block.
static int runCaseGridlock(bool buggy) {
cudaDeviceProp prop;
CUDA_CHECK(cudaSetDevice(0));
CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
// snippet: gridlock-launch
int blocksPerSm = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocksPerSm, gridlockBarrier, kThreadsPerBlock, 0));
const int resident = blocksPerSm * prop.multiProcessorCount;
// DELIBERATE BUG (day 65, case 2): one block more than can be resident.
const int blocks = buggy ? resident + 1 : resident;
// end snippet
std::printf(" capacity %d blocks (%d per SM x %d SMs), launching %d\n",
resident, blocksPerSm, prop.multiProcessorCount, blocks);
unsigned int* d_arrived = nullptr;
unsigned int* d_seen = nullptr;
CUDA_CHECK(cudaMalloc(&d_arrived, sizeof(unsigned int)));
CUDA_CHECK(cudaMalloc(&d_seen, blocks * sizeof(unsigned int)));
CUDA_CHECK(cudaMemset(d_arrived, 0, sizeof(unsigned int)));
CUDA_CHECK(cudaMemset(d_seen, 0, blocks * sizeof(unsigned int)));
gridlockBarrier<<<blocks, kThreadsPerBlock>>>(d_arrived, d_seen);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
std::vector<unsigned int> h_seen(static_cast<size_t>(blocks));
CUDA_CHECK(cudaMemcpy(h_seen.data(), d_seen, blocks * sizeof(unsigned int),
cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaFree(d_arrived));
CUDA_CHECK(cudaFree(d_seen));
int status = EXIT_SUCCESS;
for (int b = 0; b < blocks; ++b) {
if (h_seen[static_cast<size_t>(b)] !=
static_cast<unsigned int>(blocks)) {
std::fprintf(stderr, "block %d saw %u arrivals, want %d\n", b,
h_seen[static_cast<size_t>(b)], blocks);
status = EXIT_FAILURE;
break;
}
}
return status;
}
// Case 3. The fixed kernel must report exactly kExpectedSteps additions for
// every element and a final total strictly above 1.0f: the total passes
// 1.0f without ever equaling it, which is the whole reason `!=` never ends.
static int runCaseLoop(bool buggy) {
std::vector<float> h_in(kLoopElems, 0.1f);
std::vector<unsigned int> h_steps(kLoopElems, 0u);
std::vector<float> h_final(kLoopElems, 0.0f);
const size_t inBytes = kLoopElems * sizeof(float);
const size_t stepBytes = kLoopElems * sizeof(unsigned int);
const int blocks = static_cast<int>((kLoopElems + kThreadsPerBlock - 1) /
kThreadsPerBlock);
float* d_in = nullptr;
unsigned int* d_steps = nullptr;
float* d_final = nullptr;
CUDA_CHECK(cudaMalloc(&d_in, inBytes));
CUDA_CHECK(cudaMalloc(&d_steps, stepBytes));
CUDA_CHECK(cudaMalloc(&d_final, inBytes));
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), inBytes, cudaMemcpyHostToDevice));
if (buggy) {
stepsToOneStuck<<<blocks, kThreadsPerBlock>>>(d_in, d_steps,
kLoopElems);
} else {
stepsToOneConverging<<<blocks, kThreadsPerBlock>>>(d_in, d_steps,
d_final, kLoopElems);
}
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_steps.data(), d_steps, stepBytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(
cudaMemcpy(h_final.data(), d_final, inBytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_steps));
CUDA_CHECK(cudaFree(d_final));
int status = EXIT_SUCCESS;
for (size_t i = 0; i < kLoopElems; ++i) {
if (h_steps[i] != kExpectedSteps || h_final[i] <= 1.0f) {
std::fprintf(stderr,
"element %zu: %u steps (want %u), final %.9g\n", i,
h_steps[i], kExpectedSteps, h_final[i]);
status = EXIT_FAILURE;
break;
}
}
if (status == EXIT_SUCCESS) {
std::printf(" every total passed 1.0f in %u steps, final %.9g\n",
kExpectedSteps, h_final[0]);
}
return status;
}
// Runs one case in a forked child and waits at most kWatchdogSeconds.
// Returns the child's exit code, or -1 when the watchdog killed it. The
// fork happens before the child's first CUDA call, so every child builds
// its own context, and SIGKILL plus process teardown hands whatever the
// child held back to the driver.
// snippet: watchdog
static int runWatchdogged(int (*fn)(bool), bool buggy) {
std::fflush(stdout);
std::fflush(stderr);
const pid_t pid = fork();
if (pid < 0) {
std::fprintf(stderr, "fork failed: %s\n", std::strerror(errno));
std::exit(EXIT_FAILURE);
}
if (pid == 0) {
// The child; its CUDA context starts here. Flush before _exit:
// _exit skips stdio teardown, and with stdout redirected to a
// file (block-buffered) an unflushed report never reaches disk.
const int rc = fn(buggy);
std::fflush(stdout);
std::fflush(stderr);
_exit(rc);
}
const int ticksPerSecond = 10;
for (int t = 0; t < kWatchdogSeconds * ticksPerSecond; ++t) {
int status = 0;
if (waitpid(pid, &status, WNOHANG) == pid) {
return WIFEXITED(status) ? WEXITSTATUS(status) : -1;
}
struct timespec tick = {0, 1000000000 / ticksPerSecond};
nanosleep(&tick, nullptr);
}
kill(pid, SIGKILL); // a stuck kernel dies with its process, not before
int status = 0;
waitpid(pid, &status, 0);
return -1;
}
// end snippet
struct Row {
const char* label;
int (*fn)(bool);
bool buggy;
const char* diagnosis;
};
int main() {
const Row rows[] = {
{"case 1 fixed uniform barrier", runCaseBarrier, false, ""},
{"case 1 buggy barrier in an if", runCaseBarrier, true,
"128 threads wait at a barrier for 128 threads that wait on a\n"
" flag written two lines below that barrier"},
{"case 2 fixed grid fits residency", runCaseGridlock, false, ""},
{"case 2 buggy grid exceeds it by 1", runCaseGridlock, true,
"every resident block spins for an arrival that needs the one\n"
" block the scheduler has nowhere to place"},
{"case 3 fixed loop tests x < 1.0f", runCaseLoop, false, ""},
{"case 3 buggy loop tests x != 1.0f", runCaseLoop, true,
"the total steps over 1.0f without ever equaling it, so the\n"
" exit condition never goes false"},
};
const int rowCount = static_cast<int>(sizeof(rows) / sizeof(rows[0]));
// DAY65_CASE=1|2|3 runs that buggy case in the foreground, with no fork
// and no watchdog, so nvidia-smi and cuda-gdb have a hang to look at.
const char* env = std::getenv("DAY65_CASE");
if (env != nullptr) {
int (*fn)(bool) = nullptr;
if (std::strcmp(env, "1") == 0) {
fn = runCaseBarrier;
} else if (std::strcmp(env, "2") == 0) {
fn = runCaseGridlock;
} else if (std::strcmp(env, "3") == 0) {
fn = runCaseLoop;
} else {
std::fprintf(stderr, "DAY65_CASE must be 1, 2 or 3\n");
return EXIT_FAILURE;
}
std::printf("pid %d: running buggy case %s with no watchdog.\n",
static_cast<int>(getpid()), env);
std::printf(
"This process is now expected to hang until you kill "
"it.\n");
std::fflush(stdout);
return fn(true);
}
if (runWatchdogged(probeDevice, false) != EXIT_SUCCESS) {
std::fprintf(stderr, "device probe failed\n");
return EXIT_FAILURE;
}
int status = EXIT_SUCCESS;
int buggySurvivors = 0;
for (int r = 0; r < rowCount; ++r) {
std::printf("%s\n", rows[r].label);
std::fflush(stdout);
const int result = runWatchdogged(rows[r].fn, rows[r].buggy);
if (result == -1) {
std::printf(" killed by the watchdog after %d s\n",
kWatchdogSeconds);
if (rows[r].buggy) {
std::printf(" diagnosis: %s\n", rows[r].diagnosis);
} else {
// A fixed case must never time out; that is the gate.
std::fprintf(stderr, " FAILED: a fixed case hung\n");
status = EXIT_FAILURE;
}
} else if (result == EXIT_SUCCESS) {
std::printf(" completed and passed its check\n");
if (rows[r].buggy) {
++buggySurvivors;
}
} else {
std::printf(" completed with a wrong answer (exit %d)\n", result);
if (!rows[r].buggy) {
status = EXIT_FAILURE;
} else {
++buggySurvivors;
}
}
}
// The buggy rows are reported and never gated. This program predicts
// all three hang, and a buggy case that finished is a finding to
// publish, not a failure to hide: day 27's racy kernel passed 100 runs
// and taught more by being reported than a gate could have.
if (buggySurvivors > 0) {
std::printf("\nreported, not gated: %d buggy case(s) did not hang\n",
buggySurvivors);
}
if (status == EXIT_SUCCESS) {
std::printf("\nall fixed cases completed and passed\n");
}
return status;
}