Cómo usar
- Copia el comando y revisa el paquete antes de instalarlo en tu entorno de agente.
- Consulta las herramientas, credenciales y entornos necesarios en el documento original.
- Describe las entradas y el resultado deseado; contrasta la salida con las fuentes.
npx skills add https://github.com/mohitmishra786/low-level-dev-skills --skill cuda-debuggingAntes de empezar
- La instalación no incluye suscripción al modelo, cuota API ni permisos sobre tus archivos.
- Las estrellas e instalaciones reflejan actividad, no una prueba independiente de calidad o seguridad.
Documento original del Skill
Contenido del repositorio de origen con títulos, ejemplos, código, tablas, enlaces e imágenes preservados.
CUDA Debugging
Purpose
Guide agents through debugging CUDA programs with cuda-gdb for interactive GPU thread inspection, NVIDIA Compute Sanitizer for automated memory and race detection, GPU core dump analysis, device-side printf, and triaging common CUDA runtime error codes.
When to Use
cudaErrorIllegalAddress(700) or segmentation fault on device- Intermittent correctness failures in multi-threaded GPU code
- Debugging race conditions between warps or between host and device
- Stepping through kernel code line-by-line with cuda-gdb
- Validating uninitialized memory reads with initcheck
- Kernel hang or
cudaErrorLaunchTimeout(702)
Workflow
1. Build for debugging
bash
# Debug build — disables optimizations, enables device debug
nvcc -G -g -O0 -arch=sm_80 -o app_debug main.cu
# Sanitizer-friendly build (lineinfo helps reports)
nvcc -lineinfo -g -O2 -arch=sm_80 -o app_san main.cu-G is required for cuda-gdb source-level stepping. Sanitizers work with optimized builds but -G gives clearer line numbers.
2. Compute Sanitizer — automated checks
bash
# Memory errors (OOB, misaligned, leak)
compute-sanitizer --tool memcheck ./app_san
# Shared memory and global memory races
compute-sanitizer --tool racecheck ./app_san
# Uninitialized memory reads
compute-sanitizer --tool initcheck ./app_san
# Synchronization errors (missing __syncthreads)
compute-sanitizer --tool synccheck ./app_san
# Verbose with source correlation
compute-sanitizer --tool memcheck --show-reachable=yes --log-file san.log ./app_sanTypical memcheck output:
======== Invalid __global__ write of size 4
======== at 0x1a0 in vector_add(vector_add.cu:12)
======== by thread (0,0,0) in block (0,0,0)
======== Address 0x7f... is out of bounds3. cuda-gdb interactive session
bash
# Launch under cuda-gdb
cuda-gdb ./app_debug
# Or attach to running process
cuda-gdb -p <pid>Essential commands:
gdb
# Break at kernel entry
(cuda-gdb) break vector_add
(cuda-gdb) run
# Focus on GPU threads
(cuda-gdb) info cuda kernels
(cuda-gdb) cuda kernel 0
(cuda-gdb) cuda thread (0,0,0) # block (x,y,z), thread (x,y,z)
# Inspect device memory
(cuda-gdb) print data[i]
(cuda-gdb) x/10f d_ptr
# Step in kernel
(cuda-gdb) cuda step
(cuda-gdb) cuda next
# All threads in block
(cuda-gdb) info cuda threads
(cuda-gdb) cuda thread (0,0,5)4. Device printf
c
__global__ void debug_kernel(float *data, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
if (i < 5) // limit output
printf("thread %d: data[%d] = %f\n", i, i, data[i]);
data[i] *= 2.0f;
}
}bash
# Buffer size for printf (default may truncate)
cuda-gdb) set cuda printf_buffer_size 16777216Flush with cudaDeviceSynchronize() before checking output. Excessive printf from all threads will overwhelm the buffer.
5. Error code triage
| Code | Name | Common cause |
|---|---|---|
| 700 | cudaErrorIllegalAddress | OOB access, use-after-free, bad pointer |
| 701 | cudaErrorLaunchOutOfResources | Too much shared mem or registers per block |
| 702 | cudaErrorLaunchTimeout | Infinite loop, TDR watchdog (Windows/default Linux) |
| 719 | cudaErrorLaunchFailure | Assert in kernel, stack overflow |
c
// Always check after launch
kernel<<<grid, block>>>(args);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess)
fprintf(stderr, "launch: %s\n", cudaGetErrorString(err));
cudaDeviceSynchronize();
err = cudaGetLastError();
if (err != cudaSuccess)
fprintf(stderr, "exec: %s\n", cudaGetErrorString(err));6. GPU core dumps
bash
# Enable coredump (driver 450+)
export CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1
export CUDA_COREDUMP_FILE=/tmp/cuda_coredump_%h.%p
./app_san # crash generates dump
# Analyze with cuda-gdb
cuda-gdb ./app_san /tmp/cuda_coredump_hostname.pid
(cuda-gdb) cuda coredump load /tmp/cuda_coredump_hostname.pid
(cuda-gdb) bt
(cuda-gdb) info cuda kernels7. Debugging decision tree
Crash or wrong results?
├── Consistent wrong values → logic bug; use printf or cuda-gdb
├── Intermittent / depends on size → OOB or race
│ ├── compute-sanitizer --tool memcheck
│ └── compute-sanitizer --tool racecheck
├── Hang / timeout 702 → infinite loop or barrier mismatch
│ └── synccheck; audit __syncthreads paths
└── Works in debug (-G), fails in release → uninitialized mem or race
└── initcheck + racecheck on release build8. Multi-GPU and MIG notes
bash
# Isolate GPU
CUDA_VISIBLE_DEVICES=0 compute-sanitizer --tool memcheck ./app
# MIG instances appear as separate devices
nvidia-smi -LCommon Problems
| Symptom | Cause | Fix |
|---|---|---|
| cuda-gdb can't break in kernel | Built without -G | Rebuild with nvcc -G -g -O0 |
| Sanitizer reports no errors but crash persists | Async error delayed | Add cudaDeviceSynchronize() after kernel |
printf shows nothing | Buffer full or no sync | Limit prints; increase buffer; sync |
| racecheck false positive on atomics | Non-atomic RMW | Use atomicAdd/atomicCAS |
| Attach fails | Process not in CUDA context | Break after first cudaMalloc |
| TDR timeout on Windows | Long-running kernel | Split kernel; cudaDeviceSetLimit or regedit TDR |
Related Skills
skills/gpu/cuda— kernel patterns, memory hierarchy, launch configskills/gpu/cuda-profiling— performance after correctness is verifiedskills/gpu/gpu-memory-model— understanding races and coalescingskills/debuggers/gdb— host-side GDB commands shared with cuda-gdbskills/runtimes/sanitizers— ASan/TSan concepts for host codeskills/kernel/kernel-debugging— kgdb/kprobes for driver-level issues
