如何从主机访问设备二维数组全局变量答案

【问题标题】：How to access device 2D array global variable from host如何从主机访问设备二维数组全局变量
【发布时间】：2016-03-26 16:25:02
【问题描述】：

我想在 main 方法中使用“printf”打印 d_t 全局二维数组变量。但我收到一个编译警告说：

不能在宿主函数中直接读取 __device__ 变量“d_t”

如何将全局二维数组变量从设备复制到主机，然后打印每一行的第一列？

__device__ double *d_t;

__device__ size_t d_gridPitch;


__global__ void kernelFunc()
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    double* rowt = (double*)((char *)d_t + i * d_gridPitch);
    rowt[0] = rowt[0] + 40000;

}


int main()
{
    int size = 16;
    size_t d_pitchLoc;
    double *d_tLoc;

    cudaMallocPitch((void**)&d_tLoc, &d_pitchLoc, size * sizeof(double), size);
    cudaMemset2D(d_tLoc, d_pitchLoc, 0, size * sizeof(double), size);

    cudaMemcpyToSymbol(d_gridPitch, &d_pitchLoc, sizeof(int));
    cudaMemcpyToSymbol(d_t, & d_tLoc, sizeof(d_tLoc));

    kernelFunc<<<1,size>>>();

    for(int i=0; i< size; i++){
        double* rowt = (double*)((char *)d_t + i * d_gridPitch);
        printf("%.0f, ",rowt[0]);
    }

    cudaDeviceReset();

    return 0;
}

【问题讨论】：

cudaMemcpy2D 函数用于复制到或从倾斜分配（即使用cudaMallocPitch 创建）。 Here 是 cudaMemcpy2D 的 API 文档。如果您在此 CUDA 标记上进行搜索，您会发现许多证明正确用法的问题和答案，例如 this one。使用proper CUDA error checking。

标签： cuda

【解决方案1】：

正如 cmets 中所指出的，cudaMemcpy2D API 正是为这个任务而设计的。您必须分配或静态定义一个主机内存缓冲区或容器来充当来自设备的数据的存储，然后将该主机缓冲区的间距提供给cudaMemcpy2D 调用。 API 处理音高转换，而无需调用方的任何进一步干预。

如果您将打印循环替换为以下内容：

double* h_t = new double[size * size];
cudaMemcpy2D(h_t, size * sizeof(double), d_tLoc, d_pitchLoc, 
        size * sizeof(double), size, cudaMemcpyDeviceToHost);
for(int i=0, j=0; i< size; i++){
    std::cout << h_t[i * size + j] << std::endl; 
}

[注意我在这里使用iostream 进行打印。 CUDA 使用 C++ 编译器来编译主机代码，您应该更喜欢 iostream 函数而不是 cstdio，因为它们更不容易出错并且支持在大多数平台上改进诊断]。

您可以看到API调用表单与我在您上一个问题中为您提供的cudaMemset2D调用非常相似。

【讨论】：