Day 70Module 7

in-technical-review

Checkpoint: the CUDA bug catalog

Day 27 ran a racy reduction 100 times. It produced the right answer every time, and it was still wrong. Day 14 measured a related case: its broken tile kernel was clean at 32 threads per block and wrong 17,312 times at 64.

A CUDA bug can exist long before it causes a failure. This checkpoint puts ten bugs that the course has written or taught into one labelled pipeline. Your job is to fix all ten.

For each bug, the catalog gives the symptom, the tool that names it, the day that taught it, and useful search text.

Diagram: ten bugs against five detectors. A grid: ten bug rows, five detector columns (memcheck, racecheck, synccheck, initcheck, your own gates), a mark where a detector sees the bug. Rows 3 and 7 marked under memcheck. Caption "memcheck: bounds and leaks, nothing else." Rows 4, 5 and 6 marked under racecheck, synccheck and initcheck, one each. Caption "one sanitizer tool, one hazard class." Rows 1, 2, 8, 9 and 10 marked only under gates, or not at all. Caption "five of ten: no tool. Reference, tolerance and discipline." Alt text: "Ten bug classes against five detectors. The four sanitizer tools cover five bugs between them; the other five are caught only by a CPU reference and habits, or by nothing."

1. The missing error check

Symptom: nothing until an unrelated line fails. A kernel launch returns void, so its failure surfaces at the next runtime call even when that call is valid.

No tool catches the absence of a check; the fix is the two-line habit after every launch, cudaGetLastError() then cudaDeviceSynchronize(), both wrapped. Day 6 gives the full explanation, including why the error checking macro prints the string rather than the number: the string is what people search.

2. The trusted cudaMemcpyKind flag: invalid argument

Symptom: none under CUDA 12.6, then an immediate invalid argument under CUDA 13.0 on the same T4. Day 6 measured a host-to-device copy tagged cudaMemcpyDeviceToHost under 12.6: it returned cudaSuccess and copied correctly, because unified addressing let that runtime infer the direction from the pointers.

The documentation calls the mismatch undefined behaviour, so CUDA 13 is equally allowed to reject it. The direction flag is not a portable way to catch your mistake; only reading the pointers is.

3. Out of bounds: an illegal memory access was encountered

Symptom: either the classic sticky error, an illegal memory access was encountered, reported by an innocent later call, or nothing at all, because cudaMalloc rounds allocations up and a short overrun lands in padding that never faults. This pipeline's version is the second kind: 195 threads past the end of a 1144-block launch, each reading and writing 4 bytes beyond the allocations, the farthest landing 780 bytes past the end, and built never to fault (prediction 1 says exactly that).

compute-sanitizer --tool memcheck counts every one of those accesses against the allocation's true size. Day 5 put the bounds guard in your first kernel; the guard checks the index against n, and memcheck checks the other half, the allocation.

4. The shared-memory race

Symptom: answers that come and go with the warp schedule, including long streaks of correct ones. The pipeline's blur loads a shared-memory tile and reads its neighbours' cells without a barrier in between. compute-sanitizer --tool racecheck reports the hazard on every run, whether the answer is right or wrong.

A passing test is one allowed outcome of a race. No error text prints; use racecheck's report as the search text. Day 13 built the tile, day 14 removed the barrier and measured what happens, and day 62 makes racecheck routine.

5. The barrier some threads never reach: barrier error

Symptom: on Volta and later, not a hang. An early return above a __syncthreads() lets the exited threads count as having arrived, so the kernel completes and the cells those threads owned hold garbage. compute-sanitizer --tool synccheck exists for exactly this: its manual's own example is a barrier inside a divergent branch, reported as a barrier error naming the block ( https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html , checked 2026-09-01).

The rule from day 14 stands: guard the loads and the stores, never the barrier, and never return above one.

6. The buffer nothing ever wrote

Symptom: garbage that is sometimes zero, so the program works on the runs where the driver handed back zeroed pages and fails on the others. The pipeline allocates a 611-float calibration table and, in the buggy build, never copies anything into it. compute-sanitizer --tool initcheck reports each read of global memory that no write preceded.

Nothing else does: the launch is legal, the addresses are in bounds, the values are values, and no error text ever prints.

7. The leaked allocation: out of memory

Symptom: none in a short program, then out of memory in the loop that matters. One cudaFree is missing here. compute-sanitizer --tool memcheck --leak-check full prints each unfreed allocation with its size at context teardown; without the flag, the tool stays silent about leaks, which is a separate risk.

The course rule since day 5: every cudaMalloc has a checked cudaFree on every path, which is why this program's gates fall through to one exit instead of returning early.

8. The wrong index order

Symptom: a wrong answer with clean sanitizer reports and no error text anywhere, which is the point. The pipeline's last stage sums each sensor's column but indexes the matrix sensor-major when the storage is sweep-major. Every access lands inside the buffer, so memcheck has nothing to say, and there is no race, no barrier problem, no uninitialised read.

