CUDA，Memcpy 中“遇到非法内存访问”

Question

我有这个 cuda 文件：

#include "cuda.h"
#include "../../HandleError.h"
#include "Sphere.hpp"
#include <stdlib.h>
#include <CImg.h>

#define WIDTH 1280
#define HEIGHT 720
#define rnd(x) (x*rand()/RAND_MAX)
#define SPHERES_COUNT 5

using namespace cimg_library;

__global__
void kernel(unsigned char* bitmap, Sphere* s)
{
   // Map threadIdx/blockIdx to pixel position
   int x = threadIdx.x + blockIdx.x * blockDim.x;
   int y = threadIdx.y + blockIdx.y * blockDim.y;
   int offset = x + y * blockDim.x * gridDim.x;
   float ox = x - blockDim.x * gridDim.x / 2;
   float oy = y - blockDim.y * gridDim.y / 2;
   float r = 0.2, g = 0.2, b = 0.5;
   float maxz = -INF;
   for (int i = 0; i < SPHERES_COUNT; i++) {
       float n, t = s[i].hit(ox, oy, &n);
       if (t > maxz) {
           float fscale = n;
           r = s[i].r * fscale;
           g = s[i].g * fscale;
           b = s[i].b * fscale;
           maxz = t;
       }
   }

   bitmap[offset*3] = (int)(r * 255);
   bitmap[offset*3 + 1] = (int)(g * 255);
   bitmap[offset*3 + 2] = (int)(b * 255);
}

__constant__ Sphere s[SPHERES_COUNT];

int main ()
{
    //Capture start time
    cudaEvent_t start, stop;
    HANDLE_ERROR(cudaEventCreate(&start));
    HANDLE_ERROR(cudaEventCreate(&stop));
    HANDLE_ERROR(cudaEventRecord(start, 0));

    //Create host bitmap
    CImg<unsigned char> image(WIDTH, HEIGHT, 1, 3);
    image.permute_axes("cxyz");

    //Allocate device bitmap data
    unsigned char* dev_bitmap;
    HANDLE_ERROR(cudaMalloc((void**)&dev_bitmap, image.size()*sizeof(unsigned char)));

    //Generate spheres and copy them on the GPU one by one
    Sphere* temp_s = (Sphere*)malloc(SPHERES_COUNT*sizeof(Sphere));
    for (int i=0; i <SPHERES_COUNT; i++) {
        temp_s[i].r = rnd(1.0f);
        temp_s[i].g = rnd(1.0f);
        temp_s[i].b = rnd(1.0f);
        temp_s[i].x = rnd(1000.0f) - 500;
        temp_s[i].y = rnd(1000.0f) - 500;
        temp_s[i].z = rnd(1000.0f) - 500;
        temp_s[i].radius = rnd(100.0f) + 20;
    }

    HANDLE_ERROR(cudaMemcpyToSymbol(s, temp_s, sizeof(Sphere)*SPHERES_COUNT));
    free(temp_s);

    //Generate a bitmap from spere data
    dim3 grids(WIDTH/16, HEIGHT/16);
    dim3 threads(16, 16);
    kernel<<<grids, threads>>>(dev_bitmap, s);

    //Copy the bitmap back from the GPU for display
    HANDLE_ERROR(cudaMemcpy(image.data(), dev_bitmap,
                            image.size()*sizeof(unsigned char),
                            cudaMemcpyDeviceToHost));

    cudaFree(dev_bitmap);

    image.permute_axes("yzcx");
    image.save("render.bmp");
}

编译正常，但执行时出现此错误：

an illegal memory access was encountered in main.cu at line 82

也就是说，这里：

    //Copy the bitmap back from the GPU for display
    HANDLE_ERROR(cudaMemcpy(image.data(), dev_bitmap,
                            image.size()*sizeof(unsigned char),
                            cudaMemcpyDeviceToHost));

我不明白为什么... 我知道如果删除这个：

  bitmap[offset*3] = (int)(r * 255);
  bitmap[offset*3 + 1] = (int)(g * 255);
  bitmap[offset*3 + 2] = (int)(b * 255);

这个错误没有报错，所以我以为可能是out of index错误，后来报了，但是我有这个程序的一个相同版本，不使用常量内存，并且在相同的版本下也可以正常工作内核函数的...

Answer 1

这里有两个问题。第一个是这样的：

__constant__ Sphere s[SPHERES_COUNT];

int main ()
{
    ......

    kernel<<<grids, threads>>>(dev_bitmap, s);

    ......

在主机代码中，s 是主机内存变量，它为 CUDA 运行时提供句柄以与设备常量内存符号挂钩。它不包含有效的设备指针，并且不能传递给内核调用。结果是无效的内存访问错误。

你可以这样做：

__constant__ Sphere s[SPHERES_COUNT];

int main ()
{
    ......

    Sphere *d_s;
    cudaGetSymbolAddress((void **)&d_s, s);
    kernel<<<grids, threads>>>(dev_bitmap, d_s);

    ......

这将导致符号查找获取s 的设备地址，并将其传递给内核是有效的。但是，GPU 依赖于编译器发出特定的指令来通过常量缓存访问内存。设备编译器仅在检测到内核中正在访问__constant__ 变量时才会发出这些指令，而这在使用指针时是不可能的。您可以在this Stack Overflow question and answer 中了解有关编译器如何生成用于常量变量访问的代码的更多信息。