从设备复制到主机时出现 CUDA 未知错误答案

【问题标题】：CUDA unkown error when copying from device to host从设备复制到主机时出现 CUDA 未知错误
【发布时间】：2013-08-06 14:03:45
【问题描述】：

我编写了一些 CUDA 代码，在我尝试从代码中获取结果之前，一切似乎都很好：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cstdlib> 
#include <ctime> 
#include <iostream>

#define maskSize 3 

__constant__ float masks[32*maskSize*maskSize];

__global__ void myConv(float *res, const float* mats, int mSize)
{
    extern __shared__ float curr[];
    int rSize=maskSize+mSize-1;
    int idxmod=(threadIdx.x+maskSize-1) % (mSize+2*maskSize-2); //these two map any value not within (mSize-1,mSize-1) to the boarders for padding.
    int idymod=(threadIdx.y+maskSize-1) % (mSize+2*maskSize-2);
    if (threadIdx.x < mSize && threadIdx.y < mSize) //put the value of mats in the middle of the curr matrix
        curr[(threadIdx.x+ maskSize-1)*(mSize+2*(maskSize-1)) + threadIdx.y + maskSize-1]=mats[mSize*(blockIdx.y*mSize + threadIdx.x) + threadIdx.y];
    else //zero padding
        if (threadIdx.x < mSize)
            curr[threadIdx.x*(mSize+2*(maskSize-1)) +idymod] =0;
        else 
            curr[idxmod*(mSize+2*(maskSize-1)) +threadIdx.y] =0;

    __syncthreads();
    float tmp=0;

if (threadIdx.x < mSize+maskSize-1 && threadIdx.y < mSize+maskSize-1)
{
#pragma unroll
    for (int i=0;i<maskSize;i++)
        #pragma unroll
        for (int j=0;j<maskSize;j++)

            tmp+=curr[(threadIdx.x+i)*(mSize+2*(maskSize-1)) + threadIdx.y+j]*masks[blockIdx.x*maskSize*maskSize +maskSize*i +j];
    res[blockIdx.y*rSize*rSize + threadIdx.x*rSize + threadIdx.y]=tmp;
}
}

int main()
{
    int MatSize=5;
    int bSize=2000;
    int maskNum=10;
    int resSize=MatSize+maskSize-1;
    float* ms;
    ms=(float *)malloc(maskSize*maskSize*maskNum*sizeof(float));
    float* resPtr=(float *)malloc((MatSize+maskSize-1)*(MatSize+maskSize-1)*bSize*maskNum*sizeof(float));
    for (int i=0; i<maskSize;i++)
        for (int j=0; j<maskSize; j++)
            for (int k=0; k<maskNum; k++)
                ms[k*maskSize*maskSize + j*maskSize + i]=(float)(rand() % 1000)/100;
    float* inp=(float *)malloc(MatSize*MatSize*bSize*sizeof(float));
    for (int i=0; i<MatSize; i++)
        for (int j=0; j<MatSize; j++)
            for (int k=0;k<bSize;k++)
                inp[k*MatSize*MatSize + j*MatSize + i]=(float)(rand() % 500)/100;
    float *cudams, *cudaresPtr,*cudainp;
    cudaMalloc((void **) &cudams,maskSize*maskSize*maskNum*sizeof(float));
    cudaMalloc((void **) &cudaresPtr,(MatSize+maskSize-1)*(MatSize+maskSize-1)*bSize*maskNum*sizeof(float));
    cudaMalloc((void **) &cudainp,MatSize*MatSize*bSize*sizeof(float));

    cudaMemcpy((void *)cudams,(void *)ms,maskSize*maskSize*maskNum*sizeof(float),cudaMemcpyHostToDevice);

    cudaMemcpy((void *)cudainp,(void *)inp,MatSize*MatSize*bSize*sizeof(float),cudaMemcpyHostToDevice);

    cudaMemcpyToSymbol(masks,(void *)cudams,maskSize*maskSize*maskNum*sizeof(float),0,cudaMemcpyDeviceToDevice);
    dim3 threadSize(MatSize+2*(maskSize-1),MatSize+2*(maskSize-1));
    dim3 blockSize(1, 1); //for testing purposes. should be dim3 blockSize(maskNum,bSize);
    myConv<<<blockSize, threadSize, (MatSize+2*(maskSize-1))*(MatSize+2*(maskSize-1))>>>(cudaresPtr,cudainp,MatSize);
    cudaMemcpy((void *)resPtr,(const void *)cudaresPtr,(MatSize+maskSize-1)*(MatSize+maskSize-1)*bSize*maskNum*sizeof(float),cudaMemcpyDeviceToHost);
    //The problem is here - They copying won't work!

    free(inp);
    free(ms);
    free(resPtr);
    return 0;
}

我将 printf 放在不同的地方，按照这里的建议使用错误检查，打印错误字符串...找不到任何会导致将指针内容复制回主机时出错的内容。

编辑：memcheck 结果：如果我理解正确，没有错误：

O:\CudaTst>cuda-memcheck CUDA_TST ========= CUDA-MEMCHECK

花费的时间：0.144000 秒错误：无法读取字符串错误记录 ========= 错误摘要：0 个错误

使用 -l（泄漏）重新运行 - 0 次泄漏。

【问题讨论】：

尝试使用 cuda-memcheck 运行您的代码。您的内核中可能有越界访问。
完成。没有报告错误，我已经相应地编辑了我的问题
您能否编辑尽可能短的完整示例，以便其他人可以编译并遇到您的问题？这很可能是内核中的错误（例如，参见this question）。我建议按照this question and answer set确切中描述的错误检查，它将为您提供有关错误来源的更准确的信息。
完成。我尝试了您之前提到的错误检查，但唯一给我错误的是在内核完成后尝试将内容复制回主机。宏使用的函数打印了一些文本，但即使我注释掉了 if (abort) 行，它仍然在我看到打印的内容之前退出。
@user1999728：您的代码遗漏的关键是内核启动后的彻底错误检查。看看我回答中的第二个代码 sn-p 。 cudaPeekAtLastError(); cudaDeviceSynchronize() 模式将内核参数错误与内核执行错误与后续 API 错误隔离开来。

标签： cuda

【解决方案1】：

您似乎（至少）在启动内核时动态分配的共享内存不足，无法在内核内部没有缓冲区溢出的情况下运行。

每个块的共享内存量以 bytes 为单位，所以我怀疑你想要这样的东西：

size_t shmsz = sizeof(float)*size_t((MatSize+2*(maskSize-1))*
                                    (MatSize+2*(maskSize-1));
myConv<<<blockSize, threadSize, shmz)>>>(cudaresPtr,cudainp,MatSize);

除此之外，我把调试留给你。

【讨论】：