遇到非法内存访问

Question

我是 CUDA 编程的初学者，编写了一个由单个文件组成的程序main.cu，如下所示。

#include <iostream>
#include <opencv2/opencv.hpp>

#define DEBUG(str) std::cerr << "3[1;37m" << __FILE__ << ":" << __LINE__ << ": 3[1;31merror:3[0m " << str << std::endl;

#define CUDADEBUG(cudaError)      \
    if (cudaError != cudaSuccess) \
        DEBUG(cudaGetErrorString(cudaError));

#define ERROR(str)  \
    {               \
        DEBUG(str); \
        exit(1);    \
    }

__global__ void makeGrey(
    unsigned char *&pimage,
    const int &cn,
    const size_t &total)
{
    unsigned i = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned icn = i * cn;

    printf("%u\n", i);

    if (i < total)
    {
        float result = pimage[icn + 0] * .114 +
                       pimage[icn + 1] * .587 +
                       pimage[icn + 2] * .299;
        pimage[icn + 0] = result; //B
        pimage[icn + 1] = result; //G
        pimage[icn + 2] = result; //R
        // pimage[icn + 3] *= result; //A
    }
}

int main(int argc, char **argv)
{
    if (argc != 3)
        ERROR("usage: executable in out");

    cv::Mat image;
    unsigned char *dimage;

    image = cv::imread(argv[1], cv::IMREAD_UNCHANGED);
    if (!image.data)
        ERROR("Image null");

    if (image.empty())
        ERROR("Image empty");

    if (!image.isContinuous())
        ERROR("image is not continuous");

    const size_t N = image.total();
    const int cn = image.channels();
    const size_t numOfElems = cn * N;
    const int blockSize = 512;
    const int gridSize = (N - 1) / blockSize + 1;

    CUDADEBUG(cudaMalloc(&dimage, numOfElems * sizeof(unsigned char)));
    CUDADEBUG(cudaMemcpy(dimage, image.data, numOfElems * sizeof(unsigned char), cudaMemcpyHostToDevice));

    makeGrey<<<gridSize, blockSize>>>(dimage, cn, N);
    cudaError_t errSync = cudaGetLastError();
    cudaError_t errAsync = cudaDeviceSynchronize();
    if (errSync != cudaSuccess)
        std::cerr << "Sync kernel error: " << cudaGetErrorString(errSync) << std::endl;
    if (errAsync != cudaSuccess)
        std::cerr << "Async kernel error: " << cudaGetErrorString(errAsync) << std::endl;

    CUDADEBUG(cudaMemcpy(image.data, dimage, numOfElems * sizeof(unsigned char), cudaMemcpyDeviceToHost)); //line 73
    CUDADEBUG(cudaFree(dimage));                                                                           //line 74

    cv::imwrite(argv[2], image);
    return 0;
}

当我执行程序时，我得到

Async kernel error: an illegal memory access was encountered
/path-to-main.cu:73: error: an illegal memory access was encountered
/path-to-main.cu:74: error: an illegal memory access was encountered

我检查了 CV_VERSION 宏，它是 4.5.3-dev，并且安装了 Cuda Toolkit 11.4（nvcc 版本 11.4）。另外 afaik，内核根本不执行（我使用了 Nsight gdb 调试器和 printf）。我不明白为什么我正在访问非法内存区域。我感谢任何帮助。提前谢谢你。

Answer 1

如评论中所述，您的 GPU 函数通过引用获取参数。

__global__ void makeGrey(
    unsigned char *&pimage,
    const int &cn,
    const size_t &total)

这很糟糕，传递对函数的引用或多或少意味着您传递的是可以找到值的地址，而不是值本身。在您的情况下，这些值在主机使用的内存中，NOT Device/GPU 内存中，当 GPU 尝试访问这些值时，它很可能会崩溃。

您尝试传递的类型，unsigned char*、int 和 size_t 的复制成本非常低，没有必要首先通过引用传递它们。

__global__ void makeGrey(
    unsigned char *pimage,
    const int cn,
    const size_t total)

nvidia 提供了调试 CUDA 应用程序的工具，但我不是很熟悉它们，您也可以在 GPU 函数中使用 printf，但您必须组织可能来自数千个的输出线程。

一般来说，无论何时调用 GPU 函数，都要对作为参数传递的内容非常谨慎，因为它们需要从主机内存传递到设备内存。通常你希望按值传递所有内容，任何指针都需要指向设备内存，并注意引用。

遇到非法内存访问

an illegal memory access was encountered

c++

opencv

cuda