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

Debugging a kernel with cuda-gdb

A kernel doubles a 61 by 37 image, and arithmetic on its transposed store index says 2,244 of its 2,257 pixels come out wrong. Every call around it is wrapped in CUDA_CHECK and every one returns success. Nothing faults because the bad index never leaves the buffer, so memcheck has nothing to report.

The tools from days 61 and 62 cannot identify this error. The program uses legal addresses but produces wrong output, and the relevant variable exists inside a running kernel. You will compile for cuda-gdb, stop one chosen thread, print its local variables, and find the wrong index.

What -G buys you and what it charges

gdb cannot step through code it has no line table for, and an optimized kernel barely has lines: registers get renamed, statements fuse and reorder. nvcc -g -G fixes that by generating debug information for device code and turning device optimizations off (https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#device-debug-g , checked 2026-09-01). Breakpoints bind to kernel source lines, and a local like dst exists in a place the debugger can read instead of being folded into an address calculation.

-G reduces speed because it disables device optimizations. The program times its fixed kernel in both builds to measure the difference. That is also the rule for when to use which build: -G for stepping, -lineinfo for profilers and compute-sanitizer, never -G numbers on a results page.

With tens of thousands of threads, "stop" needs a clear subject. cuda-gdb keeps one thread in focus, named by its kernel, its block and its thread coordinates, and every command that reads state reads it from that focus. When a breakpoint in device code hits, everything stops, the focus lands on the thread that hit it, and cuda kernel block thread prints where you are.

cuda thread (10,3,0) moves the focus without running anything; info cuda threads lists the whole grid, grouped into runs of threads stopped at the same spot. The coordinates are the built-in index variables you have used since day 4, which is what makes the next section's arithmetic possible.

Diagram: one thread, two indices. A 61-wide, 37-tall pixel grid on the right, the 16 by 8 threads of block (1,0,0) on the left, with two arrows leaving thread (2,3,0). Band 1, the load: thread (2,3,0) is col 18, row 3, and reads pixel row * width + col. Caption "src = 3 * 61 + 18 = 201". Band 2, the store as the bug writes it: the same thread writes pixel col * height + row. Caption "dst = 18 * 37 + 3 = 669, still inside the 2,257-pixel buffer". Band 3, the 13 pixels where the two formulas agree, marked on the line from (0,0) through (5,3) to (60,36). Caption "3 * col == 5 * row: 13 fixed points, 2,244 misses". Alt text: "Thread (2,3,0) of block (1,0,0) reads pixel 201 and writes pixel 669. Both are legal addresses, so only thirteen of the 2,257 pixels come out right and nothing crashes."

A debugger does not find bugs in code that stops all threads politely

In CPU gdb, you can run to a bad line and inspect its state. On a GPU, first choose the thread whose state you need. A breakpoint with no condition stops the first warp to arrive, which may not be the thread whose arithmetic you worked out on paper.

cuda-gdb evaluates breakpoint conditions per thread on the device, so break scaleBuggy followed by condition 1 threadIdx.x==2 && threadIdx.y==3 && blockIdx.x==1 && blockIdx.y==0 stops exactly the thread you pre-computed, and the transcript below is about that thread rather than about luck.

Once stopped, you can inspect another thread in the same warp without running the program. A warp moves in lockstep, so when the focus thread, lane 18 of its warp, sits on the store line, lane 26, which is thread (10,3,0), sits there too, with its own col and dst already computed.

Switching focus gives a second data point. If both threads show the same index error, the formula is wrong rather than one thread's state.

The debugger inspects a known failure; it does not detect one. A CPU reference first shows that this kernel's output is wrong. cuda-gdb then identifies the line and variable that caused the mismatch.

The bug is supplied, the session is scripted

Full program in code/day63-cuda-gdb/scale_debug.cu, session script in commands.gdb. The pair is built to be replayed, not just read.

The bug is deliberate, labelled, and shipped next to its fix. It is day 7's transposed index, col * height + row where row * width + col belongs, on a store into global memory. Both formulas land inside the buffer for every thread, so the result is a scramble, not a crash:

// out = 2 * in, one thread per pixel. The load index follows the course
// formula. The store index is the same formula retyped from memory, and
// retyped wrong.
__global__ void scaleBuggy(const float* in, float* out, size_t width,
                           size_t height) {
    const size_t col =
        blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    const size_t row =
        blockIdx.y * static_cast<size_t>(blockDim.y) + threadIdx.y;
    if (col < width && row < height) {
        const size_t src = row * width + col;
        // DELIBERATE BUG (day 63): transposed store index, the day 7
        // classic. Every value it produces is still inside the buffer,
        // so nothing faults; the value just lands in the wrong pixel.
        // The fix is dst = src, as scaleFixed writes it.
        const size_t dst = col * height + row;
        out[dst] = 2.0f * in[src];
    }
}

The harness gates both kernels, in opposite directions. scaleFixed must match the CPU reference everywhere. scaleBuggy must miss on exactly the count the transposition predicts, computed by a host loop over the identity rather than hard-coded:

static size_t expectedBuggyMismatches() {
    size_t fixedPoints = 0;
    for (size_t row = 0; row < kHeight; ++row) {
        for (size_t col = 0; col < kWidth; ++col) {
            if (col * kHeight + row == row * kWidth + col) {
                ++fixedPoints;
            }
        }
    }
    return kElems - fixedPoints;
}

Every debugger command lives in a file, and the session runs in batch. cuda-gdb --batch -x commands.gdb ./scale_debug_g produces the same transcript every time: one conditional breakpoint, run, the focus report, info cuda threads, five nexts to walk the kernel's five statements, then print for col, row, src, dst, in[src] and the correct row * width + col, a focus switch to thread (10,3,0), and three of those prints again (col, dst and the correct formula). A session you drove by hand teaches once; a script teaches every reader the same thing and lets CI re-derive the page.

With in[i] = i the misplaced values are self-labelling, which real data never is. The input is chosen so the transcript reads plainly, not because the technique needs it.

Results

Run on the project's Tesla T4 (driver 595.84, CUDA 12.6 V12.6.85) on 2026-09-01. Transcripts in code/day63-cuda-gdb/evidence/. Every predicted number came back exactly as the arithmetic said it would.

CUDA 13.0 re-verified only the two builds and their bare runs. Both repeated the exact 2,244-of-2,257 buggy mismatch count and the fixed-build gate. The release timing stayed close at 0.0039 ms, while the standalone -G timing moved from 0.0064 to 0.0122 ms.

No CUDA 13 cuda-gdb session was captured, so the debugger transcript and thread-level claims below remain sourced from the older session rather than being promoted to CUDA 13 evidence.

Transcript What it holds
batch session, -G build [Switching focus to CUDA kernel 0, grid 1, block (1,0,0), thread (2,3,0), device 0, sm 2, warp 1, lane 18], then the breakpoint hit in scaleBuggy<<<(4,5,1),(16,8,1)>>>
the six prints, focused thread col 18, row 3, src 201, dst 669, in[src] 201, row * width + col 201
after cuda thread (10,3,0) col 26, dst 965, row * width + col 209, focus switched with no stepping
./scale_debug, both builds scaleBuggy: 2244 of 2257 pixels wrong (transposition predicts 2244), then scaleFixed matches the CPU reference on all 2257 pixels
the two timing lines release 0.0037 ms per launch, -G 0.0064 ms, so -G costs 1.7x here

The diagnosis is the two prints sitting next to each other: dst is 669 and row * width + col, the index the store should have used, is 201. The thread loaded the right pixel and put it in the wrong place, which is the whole bug in two numbers.

info cuda threads collapsed the grid to one line, (0,0,0) (0,0,0) to (3,4,0) (15,7,0) 2560, all 2,560 threads stopped at the same PC on scale_debug.cu:73. The second thread makes the pattern rather than the instance: 209 against 965 is the same transposition at a different lane, and it cost one focus switch and two prints, no stepping, because a warp moves together and its locals were already computed.

One number the page did not predict. The -G build under the debugger reported 0.1470 ms per launch against 0.0064 ms standalone. Being watched costs about 23x more than being built for watching. Time a kernel outside the debugger.

Two toolchain notes for anyone reproducing this. The nvidia channel's CUDA 12.6 cuda-gdb package ships only a wrapper that searches for cuda-gdb-python3.X-tui binaries it never installs, so it cannot start at all; the conda-forge 13.3 build debugs a 12.6-built binary fine under this driver. That build also spawns cuobjdump from its own bin directory when it needs to disassemble, so cuda-cuobjdump has to be installed and visible there or the first next fails.

Run it yourself

Install cuda-gdb with the toolkit on a CUDA-capable system, then run the three commands in the README. The debugger reads no performance counters, so day 42's permission error does not apply, as day 61 also explains. Batch mode also works in an environment without an interactive terminal.

Exercise

Pick thread (7,1,0) of block (2,3,0). On paper, compute its col, row, src and dst, then edit the condition in commands.gdb to stop that thread and check all four with print.

Time: 20 to 30 minutes. Submit: the four predicted numbers and the matching lines of your transcript, then the day63 quiz.

Check: the quiz below is keyed to the shipped transcript, and your own session either prints your four numbers or it does not. If the breakpoint never hits, your condition names a thread outside the 61 by 37 image's live region; that failure mode is question four's subject.

Hint 1

The condition has four clauses and you are changing all four. Before touching the debugger, decide whether the thread you picked survives the kernel's guard at all.

Hint 2

Block (2,3,0) starts at col 32, row 24. Add the thread offsets, then run both formulas: row * 61 + col and col * 37 + row.

Solution

col is 2 * 16 + 7 = 39, row is 3 * 8 + 1 = 25. So src is 25 * 61 + 39 = 1564 and dst is 39 * 37 + 25 = 1468. Both are legal indices, but only src is correct for this store.

Compute the expected values before starting the debugger. If a print differs, inspect the formula or state that produced it.

Pitfalls

Your breakpoint on the kernel never hits. The build has no device debug info, so the breakpoint bound to the host side or nowhere. Compile with -g -G for a debug session; -lineinfo is for profilers and sanitizers and is not enough here. The shipped session against the release build exists to show this failure on purpose.

You stopped in the kernel but in the wrong thread. An unconditional break stops the first warp to arrive. Put the block and thread coordinates in the breakpoint condition, as commands.gdb does, and verify with cuda kernel block thread before trusting any print.

You stepped once and several threads moved. next advances the focus thread's warp in lockstep, not one thread in isolation. That is the hardware, not a debugger defect. Neighbours in the warp provide additional data points.

You timed or profiled the -G build. Device optimizations are off, so every number is about the flag, not the kernel. Time the -O3 build; the program prints which build it is on its second line so a pasted transcript cannot hide the flavour.

You expected a tool to flag the bug for you. Every index the transposed formula produces is in range, so there is no illegal access to report and no race for day 62's tools to find. Only a reference tells you the output is wrong; day 12 turns this same index swap into the intended behaviour, which is why no tool can call it a bug.

You debugged the first launch and drew timing conclusions. Under -G everything is slow and the first launch also pays module load, the day 9 cost. Keep debugging and measuring in separate runs of separate builds.

Go deeper

Next

Day 64 stays inside the kernel but swaps the debugger for printf and device-side assert. It shows how buffering can delay, reorder, or discard printf output. Day 65 takes on the case where nothing prints because the kernel never finishes.

The three days share one lesson: what a GPU shows you by default is almost nothing. Each tool exposes a different type of state and has a different cost.