CUDA Cheatsheet

Error Handling

Use this CUDA reference while you build software engineering projects, review code for technical interview prep, or polish examples for a software engineer resume.

Error Model

Every CUDA Runtime API function returns cudaError_t. Kernel launches do not return errors directly — check with cudaPeekAtLastError() or cudaGetLastError() immediately after the launch.

cudaError_t err = cudaMalloc(&d_ptr, bytes);
if (err != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed: %s\n", cudaGetErrorString(err));
    exit(EXIT_FAILURE);
}

`cudaError_t` — Common Values

Code	Constant	Meaning
0	`cudaSuccess`	No error
1	`cudaErrorInvalidValue`	Invalid argument
2	`cudaErrorMemoryAllocation`	Out of device memory
35	`cudaErrorInsufficientDriver`	Driver too old
98	`cudaErrorInvalidDeviceFunction`	Kernel not compiled for device
100	`cudaErrorNoDevice`	No CUDA-capable device found
101	`cudaErrorInvalidDevice`	Device index out of range
209	`cudaErrorNoKernelImageForDevice`	No cubin/PTX image for this SM
217	`cudaErrorPeerAccessUnsupported`	P2P not supported
700	`cudaErrorIllegalAddress`	Global memory out-of-bounds access
701	`cudaErrorLaunchOutOfResources`	Too many registers or shared mem
702	`cudaErrorLaunchTimeout`	Kernel exceeded GPU watchdog timer
710	`cudaErrorAssert`	Device-side `assert` fired
715	`cudaErrorIllegalInstruction`	Invalid instruction executed
719	`cudaErrorLaunchFailure`	Unspecified kernel launch error
999	`cudaErrorUnknown`	Unknown internal error

Error String Helpers

const char* cudaGetErrorString(cudaError_t err);   // human-readable message
const char* cudaGetErrorName(cudaError_t err);     // enum name as string

// Example:
cudaError_t e = cudaMemcpy(/*...*/);
if (e) fprintf(stderr, "[%s] %s\n", cudaGetErrorName(e), cudaGetErrorString(e));

Kernel Launch Error Checking

Kernel launches return void. Check with cudaGetLastError() (clears the error) or cudaPeekAtLastError() (does not clear).

myKernel<<<grid, block>>>(args);

// Pattern 1: check + clear immediately
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess)
    fprintf(stderr, "Kernel launch error: %s\n", cudaGetErrorString(err));

// Pattern 2: also check async execution errors after sync
cudaDeviceSynchronize();
err = cudaGetLastError();
if (err != cudaSuccess)
    fprintf(stderr, "Kernel execution error: %s\n", cudaGetErrorString(err));

CUDA_CHECK Macro (standard pattern)

#define CUDA_CHECK(call)                                                      \
    do {                                                                      \
        cudaError_t _err = (call);                                            \
        if (_err != cudaSuccess) {                                            \
            fprintf(stderr, "CUDA error at %s:%d — %s (%s)\n",              \
                    __FILE__, __LINE__,                                       \
                    cudaGetErrorString(_err), cudaGetErrorName(_err));        \
            exit(EXIT_FAILURE);                                               \
        }                                                                     \
    } while (0)

// Usage
CUDA_CHECK(cudaMalloc(&d_ptr, bytes));
CUDA_CHECK(cudaMemcpy(d_ptr, h_ptr, bytes, cudaMemcpyHostToDevice));

myKernel<<<grid, block>>>(args);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());

Driver API Errors (`CUresult`)

When mixing runtime and driver API, errors are CUresult (prefix CUDA_ERROR_).

#include <cuda.h>

CUresult res = cuInit(0);
if (res != CUDA_SUCCESS) {
    const char* msg;
    cuGetErrorString(res, &msg);
    const char* name;
    cuGetErrorName(res, &name);
    fprintf(stderr, "[%s] %s\n", name, msg);
}

CUDA-GDB — Debugging on GPU

# Compile with debug info
nvcc -g -G myfile.cu -o myapp   # -G keeps device debug info

# Launch under cuda-gdb
cuda-gdb ./myapp

