CUDA 内核因 20k+ 线程无故崩溃答案

【问题标题】：CUDA Kernel is crashing without any reason with 20k+ threadsCUDA 内核因 20k+ 线程无故崩溃
【发布时间】：2017-04-24 05:01:25
【问题描述】：

我正在从事一个 CUDA 项目，但我遇到了一些我无法找到解决方案的严重问题。

我使用 NVIDIA Quadro K2000m 在我的 PC (pA) 上实施了该项目，它可以工作。但是，当我将项目部署在具有 Nvidia Tesla GPU 的集群上，并且在另一台 PC (pB) (NVIDIA gtx 960m) 上时，它不会执行！

有趣的是，当我在 pB（第二台 PC）上使用 Visual Studio 中的 Nsight Debugger 时，它会执行并且不会显示错误：Unspecified launch failure

这是第一个内核的代码：

__global__ void calcKernel(float *dev_calcMatrix,

                        int *documentarray,
                        int *documentTermArray,
                        int *distincttermsarray,
                        int *distinctclassarray,
                        int *startingPointOfClassDoc,
                        int *endingPOintOfClassDoc,
                        int sizeOfDistinctClassarray,
                        int sizeOfTerms)
{

 int index = blockIdx.x * blockDim.x + threadIdx.x;

int term = distincttermsarray[index];

if (index <= sizeOfTerms) {

    for (int i = 0; i < sizeOfDistinctClassarray; i++)
    {
        int save = (index * sizeOfDistinctClassarray) + i;
        bool test = false;
        for (int j = startingPointOfClassDoc[i]; j <= endingPOintOfClassDoc[i]; j++)
        {
            if (term == documentarray[j])
            {
                printf("%i \t", index);
                dev_calcMatrix[save] = dev_calcMatrix[save] + documentTermArray[j];

                //printf("TermArray: documentTermArray[j] %d\n", dev_calcMatrix[save], documentTermArray[j]);

                test = true;
            }
        }

        if (!test) dev_calcMatrix[save] = 0;


    }
}
}

这是我用来创建线程和块的代码：

float blockNotFinal = data.sizeOfDistinctTerms / 1024;
int threads = 0;
int  blocks = (int)floor(blockNotFinal);

dim3 dimGrid((blocks + 1), 1, 1);
if (data.sizeOfDistinctTerms < 1024)
{
    threads = data.sizeOfDistinctTerms;
}
else
{
    threads = 1024;
}
dim3 dimBlock(threads, 1, 1);

所以，我需要创建 23,652 个线程。我正在做的是 23,652 / 1024 = 23.09。在我得到 23.09 值后，我将它舍入到 23 并添加 + 1 = 24 个块。所以我正在创建 24 个块 * 1024 个线程：24,576 个线程。

我知道有些线程会被创建，即使它们不会被使用，这就是为什么我在内核的开头添加了这个 if 语句：

int index = blockIdx.x * blockDim.x + threadIdx.x;

if (index <= sizeOfTerms (23,652 is the size)) { .... }

问题是我在 IF 语句之前和 IF 语句之后添加了一些 PRINTF()。

在 IF 语句之前，线程崩溃前的最大索引为：24479 在 IF 语句中，崩溃前的最大线程索引为：23487。

所以，从上面的信息来看，线程数并没有达到最大值。另外，在集群上它给了我另一个错误：Illegal memory access遇到。我知道这个错误意味着它的索引可能超出范围，但我给出的数组大小与线程数相同。

这是我在 GPU 中分配内存的代码：

cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
    goto Error;
}

cout << "\n Allocated GPU buffers";
// Allocate GPU buffers for input and output vectors
cudaStatus = cudaMalloc((void**)&dev_calcMatrix, data.sizeOfDistinctTerms * data.sizeOfDistinctClassarray * sizeof(float));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed!");
    goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_probMatrix, data.sizeOfDistinctTerms * data.sizeOfDistinctClassarray * sizeof(float));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed!");
    goto Error;
}

cudaStatus = cudaMalloc((void**)&classSummationTerms, data.sizeOfDistinctClassarray * sizeof(int));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed!");
    goto Error;
}

cudaStatus = cudaMalloc((void**)&documentarray, data.sizeOfTotalTermsDocsFreq * sizeof(int));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed!");
    goto Error;
}

cudaStatus = cudaMalloc((void**)&documentTermArray, data.sizeOfTotalTermsDocsFreq * sizeof(int));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed!");
    goto Error;
}

cudaStatus = cudaMalloc((void**)&distincttermsarray, data.sizeOfDistinctTerms * sizeof(int));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed!");
    goto Error;
}

