CUDA Cheatsheet

Error Handling

Use this CUDA reference while you build software engineering projects, review code, or refresh the syntax you reach for most.

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;
}