如何捕获或处理 CUDA 内核启动错误

Question

我使用 CUDA 工具包示例中的 checkCudaErrors 辅助函数。参见 "helper_cuda.h"。我很困惑为什么这个例子中的启动错误没有被 checkCudaErrors 捕获。错误是启动的线程太多 (2048)。

从调试 (linux gdb)，控制台打印（红色标准错误）"warning: Cuda API error detected: cudaLaunch returned (0x9)"。

而当我从 Bash shell 执行 Release 或 Debug 构建时，checkCudaErrors 没有打印错误。

这是为什么？

我的预期是在启动后立即在 D2H memcpy 调用中捕获并打印错误。这不正确吗？

最小可重现示例：

#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 对此进行了解释。

表示有两种错误类型，不同之处在于 API 报告 (returns) 它们的方式*。

我的结果是； 捕获内核启动错误的唯一方法 是在启动调用之后使用 cudaPeekAtLastError() 或 cudaGetLastError()。这些是 return 编辑 启动错误代码 的唯一 API 函数。其他后续 API 调用没有 return 启动错误代码，也没有清除它；稍后可以通过 cudaPeekAtLastError 或 cudaGetLastError 获得。

Answer 2

CUDA 内核启动不会return启动的错误代码。要捕获错误，您需要在启动后以及任何其他 API 调用之前执行一些明确的错误检查：

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

第一次调用应该至少捕获任何启动错误，并且在内核执行期间的第二次调用错误也将被捕获（另请参见 this answer）。由于您尚未执行此操作，因此最早要等到下一次 API 调用时才会看到错误。

如何捕获或处理 CUDA 内核启动错误

How to catch or handle CUDA kernel launch errors

cuda

error-checking