遇到非法内存访问答案

【问题标题】：an illegal memory access was encountered遇到非法内存访问
【发布时间】：2021-11-22 07:49:06
【问题描述】：

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

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

#define DEBUG(str) std::cerr << "\033[1;37m" << __FILE__ << ":" << __LINE__ << ": \033[1;31merror:\033[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）。我不明白为什么我要访问一个非法的内存区域。我很感激任何帮助。提前谢谢你。

【问题讨论】：

您的 GPU 函数 makeGrey 通过引用获取其参数，这些值位于堆栈中，而不是 GPU 内存中，而是按值获取。
谢谢@Kaldrr。您知道有关如何调试 cuda 应用程序的任何建议或教程吗？我必须手动检查数百万个线程吗？其次，很抱歉我不知道我不应该使用引用。我认为cudaMalloc 支持引用。这是为什么？最后，我应该怎么做才能提高你的声誉？如果你写了一个答案，我会接受或不惜一切代价。
CUDA调试在线教程here.
@Kaldrr 我接受了你的回答。在您的帖子中，您告诉参数值驻留在内存中，但我使用了cudaMalloc。难道不是在GPU（VRAM等）中分配一块内存而不是主内存吗？我认为普通的 c++ 引用风格也可以在那里工作，就像我们对主内存所做的那样。
@HakanDemir 您使用cudaMalloc 的方式与将参数传递给您的makeGrey 函数的方式无关。每当您在 C++ 中作为函数参数引用时，您就是在告诉编译器您不是直接传递值，而是在内存中存在该值的位置。但是这个值存在于主机/CPU 内存中，这意味着设备/GPU 代码将尝试使用引用来读取值，该引用指向主机/CPU 内存中的一个值，这在您的内存中是非常未定义的行为（很可能是崩溃）案例。

标签： c++ opencv cuda

【解决方案1】：

正如评论中提到的，您的 GPU 函数通过引用获取参数。

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

这很糟糕，传递对函数的引用或多或少意味着您传递的是可以找到值的地址，而不是值本身。在您的情况下，这些值位于主机使用的内存中，NOT 设备/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 函数时，都要对作为参数传递的内容非常谨慎，因为它们需要从主机内存传递到设备内存。通常你想通过值传递所有东西，任何指针都需要指向设备内存，并注意引用。

【讨论】：