Markdown

CUDA-GDB: device debugging and GPU core dumps¶

Scope: cuda-gdb as an instrument on a shared GPU server. What device state it exposes that nothing else can (the faulting warp, its lanes, its registers, the PC), the price it charges for that (a breakpoint on one GPU pauses every GPU running CUDA on the host), and the one mode that belongs anywhere near production: the GPU core dump, which turns an illegal memory access that only reproduces on rank 47 after six hours into a file you open on your laptop. For correctness bugs you can reproduce cheaply, reach for Compute Sanitizer first; for performance, Nsight.

What it is¶

cuda-gdb is NVIDIA's port of GDB that debugs host and device code in one session. Its GDB base moved "from GDB 14.2 to 16.3" as of the 13.2 release, so every host-side command you know still works, and it adds a device dimension: a focus (which GPU, which kernel, which block, which thread) and a set of info cuda views into hardware state that no other tool surfaces.¹

It runs in three modes, and they are not equally usable on a server:

Launch under the debugger. cuda-gdb ./app. Fine on a development box, ruinous on a shared node (see below).
Attach to a running process. CUDA-GDB "can attach to and detach from a CUDA application running on GPUs with compute capability 2.0 and beyond, using GDB's built-in commands for attaching to or detaching from a process", so attach <pid> from the prompt. Alternatively set CUDA_DEVICE_WAITS_ON_EXCEPTION=1 and the application "will run normally until a device exception occurs. The application will then wait for CUDA-GDB to attach itself to it for further debugging."⁴ Attach is the interactive mode that makes sense for a long-running job that has already gone wrong.
Post-mortem on a GPU core dump. The process faults, the driver writes a dump, the process dies, and you open the dump later with target cudacore. Nothing stays paused, nothing waits for a human. This is the mode that survives contact with a production cluster.

Remote debugging works through cuda-gdbserver on the target node, with target remote host:port from a client, including cuda-gdbserver :1234 --attach <pid> against an already-running process.⁵

What it is not: it is not a memory checker (that is Compute Sanitizer, which finds out-of-bounds writes that never fault), and it is not a profiler (that is Nsight Systems and Nsight Compute). Reaching for the debugger when a sanitizer run would have named the line is the most common way to waste a day.

Why use it¶

Two properties of CUDA make a normal stack trace worthless, and both are why device-aware debugging exists at all.

Errors are asynchronous, so the reported location is not the fault location. Kernel launches are asynchronous. A kernel that dereferences a bad pointer does not fault at the launch; the error surfaces at the next synchronizing call, which may be a cudaMemcpy in a different Python frame, thousands of instructions later. The traditional fix is CUDA_LAUNCH_BLOCKING=1, which serializes launches so the error is attributed to the launch that caused it. That lever has quietly stopped working for modern serving stacks: once the hot path is captured into CUDA graphs, vLLM reports that "even with CUDA_LAUNCH_BLOCKING=1, we can only see that there's an issue when launching the CUDA graph, but still cannot pinpoint the exact kernel that failed."¹⁵ The whole graph is one launch. Blocking on it tells you the graph broke, which you already knew.

The error is sticky, so there is nothing left to inspect afterwards. An illegal address is not a recoverable exception. NVIDIA's own error table is blunt: cudaErrorIllegalAddress means "The device encountered a load or store instruction on an invalid memory address. This leaves the process in an inconsistent state and any further CUDA work will return the same error. To continue using CUDA, the process must be terminated and relaunched."¹⁷ Every subsequent CUDA call returns the same error, so any handler you write to "print some diagnostics on failure" cannot read device memory, cannot launch a probe kernel, and cannot query anything. The state you need is in a context that is already poisoned. A core dump is taken by the driver at the moment of the exception, before the context is torn down, which is why it can show you what your exception handler cannot.

What you get from a dump is the whole faulting frame: the exception class, the PC, and the focus already set to the offending thread, for example CUDA Exception: Warp Illegal Address with [Current focus set to CUDA kernel 0, grid 9, block (17454,0,0), thread (0,0,0), ...] (the real line continues with the hardware coordinates).¹⁵ From there info cuda threads, backtrace, and print on device variables work as normal, against the dead process.

When to use it (and when not)¶

Symptom	Reach for	Why
Wrong numbers, no crash	Compute Sanitizer `memcheck` / `racecheck`	Silent OOB and shared-memory races never fault, so the debugger has nothing to stop on
`illegal memory access`, reproduces in seconds on one GPU	Compute Sanitizer first, then `cuda-gdb`	The sanitizer names the line without a human in the loop
`illegal memory access`, only at scale, only on some ranks, hours in	GPU core dump	The only mechanism that captures the fault without a human waiting at a breakpoint
Hang with no error	Attach to the process, then `info cuda kernels`	Shows which kernels are resident and where their warps are parked
Slow, but correct	Nsight	The debugger reports no timings
Node-level fault: Xid, ECC, fallen-off-the-bus	Diagnostics and DCGM	Not a software bug; do not debug the application

