从 CUDA 内核访问全局内存指针答案

【问题标题】：accessing a global memory pointer from a CUDA kernel从 CUDA 内核访问全局内存指针
【发布时间】：2014-07-21 04:11:41
【问题描述】：

我正在尝试分配设备内存并将指针存储为全局变量。但是，当我尝试从内核访问内存时，我从 cudaDeviceSynchronize() 收到此错误：cudaErrorIllegalAddress。我检查了从 cudaMalloc 和 cudaMemcpy 返回的 cudaStatus 代码，它们都成功了。

我希望下面的示例能够直截了当地说明我想要做什么。基本上，我有大量的样本数据，我希望所有内核都能从中读取，但我不想每次都将指针传递给内核调用。

我使用的是 Windows 8 x64，使用 Visual Studio 2012 和 nvcc（通过 VS 集成）编译代码。目标是 x64 调试可执行文件。我的设备是 GTX 780。

#include "cuda_runtime.h"
#include <stdio.h>
#define SIZE (1024 * 1024 * 10)

__device__ int* cData;

void Init()
{
    int* data = new int[SIZE];
    cudaError_t cudaStatus;
    cudaStatus = cudaMalloc(&cData, SIZE * sizeof(int));
    for (int i = 0; i < SIZE; i++)
        data[i] = i;

    cudaStatus = cudaMemcpy(cData, data, SIZE * sizeof(int), cudaMemcpyHostToDevice);
    delete data;
}

__global__ void kernel(int i, int* output)
{
    *output = cData[i];
}

int main()
{
    cudaError_t cudaStatus = cudaSetDevice(0);
    cudaDeviceProp properties;
    int* result;
    cudaStatus = cudaMallocManaged(&result, sizeof(int));
    Init();

    kernel<<<1, 1>>>(1000, result); // invoke a single thread, expecting the value of *result to be 1000 afterwards
    cudaStatus = cudaGetLastError();
    cudaStatus = cudaDeviceSynchronize(); // returns cudaErrorIllegalAddress

    printf("Value is: %d", *result); // crashes the program, "In page error reading location 0x0000000D00000000"
    cudaFree(result);
    cudaStatus = cudaDeviceReset();
    return 0;
}

【问题讨论】：

cData 驻留在设备上，但cudaMalloc 要求指针位于主机上（并指向设备内存）。

标签： c++ cuda

【解决方案1】：

我们不对__device__ 变量使用cudaMalloc 和cudaMemcpy。

阅读the documentation 获取__device__ 变量，其中说明要使用的API 调用：

 cudaMemcpyToSymbol();
 cudaMemcpyFromSymbol();

如果您想在动态分配的设备数组上使用cudaMalloc，但将返回的指针存储在__device__ 变量中，则必须执行以下操作：

void Init()
{
    int* data = new int[SIZE];
    int* d_data;
    cudaError_t cudaStatus;
    cudaStatus = cudaMalloc(&d_data, SIZE * sizeof(int));
    for (int i = 0; i < SIZE; i++)
        data[i] = i;

    cudaStatus = cudaMemcpy(d_data, data, SIZE * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpyToSymbol(cData, &d_data, sizeof(int *));
    delete data;
}

当我按原样编译您的代码时，我从 CUDA 6 nvcc 收到以下编译器警告：

t411.cu(15): warning: a __device__ variable "cData" cannot be directly read in a host function

不应忽视这些警告。

如果 SIZE 在编译时已知，就像在您的示例中一样，您也可以执行以下操作：

__device__ int cData[SIZE];

void Init()
{
    int* data = new int[SIZE];
    cudaError_t cudaStatus;
    for (int i = 0; i < SIZE; i++)
        data[i] = i;
    cudaStatus = cudaMemcpyToSymbol(cData, data, SIZE * sizeof(int));
    delete data;
}

【讨论】：