CUDA Error Handling

The Two Faces of CUDA Errors

When writing CUDA code, you'll encounter two primary categories of errors. Understanding their differences is key to handling them correctly.

1. Synchronous Errors: These errors occur in blocking API calls, and the function returns an error code immediately. For example, if you call cudaMalloc() and request more memory than is available on the GPU, the function will fail and return an error code right away. Most CUDA runtime API functions return a value of type cudaError_t. A return value of cudaSuccess indicates the call was successful. Any other value signals a problem.

2. Asynchronous Errors: These are more subtle. CUDA operations like kernel launches (myKernel<<<...>>>()) are asynchronous by design. The CPU launches the kernel and immediately continues to the next line of code without waiting for the kernel to finish. An error might occur much later, during the kernel's execution on the GPU (e.g., an illegal memory access that wasn't caught by your bounds check). This error won't be reported by the launch call itself. Instead, the CUDA runtime will record the error and report it back to you on the next synchronous API call your host code makes, such as cudaMemcpy() or cudaDeviceSynchronize().

This "delayed" reporting of asynchronous errors can be a major source of confusion during debugging. An error reported by cudaMemcpy might have actually been caused by a kernel that was launched long before it.

The following reading from NVIDIA's official guidelines provides excellent context on this distinction and includes clear "Bad" vs. "Good" examples that demonstrate why immediate checking is so important.

A Systematic Approach to Error Handling

Checking the return code of every single CUDA call with a unique if statement is tedious and clutters your code. A much cleaner and more robust approach is to use a helper function or, more commonly, a C/C++ preprocessor macro. This standardizes error checking and provides rich diagnostic information when an error occurs.

Handling API Call Errors

Let's start with API calls. We need a way to wrap a CUDA call and check its return value. A simple wrapper function might look like this:

#include <iostream>

void checkCuda(cudaError_t result, char const *const func, const char *const file, int const line) {
    if (result != cudaSuccess) {
        std::cerr << "CUDA error at " << file << ":" << line << " code=" << static_cast<unsigned int>(result)
                  << " \"" << cudaGetErrorString(result) << "\" in " << func << std::endl;
        // Make sure we exit on failure
        exit(EXIT_FAILURE);
    }
}

// This is the macro we'll use



#define checkCudaErrors(val) checkCuda((val), #val, __FILE__, __LINE__)

// Example Usage:
// checkCudaErrors(cudaMalloc(&d_data, size));