# Useful commands inside cuda-gdb:
# cuda thread (blockIdx.x, threadIdx.x)   — switch focus
# cuda kernel / block / warp / lane       — inspect hierarchy
# info cuda threads                        — list active threads
# info cuda kernels                        — list running kernels
# print threadIdx / blockIdx               — built-in vars

`cuda-memcheck` / Compute Sanitizer

# Legacy (CUDA < 11.x)
cuda-memcheck ./myapp

# Modern (CUDA 11+)
compute-sanitizer ./myapp                     # default: memcheck
compute-sanitizer --tool memcheck ./myapp     # out-of-bounds, use-after-free
compute-sanitizer --tool racecheck ./myapp    # shared memory race conditions
compute-sanitizer --tool initcheck ./myapp    # uninitialized device memory reads
compute-sanitizer --tool synccheck ./myapp    # invalid __syncthreads usage

Common output patterns:

CUDA Error: Global memory read of size 4 bytes out of bounds
 at 0x... in myKernel(float*, int) [src/main.cu:42]

NVTX — Profiler Annotations

Mark regions for Nsight Systems / Nvtx.

#include <nvtx3/nvToolsExt.h>

nvtxRangePushA("Data upload");
cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, stream);
nvtxRangePop();

nvtxRangePushA("Kernel");
myKernel<<<grid, block, 0, stream>>>(d_in, d_out, n);
nvtxRangePop();

Link with -lnvToolsExt.

Assertions in Device Code

#include <cassert>

__global__ void safeKernel(float* data, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    assert(i < n && "index out of range");
    data[i] = sqrtf(data[i]);
}
// Compiles out with -DNDEBUG

When assertion fires: the kernel terminates and the next sync/API call returns cudaErrorAssert (error 710) — a sticky error.

Handling OOM Gracefully

Query available memory first with cudaMemGetInfo, and treat cudaErrorMemoryAllocation as recoverable.

size_t freeBytes, totalBytes;
cudaMemGetInfo(&freeBytes, &totalBytes);   // current free / total device memory
printf("Free: %zu MB of %zu MB\n", freeBytes >> 20, totalBytes >> 20);

float* d_ptr = nullptr;
cudaError_t err = cudaMalloc(&d_ptr, hugeBytes);
if (err == cudaErrorMemoryAllocation) {
    cudaGetLastError();   // clear the error
    // Try smaller allocation, use host, etc.
} else {
    CUDA_CHECK(err);
}

Sticky vs. Non-Sticky Errors

Sticky errors (cudaErrorLaunchFailure, cudaErrorIllegalAddress, etc.): contaminate the CUDA context — all subsequent API calls return the same error. Only cudaDeviceReset() clears them.
Non-sticky errors (e.g., cudaErrorInvalidValue): individual call fails; context is still usable. cudaGetLastError() clears them.

// After a sticky error:
cudaDeviceReset();   // destroys and resets the entire context

`cudaGetLastError` vs `cudaPeekAtLastError`

Function	Clears error?	Use case
`cudaGetLastError()`	Yes	Check and consume the error
`cudaPeekAtLastError()`	No	Read error without clearing (e.g., in a loop)

// In a loop: peek to inspect, then get to clear at the end
for (...) {
    myKernel<<<g,b>>>(args);
    if (cudaPeekAtLastError() != cudaSuccess) break;
}
cudaError_t final = cudaGetLastError();  // clear

Error Propagation in Libraries

When writing a reusable library, return errors instead of exiting:

cudaError_t myLibFunc(float* in, float* out, int n) {
    cudaError_t err;

    float* tmp;
    err = cudaMalloc(&tmp, n * sizeof(float));
    if (err != cudaSuccess) return err;

    myKernel<<<(n+255)/256, 256>>>(in, tmp, n);
    err = cudaGetLastError();
    if (err != cudaSuccess) { cudaFree(tmp); return err; }

    err = cudaMemcpy(out, tmp, n * sizeof(float), cudaMemcpyDeviceToDevice);
    cudaFree(tmp);
    return err;
}