Ivan's wiki

This page is read only. You can view the source, but not change it. Ask your administrator if you think this is wrong.
# 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.

```c
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:

```c
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:

```c
#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.