Where the Python stack goes dark

Almost all LLM code is Python, and almost all of the time a Traceback or a pdb breakpoint is all you need. The debugger I'm about to use earns its keep at exactly the moment that stops being true — when the failure is below the framework, in the native CUDA / C++ / collective layers where a Python stack trace just says RuntimeError and goes dark.

Those are the bugs that actually halt real work: an illegal memory access inside a custom or vendor attention kernel, a NCCL hang where one rank diverged, a segfault in a C++ extension linked into PyTorch. The good news is that the tools form a ladder, and every rung is GDB or a GDB dialect — so the skill transfers as you climb from one CPU thread to thousands of GPU threads to thousands of MPI ranks. Here are the bottom two rungs, on real hardware (an NVIDIA T4 via Brev), with the transcripts.

Rung 1 — CPU: gdb finds the fault and the caller

The deliberately-broken program passes a null pointer into a summation loop:

static long compute_total(const int *data, int n) {
    long total = 0;
    for (int i = 0; i < n; i++)
        total += data[i];          /* SIGSEGV here when data == NULL */
    return total;
}
int main(void) {
    int *data = NULL;              /* BUG: never allocated */
    printf("total = %ld\n", compute_total(data, 8));
}

Build with symbols and no optimization (gcc -g -O0), run under gdb, and the backtrace names both the crash site and the call that caused it:

Program received signal SIGSEGV, Segmentation fault.
0x...519d in compute_total (data=0x0, n=8) at cpu/segfault.c:16
16          total += data[i];
#0  compute_total (data=0x0, n=8) at cpu/segfault.c:16
#1  main () at cpu/segfault.c:23

data=0x0 in the frame is the whole diagnosis — and bt shows it came from line 23. (For a native crash inside Python, gdb --args python train.py plus the py-bt helper interleaves the Python and C stacks — invaluable when a torch op faults.)

Rung 2 — NVIDIA GPU: sanitizer first, then cuda-gdb

Now a CUDA kernel with the most common GPU bug — a missing bounds check:

__global__ void vadd(const float *a, const float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];            /* BUG: no `if (i < n)` */
}
// launched with 1024 threads (4×256) for n = 1000 → 24 threads run off the end

For an illegal-access bug, reach for compute-sanitizer before you reach for a debugger — it pinpoints the faulting access without any stepping:

========= Invalid __global__ read of size 4 bytes
=========     at vadd(...)+0x330 in vector_add_oob.cu:17
=========     by thread (232,0,0) in block (3,0,0)
=========     Access ... is out of bounds
=========     and is 1 bytes after the nearest allocation ... of size 4000 bytes

Exact line (17), exact lane (block 3, thread 232 → global index 1000), and the 4000-byte (= 1000-float) allocation it overran. When you do want to step in, cuda-gdb lets you focus a single GPU thread out of the grid and read its state:

CUDA thread hit Breakpoint 1, vadd<<<(4,1,1),(256,1,1)>>> (... n=1000) at vector_add_oob.cu:17
$1 = 1000      # print n
$2 = 0         # print i      — focus on block (0,0,0) thread (0,0,0)
[Switching focus to block (3,0,0), thread (232,0,0)]
$3 = 1000      # print i      — the out-of-bounds lane: i == n

cuda thread (232,0,0) block (3,0,0) moves the lens to the bad lane, and i == n == 1000 is the off-by-the-grid-tail bug made visible. (Two version gotchas from the real run: set cuda memcheck was removed in CUDA 12.x — use compute-sanitizer; and break on the use line, not the kernel's first line, or i reads <optimized out> because it isn't computed yet.)

The rungs above (same DNA, more parallelism)

The ladder keeps going; only the axis of parallelism changes:

  • AMD GPUs — rocgdb. ROCm's GDB for HIP kernels; "warps" become "wavefronts," CUDA becomes HIP, but the focus-a-thread workflow is identical. Relevant as AMD (MI300X-class) takes more LLM share.
  • MPI / multi-node — gdb4hpc. HPE's GDB-based parallel debugger: address hundreds of ranks at once via process sets (bt $a{0..7}), attach to a running srun job, and use comparative debugging to find the first rank where two runs diverge — the scalable way to chase "it broke after we changed the parallelism mapping."
  • At scale — Linaro DDT. The commercial, GUI parallel debugger (lineage: Allinea → Arm → Linaro) that unifies CPU + NVIDIA + AMD across thousands of ranks, when stitching the open tools together by hand gets painful.

The takeaway

When an LLM job dies below the framework, the move isn't to add print statements to Python — it's to drop to the right rung of the debugger ladder: sanitizer to localize a memory error, gdb/cuda-gdb/rocgdb to inspect the offending thread, gdb4hpc/DDT to find the one rank that went wrong. Build with -g (and -G/-ggdb for device code), reproduce small, and read the backtrace.

The runnable demos and full transcripts are open at codeberg.org/srinathv/llm-masters (projects/05-debugging/); the cross-tier playbook is in that repo's debugging tutorial. If your team is fighting a crash or hang below the framework, that's the work I do.

References

  1. cuda-gdb: https://docs.nvidia.com/cuda/cuda-gdb/ · compute-sanitizer: https://docs.nvidia.com/compute-sanitizer/
  2. rocgdb (ROCm): https://rocm.docs.amd.com/projects/ROCgdb/en/latest/
  3. gdb4hpc (HPE Cray PE): https://cpe.ext.hpe.com/docs/latest/debugging-tools/gdb4hpc/index.html
  4. Linaro Forge (DDT/MAP): https://www.linaroforge.com/