CUDA 超过最大线程数没有错误？答案

【问题标题】：CUDA more than max threads without errors?CUDA 超过最大线程数没有错误？
【发布时间】：2013-04-17 08:15:29
【问题描述】：

最初的问题是启动更多线程，这可能是这样的：

someKernel<<<1 , 1025>>> ( ... );

并没有检测到错误，因为我不知道如何检测内核调用错误。这在这个问题的 talonmies 答案中得到了很好的解释：

What is the canonical way to check for errors using the CUDA runtime API?

为了简洁起见，我自己编写了自己的代码，而不是修改我提供的代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t cudaError, char *file, int line, bool abort=true)
{
   if (cudaError != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(cudaError), file, line);
   }
}

__global__ void addKernel(const int *dev_a, const int *dev_b, int *dev_c)
{
    int i = threadIdx.x;
    if ( i < 5 )
        dev_c[i] = dev_a[i] + dev_b[i];
}

int main()
{
    const int arraySize = 5;
    const int a[arraySize] = { 1, 2, 3, 4, 5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };

    int *dev_a(nullptr), *dev_b(nullptr), *dev_c(nullptr);

    gpuErrchk( cudaMalloc((void**)&dev_a, arraySize * sizeof(int)) );
    gpuErrchk( cudaMalloc((void**)&dev_b, arraySize * sizeof(int)) );
    gpuErrchk( cudaMalloc((void**)&dev_c, arraySize * sizeof(int)) );

    gpuErrchk( cudaMemcpy(dev_a, a, arraySize * sizeof(int), cudaMemcpyHostToDevice) );
    gpuErrchk( cudaMemcpy(dev_b, b, arraySize * sizeof(int), cudaMemcpyHostToDevice) );

    const int testMax1D = 1025; 
    dim3 testMax2D ( 32, 33 );

    addKernel<<<1, testMax2D>>> ( dev_a , dev_b, dev_c );
    gpuErrchk( cudaPeekAtLastError() );
    gpuErrchk( cudaDeviceSynchronize() );

    gpuErrchk(  cudaMemcpy( c, dev_c, arraySize * sizeof(int), cudaMemcpyDeviceToHost) );

    printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
        c[0], c[1], c[2], c[3], c[4]);

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    return 0;
}

我现在得到正确的错误报告。感谢您的耐心等待。

看不懂gpuAssert函数中的这个调用，所以省略了：

if (abort) exit(code);

exit 是自定义编写的函数还是我错过了什么？

【问题讨论】：

您需要在每次内核启动后check the CUDA error。当您超过启动配置中的最大块数或线程数时，您将收到“CUDA 未指定的启动错误”等。
@JaredHoberock 我添加了整个代码。我（或者更确切地说是 nVidia 的示例开发人员）似乎在每个可能的步骤上检查错误。我错过了一个？
是的。您不了解如何检查内核启动错误。仔细阅读 Jared 建议的链接。
我不确定您从哪里获得此代码或为什么将其归于 NVIDIA。它可能来自 NVIDIA，我不知道。但是 NVIDIA vectorAdd 示例是 here，它包含适当的内核错误检查。
@DanS：请阅读this answer的后半部分，直到您理解为止，然后根据链接答案中的建议修改此代码，然后运行您的代码。

标签： visual-studio-2010 cuda parallel-processing

【解决方案1】：

内核启动可能会出现两类错误，需要按照特定顺序在单独的步骤中检查它们。

第一类错误是在内核调用时和内核实际在设备上启动之前同步报告的，即这些是“启动前”错误。这些错误通常涉及请求比可用资源更多的特定资源（例如，共享内存过多、线程过多）。通过在内核调用后立即调用 cudaGetLastError() 来检查这些。

第二类错误是在设备上启动内核后的某个时间点发生的错误（例如内存访问冲突、看门狗定时器超时）。这些是“发布后”错误。内核调用后一段时间报告它们的原因是内核启动异步发生的自然结果。下次有机会报告它们，通常是下一次同步 API 调用。通过调用 cudaDeviceSynchronize() 并检查其状态返回来检查这些。

发布的代码缺少对第一类错误的检查。

【讨论】：