Only the CPU reference notices. Day 4 is where index arithmetic starts, and this is why every stage in this course carries a reference, however trivial: the tools check memory discipline, not meaning.

9. The assert that is not in the binary

Symptom: a stage that is checked in your debug build and unchecked in the build you ship. assert() compiles to nothing under -DNDEBUG, which CMake's Release configuration adds for you, and this pipeline's buggy build line adds it on purpose. Bug 5 corrupts one block's QA figure and the assert that would have reported it does not exist.

No tool reports an absent check, and nothing prints. The course gates with branches that return EXIT_FAILURE for this reason, and day 64 covers what device-side assert is actually for.

10. The tolerance-free float compare

Symptom: a gate that rejects a correct kernel. The device fuses a * x + b into one FMA and the host reference does not, so the two sides differ in the last bit on some elements, which day 47 demonstrated with the programming guide's own worked example. An exact != compare turns that rounding difference into a bug report; the only error text here is the one your own gate prints.

Day 5's relative tolerance defines agreement for floats. Day 68 covers the case where you really do want bits.

One program, both variants

The full program is code/day70-bug-catalog/bug_catalog.cu, about 300 lines from raw sensor sweeps to per-sensor totals, built so the answer key and the exercise can live in one file without spoiling each other.

The buggy and fixed variants are one file. Every bug site carries a BUG n (deliberate, day 70) comment, and the repair sits beside it behind #if DAY70_FIXED, so the two binaries differ by one compile flag and nothing can drift between them. Building the buggy variant adds -DNDEBUG, because that is what a Release build does and bug 9 needs it.

Every stage has a reference and a gate. The gates fall through to a single exit so one failure cannot hide another and every allocation is freed on every path. The buggy build's gates are part of the exercise: one is an assert that is not there, one compares floats exactly.

The bug sites read like the catalog. Three of them:

    blurSamples<<<blocks, kThreadsPerBlock>>>(d_raw, d_blur, kElems);
#if DAY70_FIXED
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
#else
    // BUG 1 (deliberate, day 70): no cudaGetLastError, no synchronize. This
    // launch happens to be legal, so nothing is lost today. The day the
    // grid or the kernel goes wrong, the failure surfaces at whichever call
    // looks next, and that call is innocent. Day 6's two lines, always.
#endif

__global__ void calibrateSamples(const float* __restrict__ in,
                                 const float* __restrict__ offset,
                                 float* __restrict__ out, size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    const int c = static_cast<int>(i % static_cast<size_t>(kCols));
#if DAY70_FIXED
    if (i < n) {
        out[i] = kGain * in[i] + offset[c];
    }
#else
    // BUG 3 (deliberate, day 70): no guard, so each of the 195 threads past
    // the end reads and writes 4 bytes beyond in[] and out[], the farthest
    // landing 780 bytes past the allocations. cudaMalloc's rounding means
    // none of it should fault; memcheck counts every access anyway.
    out[i] = kGain * in[i] + offset[c];
#endif
}

#if DAY70_FIXED
    if (energyBad != 0) {
        std::fprintf(stderr,
                     "stage 1 energy: %zu of %d blocks wrong, "
                     "first at %zu\n",
                     energyBad, blocks, firstBad);
        status = EXIT_FAILURE;
    } else {
        std::printf("stage 1 energy: all %d blocks match\n", blocks);
    }
#else
    // BUG 9 (deliberate, day 70): the QA stage's only gate is an assert, and
    // the buggy build line defines NDEBUG, which expands it to nothing. Bug
    // 5 corrupts exactly one block's energy and this line is the one that
    // would have said so. The extra parentheses keep the course linter's
    // no-runtime-assert rule from firing on the line that exists to show why
    // that rule exists.
    (assert(energyBad == 0));
    std::printf("stage 1 energy: checked\n");
#endif

Because the answer key lives in the same file, the exercise only works if you fix the buggy branches yourself before reading the #if DAY70_FIXED sides. The comments name each bug's class, not its repair.

Results

Measured on a Tesla T4 with driver 595.84, CUDA 12.6 (V12.6.85), and compute-sanitizer 12.6 on 2026-09-01. The buggy binary printed pipeline: FAIL and exited 1. The fixed binary printed pipeline: PASS, exited 0, and all four sanitizer passes were clean.

CUDA 13.0 changes how far the buggy plain run gets. It rejects bug 2's deliberately mismatched cudaMemcpyKind immediately with invalid argument, so the buggy build exits 1 before the later planted bugs are reachable. That call is documented undefined behavior, and the specimen remains deliberately wrong.

The CUDA 13 fixed build passes every stage and exits 0. No CUDA 13 sanitizer passes were captured for this day, so the detector mapping below remains the CUDA 12.6 evidence rather than a claim that the later bugs were reachable under CUDA 13.

The buggy reports split this way:

Tool	Measured result
memcheck	196 errors; 1,170,676 bytes leaked in 1 allocation
racecheck	1 displayed race, with 1,156,852 and 13,828 hazards at the two reads in `blurSamples`
synccheck	0 errors
initcheck	292,864 uninitialised-read errors