The hard constraint on a shared server is in one sentence of the manual: "Any GPU hitting a breakpoint will pause all the GPUs running CUDA on that system," and "When any GPU is resumed, all the GPUs are resumed."⁶ On an 8-GPU node with seven other tenants, an interactive breakpoint is a node-wide outage for as long as you sit at the prompt reading a backtrace. The manual gives the containment: "If the CUDA_VISIBLE_DEVICES environment is used, only the specified devices are suspended and resumed."⁶ So the rule for interactive use on shared hardware is: never without CUDA_VISIBLE_DEVICES pinning the debugged process to the GPUs you own, and preferably not at all. Use dumps.

Two more places it does not go:

Core dumps under MPS. The blast radius is the whole MPS server, not the faulting client: "If an MPS client triggers a core dump, every other client running on the same MPS server will fault."⁷ One tenant's bug becomes everyone's, so do not arm dumps on an MPS-shared GPU.
Alongside another tool. "GPU core dump generation is unsupported when other CUDA developer tools, including CUDA-GDB, are interacting with the application, unless explicitly documented as a supported use case (e.g., generate-core-file command)."⁸ You do not get to run the sanitizer and collect a dump in the same process. CUPTI states the same rule from its side: it "still does not support tracing or profiling along with other NVIDIA Developer Tools like cuda-gdb or Compute Sanitizer."¹⁶

Architecture¶

flowchart TB
    subgraph INT["Interactive (a human waits)"]
      APP1["CUDA process"] -->|"ptrace + debug API"| CG["cuda-gdb"]
      CG -->|"breakpoint hit"| PAUSE["ALL GPUs running CUDA<br/>on this host are paused"]
      CG --> FOCUS["focus: device / sm / warp / lane<br/>kernel / block / thread"]
    end
    subgraph PM["Post-mortem (nobody waits)"]
      APP2["CUDA process"] -->|"GPU exception"| DRV["driver writes dump<br/>before context teardown"]
      DRV --> FILE["core_TIME_HOSTNAME_PID.nvcudmp"]
      DRV --> ABRT["abort() unless skip_abort"]
      FILE -->|"target cudacore"| CG2["cuda-gdb (any host, later)"]
      CG2 --> FOCUS
    end

The left path is what most people try first and what makes them unwelcome on a shared cluster. The right path is what you actually deploy: the driver serializes device state to a file at the instant of the exception, and the debugging happens on a workstation, hours later, from an artifact.

How to use it: the focus model¶

Everything device-side in cuda-gdb is relative to a focus, and there are two coordinate systems for the same thread:

Software focus: cuda kernel 0 block (17454,0,0) thread (0,0,0)
Hardware focus: cuda device 0 sm 124 warp 0 lane 0

The views (info cuda devices, info cuda kernels, info cuda blocks, info cuda warps, info cuda lanes, info cuda sms, info cuda barriers) all report relative to the current focus, so a wrong focus produces confidently wrong output.¹¹

The hardware coordinates are not computable from the software ones, and it is worth being blunt about that because it is an inviting mistake. The manual defines them physically: "warpid Warp index inside the current SM" and "laneid Lane index (0-31) inside the current warp."² The SM a block lands on, and the warp slot it occupies once there, are assigned by the hardware scheduler; several blocks are resident on an SM at once, so a block's warps do not start at slot 0. The vLLM walkthrough shows it plainly: the very first thread of the very first block, block (0,0,0), thread (0,0,0), reports device 0, sm 124, warp 0, lane 0.¹⁵ No arithmetic recovers sm 124. Only the debugger knows, and info cuda threads is how you ask.

What the programming model does fix is the lane. A block's threads are linearized, "for a two-dimensional block of size (Dx, Dy), the thread ID of a thread of index (x, y) is (x + y Dx); for a three-dimensional block of size (Dx, Dy, Dz), the thread ID of a thread of index (x, y, z) is (x + y Dx + z Dx Dy)",¹⁸ and warps are cut from that linear order: "The way a block is split into warps is always the same; each warp contains threads of consecutive, increasing thread IDs with the first warp containing thread 0."¹⁹ So lane = tid % 32 is exact, and tid // 32 gives which warp of the block a thread falls in, which is a genuinely useful quantity as long as you never mistake it for the debugger's warp.

The practical consequence: lane == threadIdx.x % 32 is true only for one-dimensional blocks, and quietly false the moment blockDim.x is not a multiple of 32. That is the mistake that has an operator staring at lane 23 while the debugger is telling them about lane 15. The model below computes exactly what is derivable, refuses to compute what is not, and makes the tail-warp case explicit.

"""What you CAN and CANNOT derive about cuda-gdb's focus from (block, thread).

cuda-gdb's hardware focus is (device, sm, warp, lane), where the manual defines
"warpid: Warp index inside the current SM" and "laneid: Lane index (0-31) inside
the current warp". The SM and the warp SLOT are assigned by the hardware scheduler,
so they are NOT functions of (block, thread) and nothing below tries to compute them.
What IS fixed by the programming model:
  - the thread's linear ID in its block: x + y*Dx + z*Dx*Dy
  - its lane: tid % 32, because a block is split into warps of consecutive tids
  - which warp OF THE BLOCK it falls in: tid // 32 (a block-relative index, and NOT
    the debugger's `warp`, which counts warp slots on an SM)
"""
import numpy as np

