Error handling (CUDA)

Error handling (CUDA)

A kernel launch is asynchronous and returns no error code. Error handling in CUDA has to be explicit, and code that skips it fails silently: a kernel that never ran produces a program that exits with status 0 and wrong or absent output.

kernel<<<blocks, threads>>>(d_a);   // returns void, always

Two checks are needed after a launch. cudaGetLastError() catches launch-time failures such as an invalid block size, and a synchronise catches execution failures such as an out-of-bounds write:

kernel<<<blocks, threads>>>(d_a);
cudaError_t err = cudaGetLastError();          // launch failure
if (err != cudaSuccess)
    fprintf(stderr, "launch: %s\n", cudaGetErrorString(err));
 
err = cudaDeviceSynchronize();                 // execution failure
if (err != cudaSuccess)
    fprintf(stderr, "exec: %s\n", cudaGetErrorString(err));

Since every runtime call returns a code that gets ignored in most example code, the usual practice is a checking macro applied to all of them:

#define CUDA_CHECK(call)                                              \
    do {                                                              \
        cudaError_t e = (call);                                       \
        if (e != cudaSuccess) {                                       \
            fprintf(stderr, "%s:%d: %s\n",                            \
                    __FILE__, __LINE__, cudaGetErrorString(e));       \
            exit(EXIT_FAILURE);                                       \
        }                                                             \
    } while (0)
 
CUDA_CHECK(cudaMalloc(&d_a, bytes));

Errors are sticky and reported late. Because they surface at whatever call happens to synchronise next, the reported location is often far from the real fault. CUDA_LAUNCH_BLOCKING=1 makes launches synchronous so the error appears at the right line, at the cost of destroying any overlap.

For memory bugs specifically, compute-sanitizer ./prog reports out-of-bounds accesses, races, and leaks with a line number, in the same role Valgrind fills on the host. It is worth running before spending time reading the kernel.

Table of Contents

Error handling (CUDA)