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

Finding memory bugs with compute-sanitizer

The program for this page runs four kernels and compares all 6,151 elements of each against a CPU reference. Its bare run prints all 6151 elements match four times and exits 0. Yet it contains three labelled memory bugs: a read past an allocation, a write past a shared tile, and a cudaMalloc that nothing frees.

Day 27 already showed a racy kernel passing 100 runs in a row. Here, the output check misses invalid memory accesses instead of a timing error. You will use a compute-sanitizer report to find each error class, address, thread, block, and launch stack.

What the hardware checks, and what it does not

Out of bounds does not mean crash. cudaMalloc carves your request out of larger mapped regions, with alignment padding after it.

The hardware raises cudaErrorIllegalAddress only when a warp touches an address with no mapping behind it at all. Everything between the end of your 24,604-byte buffer and the end of the mapped region may remain mapped. A load there can complete and return data without raising an error.

memcheck checks that allocation boundary in software. It is one of four tools inside the compute-sanitizer binary; the others are racecheck, initcheck, and synccheck. It instruments every global, shared and local access in your kernel, checking each against a table of what you allocated.

The instrumented run is much slower. In return, the report names the access, kernel, thread, block, byte count, and distance past the allocation.

The same allocation table also supports leak checking. Run with --leak-check full and at context destruction the tool prints every device allocation that was never freed, with its size and the host call stack of the cudaMalloc that created it. The manual defines the leak set as allocations still live when the context dies, whether by process exit, cudaDeviceReset() or cuCtxDestroy() (https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html , checked 2026-09-01).

These checks do not use profiler counters. Unlike day 42's ncu on a restricted driver, every command on this page runs as an ordinary user.

Diagram: where four bytes past the end actually lands. Three horizontal bands, each an address line running left to right. Band 1, the hardware's view: a 24,604-byte allocation, then a stretch of padding, then the end of the mapping far to the right. An arrow for the read at byte 24,604 lands in the padding. Caption "no fault until the mapping ends, and the mapping ends nowhere near your buffer." Band 2, memcheck's view: the same line with a wall drawn at exactly byte 24,604. The same arrow hits the wall. Caption "one report: 4 bytes, out of bounds, thread (6,0,0), block (24,0,0)." Band 3, the shared tile: 258 cells ending at byte 1,032 inside a 1,280-byte block allocation. An arrow writes cell 258, in the padding. Caption "25 blocks, 25 stray writes, zero wrong answers." Alt text: "A read four bytes past a 24,604-byte buffer lands in padding the hardware never checks. memcheck draws the bound where the allocation ends and reports the thread and block that crossed it."

Reading one report block

Every memcheck finding is a block of ========= lines with the same anatomy. The manual's own example, from a kernel handed a deliberately odd pointer (same URL as above, checked 2026-09-01; this output is the manual's, not ours):

========= Invalid __global__ write of size 4 bytes
=========     at unaligned_kernel()+0x100 in memcheck_demo.cu:34
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x7cc43aa00001 is misaligned

Line one gives the class: the direction (read or write), memory space (__global__, __shared__, __local__) and the access size. Line two gives the source location, with a file and line number when you compile with -lineinfo. Line three gives the exact thread and block coordinates.

Line four gives the address and its relation to the nearest allocation. This example shows another report class: an address can be in bounds but misaligned, because a 4-byte word must start on a 4-byte boundary. The saved launch backtrace below the block points from library code to your calling code.

The number of reports helps locate the bug. One bad access points at a special thread, often an edge case. One report per thread block points at a line every block runs, such as a halo load.

Hundreds of reports per warp suggest a broken index formula. Use that pattern before opening the source.

A program built to be diffed

Full program in code/day61-sanitizer/memory_bugs.cu. Everything about its design serves the diff between the supplied build and the clean one.

Every bug ships next to its fix. The file holds four kernels: an adjacent difference and a shared-memory windowed sum, each in a correct variant and a broken one, plus a kRunBuggyPass constant that removes the broken pass entirely. The sanitizer transcript of the supplied build and the transcript of the clean build differ only by the planted bugs, so every line of noise has a known cause.