cudaStatus = cudaMalloc((void**)&distinctclassarray, data.sizeOfDistinctClassarray * sizeof(int));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed!");
    goto Error;
}

cudaStatus = cudaMalloc((void**)&startingPointOfClassDoc, data.sizeOfDistinctClassarray * sizeof(int));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed!");
    goto Error;
}

cudaStatus = cudaMalloc((void**)&endingPOintOfClassDoc, data.sizeOfDistinctClassarray * sizeof(int));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed!");
    goto Error;
}

cout << "\n Copied input vectors from host to GPU";
// Copy input vectors from host memory to GPU buffers.
cudaStatus = cudaMemcpy(documentarray, data.documentarray, data.sizeOfTotalTermsDocsFreq * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMemcpy failed!");
    goto Error;
}

cudaStatus = cudaMemcpy(documentTermArray, data.documentTermArray, data.sizeOfTotalTermsDocsFreq * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMemcpy failed!");
    goto Error;
}

cudaStatus = cudaMemcpy(distincttermsarray, data.distincttermsarray, data.sizeOfDistinctTerms * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMemcpy failed!");
    goto Error;
}

cudaStatus = cudaMemcpy(classSummationTerms, data.classSummationTerms, data.sizeOfDistinctClassarray * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMemcpy failed!");
    goto Error;
}

cudaStatus = cudaMemcpy(distinctclassarray, data.distinctclassarray, data.sizeOfDistinctClassarray * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMemcpy failed!");
    goto Error;
}

cudaStatus = cudaMemcpy(startingPointOfClassDoc, data.startingPointOfClassDoc, data.sizeOfDistinctClassarray * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMemcpy failed!");
    goto Error;
}

cudaStatus = cudaMemcpy(endingPOintOfClassDoc, data.endingPOintOfClassDoc, data.sizeOfDistinctClassarray * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMemcpy failed!");
    goto Error;
}


cout << "\n Now we call the CALCKERNL()";
// Launch a kernel on the GPU with one thread for each element.
calcKernel <<<dimGrid, dimBlock >>>(dev_calcMatrix,
                            documentarray, 
                            documentTermArray, 
                            distincttermsarray, 
                            distinctclassarray, 
                            startingPointOfClassDoc, 
                            endingPOintOfClassDoc,
                            sizi,
                            sizeOfTerms);

//// cudaDeviceSynchronize waits for the kernel to finish, and returns
//// any errors encountered during the launch.
//cudaStatus = cudaDeviceSynchronize();
//if (cudaStatus != cudaSuccess) {
//  fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
//  goto Error;
//}

cudaStatus = cudaStreamSynchronize(0);
if (cudaStatus != cudaSuccess) {
    //fprintf(stderr, "calcKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
    cout << "\n Synchronization failed: " << cudaGetErrorString(cudaStatus);
    goto Error;
}
// Check for any errors launching the kernel
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "calcKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
    goto Error;
}

知道为什么会这样吗？

【问题讨论】：

不，23,652 就可以了。问题是，他正在运行 23,653 个线程。
我认为您将很难制造minimal reproducible example。
GPU 的线程数有限制吗？对于不同的 GPU，这个限制是否不同？
您可以使用here 描述的方法将非法内存访问错误定位到一行代码。如有必要，您可以使用内核中的printf 或其他方法（例如调试器）来帮助了解为什么该行代码会生成非法访问。

标签： c++ visual-studio cuda gpu nvidia

【解决方案1】：

没有Minimal, Complete, and Verifiable example，甚至没有完整的代码，是不可能回答的。但是您的内核的开头已经有两个错误可能导致内存访问越界：

    int index = blockIdx.x * blockDim.x + threadIdx.x;

    int term = distincttermsarray[index];

    if (index <= sizeOfTerms) {

首先，使用index 作为数组索引在检查它是否在所需范围内之前是不安全的。其次，如果sizeOfTerms 是数组元素的数量，则需要检查index < sizeOfTerms（而不是<=）。

【讨论】：

哇，我确实犯了一些严重的错误:(。我修复了你告诉我的问题，但我仍然有同样的问题:(我在 Visual Studio 上做了一个清理，因为这有时会给我真的很难的问题，但它仍然没有工作。我编译它并在我提到的集群中运行它，但仍然是同样的问题:(
在 cuda-memcheck 下运行代码并修复指出的问题。准备一个 MCVE。那我们拭目以待。

【解决方案2】：

找到并修复此错误的一种简单方法是按照@tera 的建议打开 cuda-memcheck 并使用 Cuda 调试器 运行代码，而无需任何操作命中点。调试器应该在错误发生的那一刻停止。

我的建议是关闭 TDR 的 Nsight + Visual Studio，所以如果非法错误需要一些时间才会发生，这不是问题。

【讨论】：