WARP = 32

def thread_id(x, y, z, bd):
    """Linear thread ID in a block: x + y*Dx + z*Dx*Dy. Works on scalars or arrays."""
    dx, dy, _ = bd
    return x + y * dx + z * dx * dy

def lane(x, y, z, bd):
    """The debugger's `lane`: position within the warp. This one IS derivable."""
    return thread_id(x, y, z, bd) % WARP

def warp_in_block(x, y, z, bd):
    """Which warp of the block. NOT the debugger's `warp` (an SM slot)."""
    return thread_id(x, y, z, bd) // WARP

def block_grid(bd):
    """Every threadIdx in a block, in launch order (x fastest, then y, then z)."""
    dx, dy, dz = bd
    z, y, x = np.meshgrid(np.arange(dz), np.arange(dy), np.arange(dx), indexing="ij")
    return x.ravel(), y.ravel(), z.ravel()

# --- Happy path: 1-D block, where lane really is threadIdx.x % 32 ----------
bd = (256, 1, 1)
assert (warp_in_block(37, 0, 0, bd), lane(37, 0, 0, bd)) == (1, 5)
print("1-D block (256,1,1): thread (37,0,0) -> warp 1 of the block, lane 5")

# --- Adversarial 1: 2-D block whose x extent is not a multiple of 32 -------
# The folk rule "lane == threadIdx.x % 32" breaks, and cuda-gdb will disagree
# with the operator who believes it.
bd = (24, 4, 1)
w, l = warp_in_block(23, 1, 0, bd), lane(23, 1, 0, bd)
folk = 23 % WARP
assert (w, l) == (1, 15)
assert folk == 23 and folk != l
print(f"2-D block (24,4,1): thread (23,1,0) -> warp {w} of the block, lane {l} "
      f"(folk rule threadIdx.x%32 says lane {folk}: wrong)")