The bare run must pass. Each bug is engineered to miss every byte the harness checks, every CUDA call is wrapped in the course's error macro from day 6, and the checks are real branches that return EXIT_FAILURE. The deliberate design makes the bare run pass. Bug 1:

// The same kernel with the guard moved after the load.
//
// Deliberate bug 1 (day 61): the load of in[i + 1] is unconditional, so
// the last thread (i == n - 1, thread (6,0,0) of block (24,0,0) at these
// sizes) reads in[n], four bytes past the end of a 24,604-byte
// allocation. The bare run cannot see it: cudaMalloc hands out more than
// it promises, the hardware only faults at the edge of a mapping, and
// the select below throws the loaded value away for exactly that thread.
// The store through `tap` is why the compiler cannot quietly guard the
// load for us: the value is always used, just never checked.
__global__ void diffNextPastEnd(const unsigned int* __restrict__ in,
                                unsigned int* __restrict__ out,
                                unsigned int* __restrict__ tap, size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        const unsigned int next = in[i + 1];  // reads in[n] when i == n - 1
        tap[i] = next;
        out[i] = (i + 1 < n) ? next - in[i] : 0u;
    }
}

Bug 2 is one character. The tile staging loop of the windowed sum, from day 13's pattern, runs one cell too far in every block:

    for (int c = static_cast<int>(tid); c <= kTileCells;
         c += kThreadsPerBlock) {  // <= runs one cell too far
        const size_t g = blockStart + static_cast<size_t>(c);
        tile[c] = (g >= 1 && g <= n) ? in[g - 1] : 0u;
    }

Cleanup runs on every path. Checks record a status and fall through to one free block at the bottom, because a program that exits early under --leak-check full reports every live buffer and buries the real leak. Bug 3 is the one allocation missing from that block:

        // Deliberate bug 3 (day 61): a scratch buffer somebody added while
        // debugging bug 1 and never removed. It is written by every thread
        // of diffNextPastEnd, read by nobody, and freed nowhere on any
        // path, so --leak-check full must report exactly one leaked
        // allocation of 24,604 bytes.
        unsigned int* d_tap = nullptr;
        CUDA_CHECK(cudaMalloc(&d_tap, bytes));

tap exists so the compiler cannot legally move the unguarded load into the branch that uses it, which would fix bug 1 in the binary while it stayed in the source. Whether -O3 would really do that is exactly the kind of question this course settles by running, and the Results section bets on it.

Note. initcheck, the fourth tool, finds reads of memory nothing ever wrote. Day 35's radix sort carries a cudaMemset whose comment says plainly it exists "so compute-sanitizer --tool initcheck stays quiet on the first pass": the scanned slots above the active digit range never influence the answer, but the tool cannot know that, and zeroing them costs less than arguing. Keep the tools clean even where the noise is provably harmless, so a real finding never hides in accepted noise.

Results

Run on the project's Tesla T4 (driver 595.84, CUDA 12.6 V12.6.85) on 2026-09-01. Transcripts in code/day61-sanitizer/evidence/. All five predictions held.

CUDA 13.0 with compute-sanitizer 13.0.85 reproduced the contract on 2026-09-02: the bare run exited 0, default memcheck stopped after the invalid global read with 2 errors, kernel-destroy mode reported 27 errors including 25 shared writes and the 24,604-byte leak, and the fixed build reported 0 errors and 0 leaked bytes. No finding count or verdict drifted.

  1. Held. The bare run of the supplied build printed four all 6151 elements match lines and exited 0. Bug 2's stray shared write landed in the block's own padding, exactly as the construction intended, so nothing observable went wrong.

  2. Held. Plain compute-sanitizer ./memory_bugs reported one Invalid __global__ read of size 4 bytes, and the tool named the distance rather than making you compute it: is 1 bytes after the nearest allocation at 0x74ace3000000 of size 24604 bytes.

  3. Held, and the error code is worth keeping. That default run never reached bug 2. After the invalid read the tool tore the context down and the next checked call failed: Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.

    Two errors appear in the summary: the bug and the later API failure. Error 719 is what a learner without the sanitizer would have seen, with none of the information.

  4. Held to the count. With --destroy-on-device-error kernel --leak-check full, the report carried 1 invalid global read, 25 Invalid __shared__ write of size 4 (one per block, every block seen before truncation), Leaked 24604 bytes at 0x7fac6300c400, and LEAK SUMMARY: 24604 bytes leaked in 1 allocations. Total: 27 errors.

  5. Held. The kRunBuggyPass = false build reported ERROR SUMMARY: 0 errors, LEAK SUMMARY: 0 bytes leaked in 0 allocations, and exited 0.