Here, checkCudaErrors is a macro that takes a CUDA API call as its argument. It passes the return value of that call, along with the source code of the call itself (#val), the filename (__FILE__), and the line number (__LINE__), to our checkCuda helper function. The helper function checks the result and, if it's not cudaSuccess, prints a detailed error message and terminates the program.

Handling Kernel Launch Errors

This wrapper works for API calls, but what about kernel launches, which return void? As we discussed, a kernel launch can produce both synchronous errors (e.g., invalid launch configuration) and asynchronous errors (e.g., a crash during execution).

To catch these, we must explicitly ask CUDA for the last error that occurred. The function for this is cudaGetLastError(). It returns the last error that has been recorded for the host thread, and importantly, it resets the error state to cudaSuccess (for non-sticky errors). Because it handles both API errors (which are "sticky" and will be returned by cudaGetLastError as well) and kernel launch errors, it's a universal tool.

We can create a single, powerful macro that works for both cases.

The macro from the presentation is an excellent template for our own projects:

#include <stdio.h>




#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

You would use this macro after any API call or kernel launch you want to check.

Example Usage:

// Allocate memory
float *d_a;
cudaMalloc(&d_a, size);
cudaCheckErrors("cudaMalloc for d_a failed"); // Checks the result of cudaMalloc

// Launch kernel
vectorAdd<<<numBlocks, threadsPerBlock>>>(d_a, d_b, d_c, N);
cudaCheckErrors("vectorAdd kernel launch failed"); // Checks for launch or async errors

// Copy result back
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
cudaCheckErrors("cudaMemcpy DtoH failed"); // Checks the result of cudaMemcpy

This approach ensures that if any step fails, the program halts immediately with a precise error message, pointing you directly to the source of the problem.

Deeper Dive: Sticky Errors and Final Synchronization

Your interest in low-level systems warrants a closer look at what happens when an unrecoverable error occurs.

CUDA distinguishes between non-sticky and sticky errors.

• Non-sticky errors are recoverable. A cudaMalloc failure is a good example. The call fails, but the CUDA context is still healthy, and you can continue to make other API calls (perhaps to try allocating a smaller amount of memory).
• Sticky errors are unrecoverable. These are typically caused by serious problems during a kernel's execution, like an invalid memory access, an illegal instruction, or a device-side assert. When a sticky error occurs, the CUDA context becomes corrupted. From that point on, all subsequent CUDA API calls will return the same sticky error. The only way to recover is to terminate the application.

This behavior is why our cudaCheckErrors macro calls exit(1). For many critical CUDA errors, there is no path to recovery within the same process.

There's one final "pro-level" tip for ensuring you don't miss any errors. Because kernels run asynchronously, it's possible for your main function to finish and your program to exit before a launched kernel has completed (or failed). If the kernel fails after main has returned, you will never see the error.

To prevent this, it's a best practice to explicitly synchronize with the device just before your program terminates.

int main() {
    // ... all your CUDA work ...

    // Wait for all outstanding CUDA tasks to complete
    cudaDeviceSynchronize();
    cudaCheckErrors("Final synchronization failed");

    return 0;
}

cudaDeviceSynchronize() is a blocking call that waits for all preceding commands in all streams on the current device to complete. By placing it at the end of main, you force the host to wait for any running kernels to finish, ensuring that any asynchronous errors they generate will be caught by your final cudaCheckErrors call.

Automated Checking with Compute-Sanitizer

While manual error checking is a fundamental skill, NVIDIA also provides powerful tools to automate error detection. The primary tool for this is compute-sanitizer. It's a command-line utility that you run your compiled application through. It instruments your code at runtime to detect a wide range of issues, even if you don't have cudaCheckErrors calls everywhere.

Its memcheck tool, for instance, can detect:

• Out-of-bounds global, shared, and local memory accesses.
• Misaligned memory accesses.
• Invalid __syncthreads() usage.

To use it, you compile your code with debug information (-lineinfo) and run it like so:
$ nvcc -o my_app my_app.cu -lineinfo
$ compute-sanitizer ./my_app

If it finds an error, it will print a detailed report, including a stack trace pointing to the exact line of kernel code that caused the problem. This is an indispensable tool for debugging complex kernels, and we will rely on it more as we progress. However, it's a supplement to, not a replacement for, good manual error-handling practices.

Conclusion

You have now added another essential layer of robustness to your CUDA programming toolkit. While bounds checking protects your kernel's internal logic, systematic API error handling protects your program's interaction with the entire CUDA ecosystem.

Key Takeaways:

• CUDA API calls can produce synchronous errors (returned immediately) or asynchronous errors (reported by a later call).
• Kernel launches (<<<...>>>) return void and must be checked for errors using cudaGetLastError().
• A unified error-checking macro that uses cudaGetLastError() is a robust and clean way to check both API calls and kernel launches. This macro should provide detailed diagnostic information (error string, file, line number).
• Severe kernel execution errors are often sticky, corrupting the CUDA context and requiring the program to terminate.
• Always call cudaDeviceSynchronize() before your program exits to ensure all pending asynchronous errors are caught.
• Tools like compute-sanitizer can automate error detection and are invaluable for debugging.

With the ability to write code that is both safe from out-of-bounds access and resilient to API failures, we can now confidently move on to the next major topic: performance. In our next lesson, we will begin our deep dive into optimization by tackling the single most important factor for performance in most kernels: global memory access patterns and coalescing.