How to catch or handle CUDA kernel launch errors

Question

I use the checkCudaErrors helper function from the CUDA Toolkit Samples. See "helper_cuda.h". I am perplexed as to why the launch error from this example is not caught by checkCudaErrors. The error is too many threads are launched (2048).

From Debug (linux gdb), the console prints (stderr in red) "warning: Cuda API error detected: cudaLaunch returned (0x9)".

Whereas when I execute either the Release or Debug builds from a Bash shell, no error is printed by checkCudaErrors.

Why is this?

My expectation is the error would be caught and printed at the D2H memcpy call immediately proceeding the launch. Is this incorrect?

Minimal reproducible example:

#include <cuda.h>
#include "helper_cuda.h"

__global__ void BusyIncrementKernel( const size_t increments, float * result){
    float tmp = 0;
    for ( size_t i = 0; i < increments; ++i ){ tmp += 1; }
    const int j = threadIdx.x + blockIdx.x*blockDim.x;
    if ( j == 0 ){ *result = tmp; }
}

int main( int argc, char * argv[] ){
    unsigned int blockDim = 2048;
    dim3 block{ blockDim, 1, 1};
    dim3 grid{ 1, 1, 1};
    float * dResult;
    checkCudaErrors( cudaMalloc( &dResult, sizeof(float) ));
    BusyIncrementKernel<<< grid, block >>>( 10000000, dResult );
    float result;
    checkCudaErrors( cudaMemcpy( &result, dResult, sizeof(float), cudaMemcpyDeviceToHost ));
    checkCudaErrors( cudaFree( dResult ));
    checkCudaErrors( cudaDeviceSynchronize() );
    fprintf( stderr,"result: %f\n", result );
    return 0;
}

Answer 1

This answer by talonmies specifically states kernel launches require a slightly different pattern to handle. The CUDA API documentation 3.2.9. on Error Checking explains this.

This answer by Robert Crovella indicates there are two error types, that differ in how the API reports (returns) them*.

My results are; the only way to catch kernel launch errors is with cudaPeekAtLastError() or cudaGetLastError() after the launch call. These are the only API functions that returned the launch error code. Other subsequent API calls did not return the launch error code, nor did they clear it; it could be obtained later by cudaPeekAtLastError or cudaGetLastError.

Answer 2

CUDA kernel launches do not return an error code for the launch. To catch the error, you need to perform some explicit error checking after the launch, and before any additional API calls:

checkCudaErrors( cudaPeekAtLastError() );
checkCudaErrors( cudaDeviceSynchronize() );

The first call should catch at least any launch errors, and by the second call errors during kernel execution will have been caught as well (see also this answer). Since you haven't done this, you don't see the error until the next API call at the earliest.