# --- Adversarial 2: vectorized scan vs a scalar reference, 6 block shapes --
for bd in [(256, 1, 1), (24, 4, 1), (32, 8, 4), (100, 1, 1), (7, 5, 3), (1024, 1, 1)]:
    dx, dy, dz = bd
    n = dx * dy * dz
    xs, ys, zs = block_grid(bd)
    warps, lanes = warp_in_block(xs, ys, zs, bd), lane(xs, ys, zs, bd)

    assert np.array_equal(thread_id(xs, ys, zs, bd), np.arange(n))   # launch order
    slots = warps * WARP + lanes
    assert len(np.unique(slots)) == n              # no two threads share a warp/lane slot
    assert warps.max() + 1 == -(-n // WARP)        # the block spans ceil(n/32) warps
    ref = [(int(t) // WARP, int(t) % WARP) for t in range(n)]        # slow reference
    assert ref == list(zip(warps.tolist(), lanes.tolist()))
print("launch order, lane bijection, warp count and scalar equivalence hold for 6 block shapes")

# --- Adversarial 3: the tail warp is where an unguarded kernel faults ------
# `if (i < n)` omitted; n = 1000 elements, blockDim.x = 256 -> 1024 threads launched.
n, bdx = 1000, 256
blocks = -(-n // bdx)
assert blocks == 4
last = blocks - 1
oob = np.array([t for t in range(bdx) if last * bdx + t >= n])
assert len(oob) == 1024 - n == 24
fw = warp_in_block(oob, 0, 0, (bdx, 1, 1))
fl = lane(oob, 0, 0, (bdx, 1, 1))
assert np.unique(fw).tolist() == [7]              # a single warp of the block: its last
assert fl.tolist() == list(range(8, 32))          # lanes 8-31; lanes 0-7 stay in bounds
print(f"unguarded n={n}, blockDim.x={bdx}: {len(oob)} out-of-bounds threads sit in block "
      f"{last}, warp {fw[0]} of that block, lanes {fl[0]}-{fl[-1]} (lanes 0-7 are in bounds)")

# --- Adversarial 4: why the block-relative warp is NOT the debugger's warp -
# Every block has a warp 0. Several blocks are resident on one SM at once, so the
# block-relative index cannot name a physical warp slot: it collides across blocks.
resident = [0, 1, 2]                              # three blocks co-resident on one SM
first_warp_of_each = {b: warp_in_block(0, 0, 0, (256, 1, 1)) for b in resident}
assert set(first_warp_of_each.values()) == {0}    # all three claim "warp 0"
assert len(set(first_warp_of_each.values())) == 1 < len(resident)   # not injective
# Only (block, warp_in_block) identifies a warp; the SM slot is the scheduler's to assign.
keys = {(b, warp_in_block(0, 0, 0, (256, 1, 1))) for b in resident}
assert len(keys) == len(resident)
print(f"{len(resident)} blocks co-resident on one SM each have a 'warp 0': the block-relative "
      f"index is not the debugger's SM warp slot, and no arithmetic can recover it")

Executed output:

1-D block (256,1,1): thread (37,0,0) -> warp 1 of the block, lane 5
2-D block (24,4,1): thread (23,1,0) -> warp 1 of the block, lane 15 (folk rule threadIdx.x%32 says lane 23: wrong)
launch order, lane bijection, warp count and scalar equivalence hold for 6 block shapes
unguarded n=1000, blockDim.x=256: 24 out-of-bounds threads sit in block 3, warp 7 of that block, lanes 8-31 (lanes 0-7 are in bounds)
3 blocks co-resident on one SM each have a 'warp 0': the block-relative index is not the debugger's SM warp slot, and no arithmetic can recover it

Read the third case again, because it is the shape of most real reports. A missing if (i < n) guard does not fault in every thread and it does not fault in a random-looking scatter. It faults in exactly one warp, the last warp of the last block, in the lanes above the remainder, and it leaves the lanes below the remainder perfectly healthy. When a dump lands you in block (3,0,0), thread (232,0,0) with the neighbouring lanes fine, the grid tail is the first place to look, not memory corruption.

The fourth case is the guardrail. Every block has a warp 0, and several blocks share an SM, so the block-relative warp index cannot name a physical warp slot. If you catch yourself computing which sm or warp the debugger "should" be showing, stop: ask it.

How to develop with it: build flags, breakpoints, attach¶

Build for attribution. -G is the full device debug build: it implies -O0, gives you real source stepping, and changes scheduling enough to move or mask the very race you are chasing. -lineinfo keeps optimization and adds the line tables, which is enough to attribute a fault to a source line in a dump. For third-party build systems you rarely control the nvcc line directly, so inject it: vLLM's own guidance is export NVCC_PREPEND_FLAGS='-lineinfo' (or '-G' when you need to step).¹⁵ On a serving stack, prefer -lineinfo; a -G build of a large kernel library is both enormous and slow enough to change the bug's behaviour.

Stop where it matters.

# Break on entry to every kernel the application launches.
(cuda-gdb) set cuda break_on_launch application

# Break in one thread only: the guard against a breakpoint firing 8192 times.
(cuda-gdb) break kernel.cu:185 if blockIdx.x == 17454 && threadIdx.x == 0

# Stop on any CUDA API call that returns an error, not just a fatal one.
(cuda-gdb) set cuda api_failures stop_all

set cuda break_on_launch takes none (the default: "no kernel, application or system"), application ("kernel launched by the user application"), system ("any kernel launched by the driver, such as memset"), or all; system is deprecated and "will be dropped in an upcoming release".¹⁰

set cuda api_failures has five modes, and the two that stop are not interchangeable: hide, ignore (the default, which warns on every fatal failure), stop ("The application is stopped when a CUDA API call returns a fatal error"), ignore_all, and stop_all ("The application is stopped when a CUDA API call returns any error").¹⁰ Be clear about what this buys you: the manual says CUDA-GDB "can automatically check the return code of any driver API or runtime API call."¹⁰ For an asynchronous kernel fault, the call that returns the error is the later synchronizing call, which is precisely the misleading location described above. So api_failures catches a bad API call at the API call; it does not pinpoint the faulting kernel. For that you need autostep, which single-steps a designated region so an exception raised inside it is attributed precisely rather than to the next synchronization point,¹⁰ or a core dump.

Attach to something already running. Attach uses GDB's own attach <pid>, and it needs ptrace permission, which most distributions restrict: "By default on some Linux distributions, the debugger cannot attach to an already running processes due to security settings. In order to enable the attach feature of the CUDA debugger, either cuda-gdb should be launched as root, or /proc/sys/kernel/yama/ptrace_scope should be set to zero, using the following command..."¹² In a container you additionally need the capability; NVIDIA's own mixed-stack debugging walkthrough runs its container as docker run -it --rm --cap-add sys_admin --cap-add sys_ptrace --ipc shareable --net host --gpus all ....¹⁴ On Kubernetes that means a securityContext.capabilities.add: ["SYS_PTRACE"] on a debug pod, which is a privilege escalation your platform policy should be deliberate about granting; see agent sandboxing and isolation for the general posture and container toolkit for how the runtime injects devices.

How to run it in production: GPU core dumps¶

This is the only part of cuda-gdb that belongs on a production node, and it is enabled with environment variables, not a debugger session. The dump is written by the driver when the GPU raises an exception, before the context is torn down.

Enable it and control where it lands:

export CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1          # off by default
export CUDA_COREDUMP_SHOW_PROGRESS=1                # dumps are slow; log the progress
export CUDA_COREDUMP_FILE="/mnt/dumps/cuda_%h.%p.%t"   # %h host, %p pid, %t timestamp
export CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory'

Three things in that block are load-bearing on a real cluster.

Size. A full dump contains device global memory. On an 80 GB GPU running a model that fills it, that is what you get: vLLM's guidance is blunt, "For programs like large model inference that occupy almost all GPU memory, a full coredump is impractical (hundreds of GiB of data)," and the recipe above is theirs.¹⁵ The skip_* flags are documented individually (skip_global_memory "Disables dumping of GPU global and constbank memory segments", skip_local_memory, skip_shared_memory, skip_nonrelocated_elf_images, skip_constbank_memory), with gzip_compress and faulted_contexts_only ("Dump only the contexts with exceptions") available on top.⁸ What you keep is the register state, the PC, and the warp/lane focus, which is what names the bug. What you throw away is the tensor content, which you mostly cannot interpret anyway.

Path. The default is the current working directory, named core_TIME_HOSTNAME_PID.nvcudmp.⁸ In a container that is an ephemeral layer that vanishes with the pod, and the fault is exactly the event that kills the pod. Point CUDA_COREDUMP_FILE at a mounted volume. The %h and %p expansions mean a multi-rank job whose ranks share a host will not collide.⁸

Survival, and a flag not to reach for. By default the driver calls abort() at the end of dump generation.⁸ That is usually what you want, since the context is poisoned anyway per the sticky-error rule above. The skip_abort flag suppresses it so a supervisor can run first, and it is tempting. Do not use it casually: vLLM tested it and reports that "experiments have shown that this feature has a significant bug, which may cause illegal memory access errors on the GPU to be ignored. In such cases, subsequent code may continue to run normally, but the program's memory data might already be corrupted... Therefore, this feature is generally unreliable and not recommended."¹⁵ A silently-corrupted training run is a worse outcome than a dead process. Note also that CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION=0 is documented as "equivalent to CUDA_COREDUMP_GENERATION_FLAGS="skip_abort"", so you can arrive at the same hazard by the back door.⁸

log_only is the lighter option: it writes no corefile and only logs the exception details (PC, stack trace, kind). It still requires dump generation to be enabled, so it does not escape the standing cost below.

Cost. This is not free even when nothing faults, and the manual says so directly: "Enabling core dump generation can impact application performance even if no exception is encountered."⁸ The mechanism is visible in the flag list: no_errbar_at_exit exists to "Disable the implicit 'make errors visible at exit' mode used by CUDA coredump",⁸ and the driver API's equivalent explains the trade, that by default "the GPU will ensure memory faults and other errors prevent warps from exiting, if possible", and setting the flag makes "it possible for faulted warps to exit" while "avoiding the potential performance hit".⁹ Disabling the barrier is not free either: the manual warns that if a warp exits before the exception is recognized, "it may result in the exception not being reported", so you can trade your dump away to buy the performance back.⁸ vLLM's conclusion is the practical one: "Enabling CUDA core dump does have some performance impact on CUDA kernels (since it needs to check for errors and attribute them when GPU threads exit). Therefore, it is not advisable to enable CUDA core dump in production environments."¹⁵

The workable posture is therefore: arm it on the retry, not on the fleet. When a job dies with an illegal address, relaunch that job (or that one rank) with the dump variables set and let it fail again. If the failure is rare enough that you cannot afford to wait for a second occurrence, arm it on a canary replica only, and accept the overhead there.

Everything above is also reachable programmatically through cuCoredumpSetAttribute and cuCoredumpSetAttributeGlobal, which is how a framework arms dumps for itself. Note the precedence rule: environment variables set before CUDA initialization take permanent precedence over the programmatic settings, so an operator's CUDA_COREDUMP_* cannot be overridden by application code.⁹

Open the result anywhere, on any host with the toolkit:

cuda-gdb
(cuda-gdb) target cudacore /mnt/dumps/cuda_node17.31337.1752... 
# focus is already on the faulting thread
(cuda-gdb) info cuda kernels
(cuda-gdb) backtrace
(cuda-gdb) info cuda lanes

To correlate host and device, capture both: a CPU core dump is generated by default when GPU core dump generation is enabled, and the two load together with target core core.cpu core.cuda.⁸

How to maintain it¶

Keep -lineinfo in release builds of your own kernels. It does not disable optimization and it is the difference between a dump that names attention.cu:412 and a dump that names a PC. Kernels shipped without line tables produce dumps that only NVIDIA can read.
Track what the toolkit has actually removed, not just deprecated. Kepler support was removed toolkit-wide in CUDA 12.0. CUDA 13.0's own cuFFT deprecation notes record "Removed support for Maxwell, Pascal, and Volta GPUs, corresponding to compute capabilities earlier than Turing", and cuSPARSE's independently record the same drop for the same three architectures; these are library-level deprecation notes rather than one general-toolkit announcement, but they agree with each other, so on a current toolkit those architectures are functionally gone rather than merely deprecated.³ For the dump variables: CUDA_ENABLE_LIGHTWEIGHT_COREDUMP was deprecated in 12.5 and its "Support ... has been removed" in the 13.0 release, so a runbook still setting it is doing nothing. CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION is deprecated but still functional: setting it to 0 is documented as equivalent to skip_abort, which is a hazard rather than a no-op (above).⁸
Do not leave dump generation armed by default in a Helm chart or a base image. It carries a standing cost even when nothing faults (above), and it is exactly the kind of setting that gets baked in during an incident and never removed. log_only reduces what you store, not what you pay, so it is not a free always-on setting either.
Watchpoints are not available on device code: "Watchpoints on CUDA code are not supported. Watchpoints on host code are supported."¹³ Do not build a procedure around them.

Failure modes¶

Symptom	Cause	Fix
The whole node freezes when you hit a breakpoint	A breakpoint pauses every GPU running CUDA on the host⁶	Pin the debugged process with `CUDA_VISIBLE_DEVICES`, which limits suspend and resume to those devices; better, use a dump
No dump is written after an OOM	An allocation failure is not a device exception; it is non-sticky and does not trigger a dump¹⁵	Debug OOM with the allocator and memory tooling, not the debugger
No dump appears, no error either	The process could not write to the CWD, or another CUDA tool was attached (sanitizer, Nsight), which makes dump generation unsupported⁸	Set `CUDA_COREDUMP_FILE` to a writable volume; run one tool at a time
The dump is hundreds of gigabytes	Default generation includes global memory¹⁵	The `skip_*` flag set above; add `gzip_compress`
The process vanished before a supervisor could react	The driver calls `abort()` at the end of dump generation⁸	`skip_abort` suppresses it, but vLLM reports it can cause illegal-access errors "to be ignored" and calls it "unreliable and not recommended"¹⁵; prefer letting the process die
Every MPS client on the node faulted at once	"If an MPS client triggers a core dump, every other client running on the same MPS server will fault."⁷	Do not arm dumps under MPS; isolate the suspect workload first
Backtrace shows PCs, no source	Built without `-lineinfo` or `-G`	`NVCC_PREPEND_FLAGS='-lineinfo'` and rebuild¹⁵
`attach <pid>` refuses to attach	Yama `ptrace_scope`, or a container without the capability	`ptrace_scope=0` or root; in containers `--cap-add sys_ptrace`¹²¹⁴
Kernel throughput dropped after an incident	Coredump generation left armed: "Enabling core dump generation can impact application performance even if no exception is encountered."⁸	Unset `CUDA_ENABLE_COREDUMP_ON_EXCEPTION`. `log_only` does not help; it cuts what you store, not the cost. `no_errbar_at_exit` cuts the cost but can lose the exception⁸
Debugger reports a lane you did not expect	`lane == threadIdx.x % 32` is false for multi-dimensional blocks (executed model above)	Linearize first: `tid = x + yDx + zDx*Dy`, then `lane = tid % 32`
You computed which `sm`/`warp` the debugger should show, and it disagrees	`warpid` is the "Warp index inside the current SM", a scheduler-assigned slot, not a block-relative index²	It is not derivable. Read it from `info cuda threads`

References¶

NVIDIA, CUDA-GDB 13.3 documentation (focus commands, info cuda views, attach, cuda-gdbserver, multi-GPU suspend semantics, GPU core dump support, known issues): https://docs.nvidia.com/cuda/cuda-gdb/index.html
NVIDIA, CUDA-GDB product page: https://developer.nvidia.com/cuda-gdb
NVIDIA, CUDA Driver API: Coredump Attributes Control API (cuCoredumpSetAttribute, CUcoredumpSettings, CUcoredumpGenerationFlags): https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__COREDUMP.html
NVIDIA, CUDA Runtime API: Data types (cudaErrorIllegalAddress and the sticky-error contract): https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html
NVIDIA, CUDA C++ Programming Guide (thread hierarchy: linearized thread ID): https://docs.nvidia.com/cuda/cuda-c-programming-guide/
NVIDIA, PTX ISA 9.3, §3.1 A Set of SIMT Multiprocessors (warp partitioning of consecutive thread IDs): https://docs.nvidia.com/cuda/parallel-thread-execution/
Kaichao You, CUDA Core Dump: An Effective Tool to Debug Memory Access Issues and Beyond, vLLM blog, 2025-08-11 (the production recipe, coredump size on large models, CUDA-graph limitation of CUDA_LAUNCH_BLOCKING, NVCC_PREPEND_FLAGS): https://vllm.ai/blog/cuda-debugging
Peter Entschev and Ben Zaitlen, Debugging a Mixed Python and C Language Stack, NVIDIA Technical Blog, 2023-04-20 (debug-capable container flags): https://developer.nvidia.com/blog/debugging-mixed-python-and-c-language-stack/

NVIDIA, CUDA-GDB 13.3 documentation. The release notes record the GDB base moving "from GDB 14.2 to 16.3" in the 13.2 release, with no further change noted for 13.3; the statement that 13.3 sits on GDB 16.3 is therefore an inference from that entry rather than a sentence the manual states outright. https://docs.nvidia.com/cuda/cuda-gdb/index.html ↩
NVIDIA, CUDA-GDB 13.3 manual, physical (hardware) coordinates: "laneid Lane index (0-31) inside the current warp." and "warpid Warp index inside the current SM." (also nwarpid, "Number of warps in the current SM"). Read from the CUDA-GDB PDF, Release 13.3. https://docs.nvidia.com/cuda/pdf/cuda-gdb.pdf ↩↩
NVIDIA, CUDA Toolkit release notes: "Kepler architecture support is removed from CUDA 12.0" (12.0 release notes, "General CUDA" section, toolkit-wide); "Removed support for Maxwell, Pascal, and Volta GPUs, corresponding to compute capabilities earlier than Turing" (13.0 release notes, under the cuFFT: Release 13.0 Deprecations subsection specifically, not a general-toolkit statement); the cuSPARSE section of the same 13.0 release notes independently records "Dropped support for pre-Turing architectures (Maxwell, Volta, and Pascal)." The CUDA-GDB 13.0 release notes separately record "Support for CUDA_ENABLE_LIGHTWEIGHT_COREDUMP has been removed." https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html · https://docs.nvidia.com/cuda/archive/12.0.0/cuda-toolkit-release-notes/index.html ↩
NVIDIA, CUDA-GDB: "CUDA-GDB can attach to and detach from a CUDA application running on GPUs with compute capability 2.0 and beyond, using GDB's built-in commands for attaching to or detaching from a process." With CUDA_DEVICE_WAITS_ON_EXCEPTION=1, "the application will run normally until a device exception occurs. The application will then wait for CUDA-GDB to attach itself to it for further debugging." Not supported on WSL. https://docs.nvidia.com/cuda/cuda-gdb/index.html ↩
NVIDIA, CUDA-GDB: remote debugging via cuda-gdbserver, including cuda-gdbserver :1234 --attach <pid> and target remote. https://docs.nvidia.com/cuda/cuda-gdb/index.html ↩
NVIDIA, CUDA-GDB, Multi-GPU Debugging: "Any GPU hitting a breakpoint will pause all the GPUs running CUDA on that system." "When any GPU is resumed, all the GPUs are resumed." "If the CUDA_VISIBLE_DEVICES environment is used, only the specified devices are suspended and resumed." https://docs.nvidia.com/cuda/cuda-gdb/index.html ↩↩↩
NVIDIA, CUDA-GDB: "If an MPS client triggers a core dump, every other client running on the same MPS server will fault." https://docs.nvidia.com/cuda/cuda-gdb/index.html ↩↩
NVIDIA, CUDA-GDB 13.3, GPU core dump support (read from the PDF): CUDA_ENABLE_COREDUMP_ON_EXCEPTION ("This option is disabled by default"), CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION, CUDA_COREDUMP_SHOW_PROGRESS, CUDA_COREDUMP_FILE, CUDA_ENABLE_USER_TRIGGERED_COREDUMP, CUDA_COREDUMP_PIPE. "Note: Enabling core dump generation can impact application performance even if no exception is encountered." Table 1 lists ten comma-delimited CUDA_COREDUMP_GENERATION_FLAGS: skip_nonrelocated_elf_images, skip_global_memory ("Disables dumping of GPU global and constbank memory segments"), skip_shared_memory, skip_local_memory, skip_abort ("Disables calling abort() at the end of the GPU core dump generation process"), skip_constbank_memory, gzip_compress, faulted_contexts_only ("Dump only the contexts with exceptions"), no_errbar_at_exit ("Disable the implicit 'make errors visible at exit' mode used by CUDA coredump"), and log_only ("Do not generate a corefile, only log the exception information (PC, stack trace, kind, etc.)"). "Note: Setting the CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION environment variable to 0 is equivalent to CUDA_COREDUMP_GENERATION_FLAGS=\"skip_abort\"." Default file name core_TIME_HOSTNAME_PID.nvcudmp in the current working directory; loaded with target cudacore <file> or target core core.cpu core.cuda. On losing an exception when warps exit early: "If the warp exits after the instruction causing exception has executed, but before the exception has been recognized and reported, it may result in the exception not being reported." "GPU core dump generation is unsupported when other CUDA developer tools, including CUDA-GDB, are interacting with the application, unless explicitly documented as a supported use case (e.g., generate-core-file command)." The CUDA-GDB 13.0 release notes record "Support for CUDA_ENABLE_LIGHTWEIGHT_COREDUMP has been removed." https://docs.nvidia.com/cuda/pdf/cuda-gdb.pdf · https://docs.nvidia.com/cuda/cuda-gdb/index.html · Note one inconsistency between two NVIDIA sources: the driver API documents the CU_COREDUMP_FILE default as core.cuda.HOSTNAME.PID, while the CUDA-GDB guide documents core_TIME_HOSTNAME_PID.nvcudmp. Do not depend on either; set CUDA_COREDUMP_FILE explicitly. ↩↩↩↩↩↩↩↩↩↩↩↩↩↩↩
NVIDIA, CUDA Driver API, Coredump Attributes Control API: cuCoredumpSetAttribute / cuCoredumpSetAttributeGlobal, and the note that environment variables set before CUDA initialization take permanent precedence over programmatic settings. On CU_COREDUMP_NO_ERRBAR_AT_EXIT: by default "the GPU will ensure memory faults and other errors prevent warps from exiting, if possible. This can potentially affect the performance of the application. Setting this flag will disable this functionality, making it possible for faulted warps to exit, but also avoiding the potential performance hit." Note the qualifier: it is faulted warps that are held, not every warp. CU_COREDUMP_LIGHTWEIGHT and CU_COREDUMP_TRIGGER_HOST are deprecated as of CUDA 12.5. https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__COREDUMP.html ↩↩
NVIDIA, CUDA-GDB 13.3 manual (read from the PDF). set cuda break_on_launch possible options: "none | no kernel, application or system (default)", "application | kernel launched by the user application", "system | any kernel launched by the driver, such as memset", "all | any kernel, application and system"; the release notes add "The break_on_launch system option is deprecated. It will be dropped in an upcoming release." set cuda api_failures: "CUDA-GDB can automatically check the return code of any driver API or runtime API call... Five modes are supported: hide (CUDA API call failures are not reported); ignore (Warning message is printed for every fatal CUDA API call failure) (default); stop (The application is stopped when a CUDA API call returns a fatal error); ignore_all (Warning message is printed for every CUDA API call failure); stop_all (The application is stopped when a CUDA API call returns any error)." The autostep command single-steps a region automatically to increase the precision of CUDA exception attribution. https://docs.nvidia.com/cuda/pdf/cuda-gdb.pdf ↩↩↩↩
NVIDIA, CUDA-GDB: focus commands (cuda device/sm/warp/lane, cuda kernel/block/thread) and the info cuda devices|kernels|blocks|threads|warps|lanes|sms|barriers|contexts|managed views. https://docs.nvidia.com/cuda/cuda-gdb/index.html ↩
NVIDIA, CUDA-GDB: "By default on some Linux distributions, the debugger cannot attach to an already running processes due to security settings. In order to enable the attach feature of the CUDA debugger, either cuda-gdb should be launched as root, or /proc/sys/kernel/yama/ptrace_scope should be set to zero, using the following command:" https://docs.nvidia.com/cuda/cuda-gdb/index.html ↩↩
NVIDIA, CUDA-GDB: "Watchpoints on CUDA code are not supported. Watchpoints on host code are supported." https://docs.nvidia.com/cuda/cuda-gdb/index.html ↩
Peter Entschev and Ben Zaitlen, Debugging a Mixed Python and C Language Stack, NVIDIA Technical Blog, 2023-04-20: the debug container is launched with docker run -it --rm --cap-add sys_admin --cap-add sys_ptrace --ipc shareable --net host --gpus all ... (the post's image tag is a 2019-era RAPIDS build; the capability flags are the transferable part). https://developer.nvidia.com/blog/debugging-mixed-python-and-c-language-stack/ ↩↩
Kaichao You, CUDA Core Dump: An Effective Tool to Debug Memory Access Issues and Beyond, vLLM blog, 2025-08-11: the CUDA_ENABLE_COREDUMP_ON_EXCEPTION / CUDA_COREDUMP_SHOW_PROGRESS / CUDA_COREDUMP_GENERATION_FLAGS / CUDA_COREDUMP_FILE recipe; "For programs like large model inference that occupy almost all GPU memory, a full coredump is impractical (hundreds of GiB of data)."; "even with CUDA_LAUNCH_BLOCKING=1, we can only see that there's an issue when launching the CUDA graph, but still cannot pinpoint the exact kernel that failed."; export NVCC_PREPEND_FLAGS='-lineinfo'; "Enabling CUDA core dump does have some performance impact on CUDA kernels (since it needs to check for errors and attribute them when GPU threads exit). Therefore, it is not advisable to enable CUDA core dump in production environments."; the skip_abort warning ("experiments have shown that this feature has a significant bug, which may cause illegal memory access errors on the GPU to be ignored... Therefore, this feature is generally unreliable and not recommended"); the observed focus line [Current focus set to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 124, warp 0, lane 0]; out-of-memory errors from allocation are non-sticky and do not trigger a core dump. https://vllm.ai/blog/cuda-debugging ↩↩↩↩↩↩↩↩↩↩↩
NVIDIA, CUPTI documentation, Multiple Subscribers: "CUPTI still does not support tracing or profiling along with other NVIDIA Developer Tools like cuda-gdb or Compute Sanitizer." https://docs.nvidia.com/cupti/main/main.html ↩
NVIDIA, CUDA Runtime API, cudaErrorIllegalAddress: "The device encountered a load or store instruction on an invalid memory address. This leaves the process in an inconsistent state and any further CUDA work will return the same error. To continue using CUDA, the process must be terminated and relaunched." https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html ↩
NVIDIA, CUDA C++ Programming Guide, Thread Hierarchy: "For a one-dimensional block, they are the same; for a two-dimensional block of size (Dx, Dy), the thread ID of a thread of index (x, y) is (x + y Dx); for a three-dimensional block of size (Dx, Dy, Dz), the thread ID of a thread of index (x, y, z) is (x + y Dx + z Dx Dy)." https://docs.nvidia.com/cuda/cuda-c-programming-guide/ ↩
NVIDIA, PTX ISA 9.3, §3.1: "The way a block is split into warps is always the same; each warp contains threads of consecutive, increasing thread IDs with the first warp containing thread 0." https://docs.nvidia.com/cuda/parallel-thread-execution/ ↩