100 DAYS OF GPU Course preview
Day 65Module 7
in-technical-review

Deadlocks and hangs

A wrong answer gives you something to inspect. A hang may leave a stuck terminal, an ignored ^C, and nvidia-smi reporting 100 percent use. That figure means only that a kernel is resident, not that it makes progress.

You get no error string or bad index to search for. Day 14 promised this page would diagnose the hang its barrier rule can cause.

This page creates three hangs on purpose. It keeps each one inside a child process that a watchdog kills, then uses tools to find who waits for whom.

Every hang is a wait cycle

A GPU program hangs when something waits for something that is, directly or through others, waiting for it. The cycle can live at three different scopes, and this page builds one of each.

Inside a block. The block's threads wait for each other. In case 1, half the block waits at a __syncthreads() that only fires when all 256 threads arrive, waiting to hand off a tile in shared memory, while the other half spins on a flag nobody will ever write, because the write sits two lines below the barrier.

Since compute capability 7.0, a warp can leave lanes stuck without stalling the whole card. A barrier still counts arrivals, and 128 of 256 is not 256.

Across blocks. Blocks wait for a block the hardware cannot start. A block, once running, holds its SM slot until it exits; nothing preempts it.

A homemade grid-wide barrier, where each block counts itself in and then spins until the count reaches gridDim.x, is correct exactly as long as every block in the grid is resident at once. Day 28 measured that capacity on the course's Tesla T4: 4 blocks per SM at 256 threads, 160 blocks across its 40 SMs. That result is one measured example, not a fixed limit for other cards.

Case 2 asks the occupancy API for that number at run time and then launches one block more. Every resident block spins waiting for an arrival that needs an SM slot none of them will give up.

Inside one thread. A loop's exit condition can stay true forever. Case 3 adds 0.1f to a total and tests x != 1.0f.

Float 0.1f is slightly above one tenth, so the tenth addition lands just past 1.0f. Once the total is large enough, adding 0.1f rounds back to the same value. One thread looping forever prevents the launch from finishing.

The widget below is day 14's block model with the barrier-in-if preset loaded, which is case 1's shape: step it and watch the barrier counter stop short of the block.

The guide promises less than a hang for case 1's form: a __syncthreads() under a non-uniform condition means "execution may hang or produce unintended side effects" (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html , checked 2026-09-01). May. Case 1 is built so that may becomes does: the threads outside the if neither exit nor reach a barrier, they block on data from beyond it.

The diagnostic ladder, and a harness that cannot hang CI

Full program in code/day65-deadlocks/deadlocks.cu. The harness must contain each hang before you run any diagnostic tool.

Every case runs in a forked child, and the parent never touches CUDA. You cannot cancel a stuck kernel from inside its process without destroying the CUDA context. The child owns the context, so SIGKILL returns its resources to the driver while the parent remains able to print the table.

The fork happens before the child's first CUDA call. A process that has initialized CUDA must not fork children that use the GPU.

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;
}

Every bug ships beside its fix, and only the fixes are gated. The fixed variants must complete and pass, on real branches returning EXIT_FAILURE. The harness expects the watchdog to kill the buggy variants and reports either outcome. A program with a hang bug may still finish, and day 27 already showed what a report is worth when the bug declines to fire. Here is case 1's kernel; the deliberate bug is the barrier's position, and the launch hangs at the program's cudaDeviceSynchronize():

__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];
    }
}

Case 2's bug is a launch size, not a kernel. Both variants run the same spin barrier; the switch is one block:

    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;

And case 3's bug is two characters:

        while (x != 1.0f) {  // deliberate bug: the test must be <, not !=
            x += in[i];
            ++count;
        }

Use the following checks when a program hangs without warning.

