What is CUDA error checking?

Every runtime call returns a status and every kernel launch fails silently unless you ask, which is what cudaGetLastError and a check macro are for.

There are two halves and they fail differently. Every function in the CUDA runtime API hands back a cudaError_t, and C does not warn on a discarded return value, so a program that throws all of them away compiles clean. The launch is the exception that matters: "Kernel launches using triple chevron notation do not return a cudaError_t." There is no value to ignore, so the runtime keeps a status of its own, one per host thread, overwritten by each failure. cudaGetLastError reads it and resets it to cudaSuccess; cudaPeekAtLastError reads it and leaves it set. Take it when you are handling the failure, because a status left set makes the next unrelated check fire on something that already happened. Peek when you only want to look, which is what a logger wants.

The other half is distance. Work you queue runs after the launch statement has returned, so a fault inside a kernel has nowhere to be reported at the moment it happens, and it goes to whichever runtime call looks next. PyTorch ships this inside its own CUDA error messages: "CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect." A framework with millions of users warns that its own stack traces point at the wrong line, and the cause is in the runtime rather than in PyTorch. That is why every launch on this site carries two checks and not one. cudaGetLastError() catches a launch the driver refused outright, such as a block size over 1024, immediately. cudaDeviceSynchronize() waits and reports what the kernel then did.

Some failures end the context and some do not, and NVIDIA marks the split in the reference rather than in prose. Eleven of the 137 cudaError enumerators carry the sentence "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." cudaErrorIllegalAddress (700) and cudaErrorLaunchFailure (719) are the two you meet first. cudaErrorInvalidValue (1) and cudaErrorMemoryAllocation (2) are not on that list: the call fails, you handle it, the context is fine. Clearing the status after a sticky one buys nothing, because the next call sets it straight back.

Know where the wall is, because one class of bug never reaches any of this. A write past the end of a buffer only faults if it lands on a page the process does not own, and cudaMalloc rounds allocations up, so a short overrun often returns cudaSuccess and quietly corrupts something. A bounds check guards the index against the element count you passed, never against the size of the allocation. Compute Sanitizer covers the other half, naming the write, the thread and the block whether or not the address happens to be unmapped.

Measured

Day 6 makes two mistakes on purpose and prints which call saw each. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), nvcc -O3 -arch=sm_75, captured 2026-08-30.

A cudaMemcpy to a null destination returned cudaErrorInvalidValue, "invalid argument", on the line that made it. Two cudaPeekAtLastError calls then reported the same code without consuming it, cudaGetLastError reported it once and cleared it, and the peek after that came back cudaSuccess. That is the entire difference between the two readers, in five lines of output. The context survived, which is what non-sticky means: the program went on to run a kernel over 611 elements and every one of them matched the CPU reference.

The second mistake travelled. A kernel wrote past the end of its buffer, and a peek straight after the launch returned cudaSuccess, because the launch is asynchronous and the fault had not happened yet. It surfaced at the next cudaMemcpy, a call whose pointers, size and direction were all correct. Then the context died: after clearing the status with cudaGetLastError, a fresh cudaMalloc returned cudaErrorIllegalAddress and so did both cudaFree calls. Only a new process gets CUDA back, which is why cudaGetLastError is a diagnostic and never a recovery.

Code

From code/day06-error-checking/error_checking.cu. One failed call, read four times, which is the program that produced the five lines above.

    const cudaError_t badCopy =
        cudaMemcpy(nullptr, h_in.data(), bytes, cudaMemcpyHostToDevice);
    const cudaError_t peek1 = cudaPeekAtLastError();
    const cudaError_t peek2 = cudaPeekAtLastError();
    const cudaError_t taken = cudaGetLastError();
    const cudaError_t peek3 = cudaPeekAtLastError();

A null destination, not a reversed cudaMemcpyKind. Day 6 tried the obvious version first and it does not work: on this card a host-to-device copy tagged cudaMemcpyDeviceToHost returned cudaSuccess and copied correctly, because unified addressing lets the runtime read the real direction off the pointer attributes and treat the kind argument as a hint it can overrule. That run is recorded at the end of the day 6 evidence file, under the superseded-first-run heading. The direction flag will not catch your mistake for you.

Related terms

Where you meet this

Day 5, vector addition, where the CUDA_CHECK macro arrives and you are told to wrap everything.
Day 6, how to check for errors in CUDA, the lesson that owns this term.
Day 8, bounds checks and grid-stride loops, for the failure no error check can report: an element nobody wrote.
an illegal memory access was encountered, the sticky one, ranked by cause.
invalid argument, the non-sticky one you can handle and carry on from.
unspecified launch failure, 700 with less information attached.

Sources

CUDA Programming Guide 2.1.7, "Error Checking in CUDA", for launches returning no cudaError_t: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html (checked 2026-08-30)
CUDA Runtime API, Error Handling, for what cudaGetLastError and cudaPeekAtLastError each do to the stored status: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__ERROR.html (checked 2026-08-30)
CUDA Runtime API, Data types, for the "inconsistent state" sentence that marks a sticky code: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html (checked 2026-08-30)
PyTorch, get_cuda_async_error_suffix, the warning it attaches to its own stack traces: https://github.com/pytorch/pytorch/blob/main/c10/cuda/CUDAMiscFunctions.cpp (checked 2026-08-30)
"What is the canonical way to check for errors using the CUDA runtime API?", 178,060 views: https://stackoverflow.com/questions/14038589/what-is-the-canonical-way-to-check-for-errors-using-the-cuda-runtime-api (checked 2026-08-29)

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.