Transcript What it holds
bare run four match lines, exit 0
memcheck, defaults 1 invalid read, then error 719 at the next sync, 2 errors
memcheck, kernel mode + leaks 1 read, 25 shared writes, 1 leak of 24,604 bytes, 27 errors
clean build 0 errors, 0 leaks, exit 0

The bare run and the clean report are two different claims. This program passed every check it makes about its own output while reading past an allocation, writing past a shared tile and leaking 24,604 bytes, and the only thing that said so was the tool.

Run it yourself

Run this on a CUDA-capable GPU where compute-sanitizer is installed. memcheck reads no performance counters, so day 42's ERR_NVGPUCTRPERM error does not apply. Compiler Explorer cannot run it because its runner executes the binary directly and cannot wrap it in another tool.

Build with -lineinfo or your reports will name code offsets instead of source lines. The exact commands, in order, are in code/day61-sanitizer/README.md.

Exercise

Run the three sanitizer commands from the README against the supplied build. Map each report block to one of the three labelled lines, naming the error class, the thread and block, and how far outside which allocation it landed. Then fix all three in place, without touching kRunBuggyPass, until the --error-exitcode 1 run exits 0.

Time: 30 to 45 minutes. Submit: the three changed lines and the clean final transcript.

Check: compute-sanitizer --tool memcheck --leak-check full --error-exitcode 1 ./memory_bugs must exit 0 with the program still printing four match lines. Before your fixes it exits nonzero; if it still does after, the report block that remains names the kernel you have not actually fixed.

Hint 1

Do the mapping before you touch code. Each block answers three questions: what kind of access, who did it, and how far from the nearest allocation. Two of those three answers are enough to pick the source line every time.

Hint 2

Count the blocks per finding. One error total points to an edge thread. One per block points to a line every block runs, and only the staging loops do that here.

The leak appears only at exit under a flag that no kernel can trigger.

Solution

For bug 1, move the boundary test before the load, as diffNext does, so in[i + 1] is read only when i + 1 < n. For bug 2, change <= to < in the staging loop, as in windowSum. Bug 3: CUDA_CHECK(cudaFree( d_tap)) in the cleanup block at the bottom of main.

The fixed kernels in the same file show the expected diff. The exercise grades the transcript rather than the code.

What generalises: memcheck enforces the bound your cudaMalloc stated, while the hardware enforces only the edge of a mapping, and every bug on this page lived in the space between those two lines.

Pitfalls

Your tests pass, so you do not run the tool. All three bugs here sit behind a green harness, the way day 27's race sat behind 100 correct runs. Run memcheck on code you believe, not just code you suspect.

One bug per run, and you thought you had one bug. By default the first invalid access tears down the context, so the next checked call fails with unspecified launch failure and everything after it never executes (/errors/unspecified-launch-failure). That API error is fallout, not a fourth bug. Fix and rerun, or pass --destroy-on-device-error kernel to collect everything in one pass.

The report names an offset, not your source line. Without -lineinfo the location line stops at kernel()+0x100. Rebuild with -lineinfo; unlike -G it costs nothing you can measure at run time and works with -O3.

Every buffer in the program shows up as a leak. --leak-check full reports what is live at context destruction, so a program that exits through an error path before its frees reports all of them. That is the status fall-through rule from day 6: cleanup must run on failing paths too, or the report may hide the one deliberate leak among unrelated leaks.

You timed the run because it was there. Instrumenting every memory access costs an order of magnitude or more, and nothing about a sanitized run's duration describes your program. Correctness numbers from tooled runs, performance numbers from bare ones, the same separation day 42 draws for the profiler.

Go deeper

Next

Day 62 points the same binary habit at time instead of space: racecheck for shared-memory races and synccheck for broken barriers, which finally makes good on day 14's promise to prove its race with real tool output. Day 63 uses an interactive cuda-gdb session to step through a failing kernel.