The leak prediction held exactly, and racecheck and initcheck named the expected kernels. Memcheck reported 196 errors rather than the predicted 390. Synccheck did not report the deliberately divergent barrier, even though the buggy pipeline failed its own output checks.

A clean tool report is evidence about that tool's findings, not proof that the kernel is correct.

The sanitizer mapping is therefore narrower than predicted: memcheck sees bugs 3 and 7, racecheck sees 4, initcheck sees 6, and synccheck misses 5. Bugs 1, 2, 5, 8, 9 and 10 appear in none of the four sanitizer reports. The program's CPU comparisons expose corrupted stages, but finding the root cause still requires the code and its gates.

Run it yourself

You need a CUDA GPU that can run the binary. You do not need root access or performance counters, as day 61 established for the whole module.

compute-sanitizer ships with the toolkit; on a conda-style install it is the cuda-sanitizer-api package. Both build lines and all eight sanitizer commands are in the repo's README in a batch you can paste.

NVIDIA's teaching notebook also shows Compute Sanitizer commands: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/cuda-cpp/notebooks/03.02-Kernels/03.02.04-Dev-Tools.ipynb (checked 2026-09-01).

Exercise

Build the buggy variant, run all four sanitizer tools on it, and fix the ten bugs until your repaired build passes memcheck, racecheck, synccheck, initcheck and its own gates. Keep a log: for each bug, one line naming what caught it, a tool, a gate, or only your eyes.

Time: 45 to 90 minutes; this is a checkpoint, not a lesson-sized exercise. Submit: your fixed bug_catalog.cu and the ten-line log.

Check: your build of the fixed program must print pipeline: PASS and exit 0, and the four sanitizer commands from the README must each report zero errors on it (--error-exitcode 1 makes that scriptable). The buggy build must fail its own gates; if it passes on your card, that run is a lucky race and worth writing down, not a pass.

Hint 1

Run the tools before you read any code, and fix in the order that unblocks observation: the bugs that corrupt data upstream (the copy that never happened, the race) make every downstream gate fail, so their counts mean nothing until the upstream stages are clean.

Hint 2

When all four tools come back clean and a gate still fails, stop looking for a memory bug. Two of the ten are in the gates themselves, and one is in an index expression the tools consider valid. What does each gate actually compare, and against what?

Solution

The ten sites are marked BUG n (deliberate, day 70) in the source, each with its repair in the #if DAY70_FIXED branch beside it, so the diff is the file's own preprocessor conditionals. The mapping the log should show: memcheck catches 3 and 7, racecheck 4, and initcheck 6. On the measured T4, synccheck reported zero errors for 5.

The CPU reference catches 8 and the corrupted stages caused by 5 and 6. Nothing names 1, 2, 9 or 10 until a person reads the code or the gate design.

General rule: a sanitizer verifies memory discipline, a reference verifies meaning, and the checks themselves need checking, because a gate that is compiled out or compares floats exactly fails in the direction you will not notice.

Pitfalls

Every call after the first failure returns the same error, so you count ten bugs where there is one. An illegal access is sticky and leaves the context; read the first report, fix it, rerun. Day 6 and the error page cover the split between sticky and recoverable codes.

A clean memcheck run gets read as a clean program. memcheck checks addresses against allocations. Bugs 8 and 10 in this very program survive it untouched, and so does every wrong-but-in-bounds index you will ever write. The reference is not optional.

The race disappears when you shrink the repro. Day 14 measured its racy kernel clean at 32 threads per block and wrong 17,312 times at 64: one warp is ordered by the hardware, so a one-warp test case cannot show a cross-warp race. Shrink the data, never the block count or block size, and trust racecheck over ten green runs.

--leak-check is off by default. Plain memcheck says nothing about bug 7; the flag is --leak-check full, and the report prints at context teardown, after your program's last line of output.

You run the sanitizer once, on the tool you like. The four tools have disjoint jobs: this program contains a race, a barrier bug and an uninitialised read, and each is visible to exactly one of racecheck, synccheck and initcheck. Four short runs, not one long one; Compute Sanitizer is four detectors in one binary. For the bug none of them finds, day 63's cuda-gdb session is the next tool to use.

Go deeper

Compute Sanitizer user manual, the four tools and their report formats: https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html (checked 2026-09-01)
CUDA Runtime API, Error Handling, for sticky against non-sticky codes: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__ERROR.html (checked 2026-09-01)
cuda-gdb user guide, for the bugs you have to watch happen: https://docs.nvidia.com/cuda/cuda-gdb/index.html (checked 2026-09-01)
The whole cudaError_t enum, filterable, with the string each code returns: /errors/

Module 7 ends here. Day 71 opens module 8 with mixed precision: FP16 storage, FP32 accumulation, and error bounds checked by a program. Tensor-core kernels still need day 70's tolerance rules when their results differ in the last bit.