The CUDA Handbook ▸ Appendix A. The CUDA Handbook Library

A.6 Error Handling

chError.h implements a set of macros that implement the goto-based error handling mechanism described in Section 1.2.3. These macros do the following:

• Assign the return value to a status variable named for the API—status_cudart for the CUDA runtime, status_cuda for the driver API, or status_nvrtc for NVRTC,

• Check that status for success and, if in debug mode, report the error to stderr,

• If it contains an error, goto a cleanup label named for the same API—Error_cudart, Error_cuda, or Error_nvrtc.

Current practice uses the concise cuda() and cu() macros, which take the function name without its API prefix and use token pasting to form the full name: cuda(Malloc( &p, N )); expands to a checked call to cudaMalloc(). Because invocations are barely longer than the unchecked calls they replace, the error handling no longer dominates a listing the way the older, more verbose CUDART_CHECK() and CUDA_CHECK() macros did in the first edition of this book. The debug build of the CUDA runtime version is as follows:

#define cuda( fn ) do { \
        (status_cudart) =  (cuda##fn); \
        if ( cudaSuccess != (status_cudart) ) { \
            fprintf( stderr, "CUDA Runtime Failure (line %d of file %s):\n\t" \
                "%s returned 0x%x (%s)\n", \
                __LINE__, __FILE__, #fn, status_cudart, chGetErrorString(status_cudart) ); \
            goto Error_cudart; \
        } \
    } while (0);

The non-debug version omits the fprintf() and simply performs the goto. The cu() macro does the same for driver API functions, pasting “cu” instead of “cuda”, checking against CUDA_SUCCESS, and using status_cuda and Error_cuda; the nvrtc() macro does likewise for the NVRTC runtime-compilation API, with status_nvrtc and Error_nvrtc.

Each API gets its own status variable and cleanup label because the return types—cudaError_t, CUresult, and nvrtcResult—are distinct enumerations that cannot share a single variable. Suffixing the names by API lets one function check runtime, driver, and NVRTC calls side by side and still goto the right cleanup code; earlier the macros shared a single status and Error, which worked only as long as a function stayed within one API. For a fuller treatment of CUDA error handling – including the reasoning behind this idiom – see CUDA Error Handling: A Definitive Guide.

The do..while is a C programming idiom, commonly used in macros, that causes the macro invocation to evaluate to a single statement. Using these macros will generate compile errors if the status variable (status_cudart, status_cuda, or status_nvrtc) or the label (Error_cudart, Error_cuda, or Error_nvrtc) for the API in use is not defined.

One implication of using goto is that all variables must be declared at the top of the block. Otherwise, some compilers generate errors because the goto statements can bypass initialization. When that happens, the variables being initialized must be moved above the first goto, or moved into a basic block so the goto is outside their scope.

Listing A-3 gives an example function that follows the idiom. The return value and intermediate resources are initialized to values that can be dealt with by the cleanup code. In this case, all of the resources allocated by the function also are freed by the function, so the cleanup code and error handling code are the same. Functions that will only free some of the resources they allocate must implement the success and failure cases in separate blocks of code.

doubleTimedReduction(     int *answer, const int *deviceIn, size_t N,     int cBlocks, int cThreads,    pfnReduction hostReduction){    double ret = 0.0;    int *deviceAnswer = 0;    int *partialSums = 0;    cudaEvent_t start = 0;    cudaEvent_t stop = 0;    cudaError_t status_cudart;     cuda(Malloc( &deviceAnswer, sizeof(int) ) );    cuda(Malloc( &partialSums, cBlocks*sizeof(int) ) );    cuda(EventCreate( &start ) );    cuda(EventCreate( &stop ) );    cuda(DeviceSynchronize() );     cuda(EventRecord( start, 0 ) );    hostReduction(         deviceAnswer,         partialSums,         deviceIn,         N,         cBlocks,         cThreads );    cuda(EventRecord( stop, 0 ) );    cuda(Memcpy(         answer,         deviceAnswer,         sizeof(int),         cudaMemcpyDeviceToHost ) );     ret = chEventBandwidth( start, stop, N*sizeof(int) ) /         powf(2.0f,30.0f);     // fall through to free resources before returningError_cudart:    cudaFree( deviceAnswer );    cudaFree( partialSums );    cudaEventDestroy( start );    cudaEventDestroy( stop );    return ret;