Check 1: nvidia-smi. A silent program may show 100 percent utilization. The number counts time during which a kernel ran, not whether it made progress (https://stackoverflow.com/questions/40937894/nvidia-smi-volatile-gpu-utilization-explanation , checked 2026-09-01). A spinning deadlock and a healthy long kernel look identical here.

Check 2: attach cuda-gdb. This tool shows where threads are waiting without counters or root access, as on day 63.

Its manual documents attaching to and detaching from a running CUDA application with GDB's ordinary attach commands (https://docs.nvidia.com/cuda/archive/12.6.2/cuda-gdb/index.html , checked 2026-09-01). Attached, info cuda kernels lists the kernels still on the device, and a backtrace shows the host thread parked in the synchronize.

The program's DAY65_CASE=1 mode gives you a live hang to inspect. The README has batch commands so the whole transcript can be captured by a script.

Check 3: add a watchdog to the harness. The next hang then fails a test instead of freezing a session. The code above does this, and day 66 folds it into the course harness.

Results

Run on the project's Tesla T4 (driver 595.84, CUDA 12.6 V12.6.85) on 2026-09-01. Transcript in code/day65-deadlocks/evidence/. All four predictions held.

The device timeout query now uses cudaDeviceGetAttribute with cudaDevAttrKernelExecTimeout, because CUDA 13 removed the old cudaDeviceProp.kernelExecTimeoutEnabled member. Fresh CUDA 12.6 and CUDA 13.0 builds produced the same table exactly: all three buggy cases reached the five-second watchdog, all fixed cases passed, capacity stayed at 160, and the runtime limit remained off. The API compatibility edit changed no lesson behavior.

case fixed variant buggy variant
1, barrier in an if completed and passed its check killed by the watchdog after 5 s
2, grid barrier, one block over capacity 160 blocks, launched 160, passed killed by the watchdog after 5 s
3, x != 1.0f 10 steps everywhere, final total 1.00000012 killed by the watchdog after 5 s
  1. Held. Three buggy cases hung and were killed at five seconds; three fixed cases completed and passed. The program exited 0, because a hang the watchdog catches is the expected result, not a failure.
  2. Held exactly. capacity 160 blocks (4 per SM x 40 SMs), launching 160, the same figure day 28 measured for a 256-thread cooperative launch on this card. One block more never returns.
  3. Held. driver run time limit on kernels: off. Nothing outside this program would have stopped the hang, which is why the harness carries its own timeout.
  4. Held, and the number is the argument. every total passed 1.0f in 10 steps, final 1.00000012. The accumulated total steps over 1.0f without ever landing on it, so x != 1.0f never goes false. A convergence test that asks for equality on a float is asking the arithmetic for a promise it does not make.

The diagnoses apply beyond this program. Case 1: 128 threads wait at a barrier for 128 threads that wait on a flag written two lines below that barrier. Case 2: every resident block spins for an arrival that needs the one block the scheduler has nowhere to place.

Case 3: the loop's exit condition can never go false. Three different causes, but the same question applies each time: who waits for whom?

Run it yourself

Run this on Linux with a supported NVIDIA GPU. The program needs a shell for fork, environment variables, and a second terminal for nvidia-smi and cuda-gdb, so Compiler Explorer cannot run the full exercise.

The watchdogged run needs no root or counters. Attaching cuda-gdb to a process you did not launch as a child may need extra permission: many distributions ship /proc/sys/kernel/yama/ptrace_scope as 1, and then cuda-gdb must run as root or the scope must be set to 0, per the cuda-gdb manual page linked above. sudo on a machine you own; on Colab, run the attach from the same notebook shell that started the hang.

Exercise

Answer the four questions, then take the quiz. Both are answerable from this page and from one run of the program.

Time: 25 to 35 minutes. Submit: four answers, each naming the case, the doc sentence or the output line that settles it.

Check: the reveal below, plus the day 65 quiz. The program's own gates only cover the fixed variants; these questions are about reading the hangs, which no exit code can grade.

  1. Case 1's consumers spin on a flag. Suppose you "fix" the kernel by moving the flag write above the barrier instead of moving the barrier. Does it still hang, and is it correct?
  2. Case 2 launches capacity plus one and hangs. Would capacity plus one ever complete on a bigger card, and what does that mean for testing this pattern?
  3. You attach cuda-gdb during DAY65_CASE=3 and every thread is inside stepsToOneStuck. How do you tell this hang from case 1's, where every thread is also inside a kernel?
  4. Your own program hangs at a cudaMemcpy. Why is the copy almost certainly innocent?
Hint 1

For each question, name the wait cycle first: who is waiting, and for what. Two of the four turn out not to contain a cycle at all.

Hint 2

For question 3, ask what kind of statement each kernel is stuck on, and what a backtrace per thread would show waiting versus looping. For question 4, ask which earlier call was allowed to return before its work finished.

Solution

1. It stops hanging and starts racing. With the flag written before the barrier, the consumers are released, but nothing orders the producers' tile writes against the consumers' tile reads anymore; the barrier that would have done it still sits where only half the block goes. You have replaced a visible hang with day 14's silent wrong answer, and compute-sanitizer's racecheck is the tool that would object.

2. Yes. Capacity depends on the card and the kernel's resource use, so 161 blocks hang a 160-block GPU but complete on a card that holds more. Test the pattern at a grid the smallest supported card cannot hold, or better, use a cooperative launch, which refuses an oversized grid with an error code instead of wedging.

3. By what the threads are stuck on. In case 1 the producers are parked at a barrier instruction and the consumers in a tight flag loop, two distinct positions splitting the block in half; in case 3 every thread loops through the same few lines with no barrier anywhere. That difference is rung 2's whole value; nvidia-smi cannot see it.

4. Because kernel launches return before the kernel runs. The launch is asynchronous, the copy is the first call afterwards that must wait for the device, so the hang surfaces there, whatever the copy's arguments say.

A hang is a wait cycle. Once you can say who waits for whom, its scope (block, grid, or one thread) narrows the fix.

Pitfalls

Your program hangs at cudaMemcpy or at a sync, and you debug that call. Launches are asynchronous; the first blocking call after a stuck kernel inherits the wait. Attach and look at the device before blaming the line the host is parked on.

You read 100 percent GPU utilization as progress. The figure counts time a kernel was executing on the device. A deadlocked spin scores the same as productive work, so it can confirm a kernel is resident and nothing else.

You expect the hang from cudaStreamWaitEvent on an unrecorded event. An event before its first record "represents an empty set of work" (https://docs.nvidia.com/cuda/archive/12.6.2/cuda-runtime-api/group__CUDART__EVENT.html , checked 2026-09-01), so the stream's wait passes immediately. The bug this produces is missing synchronization, not a hang: put the cudaEventRecord before the wait that names it in program order.

You conclude divergent code around __syncthreads() is safe because it did not hang for you. The guide's wording is "may hang or produce unintended side effects", and day 14 showed the non-hanging outcome is a silent race: a barrier waits only for threads that have not exited, so an early return above it counts as arrival and the kernel completes with unordered writes. Case 1 hangs precisely because its spinning consumers never exit. Keep barrier conditions uniform across the block; day 62's synccheck tool exists for the ones you cannot eyeball.

You ship a homemade spin barrier that passed every test. It passes while the grid fits residency and wedges the first time it does not, and the threshold differs per card. A grid-wide barrier wants day 28's cooperative launch, which fails loudly at launch time instead of hanging at run time.

You kill the process and assume the card is wedged because utilization lingered. Process exit destroys the context and the driver reclaims the device; a stuck kernel does not outlive its process. Persistent load after your process is gone belongs to someone else; nvidia-smi lists pids.

Go deeper

Next

Day 66 takes this page's watchdog and the course's pass-fail harness and turns them into a test suite with property tests, tolerances and the sanitizers wired into CI, so the next hang or race fails a check instead of a demo. Behind it sit day 63 for driving cuda-gdb properly and day 62 for the tools that grade barrier bugs.