如何将设备内存中分配的结构化数据从设备复制到主机

Question

我是 GPU 和 CUDA 编程的新手。我正在尝试将设备上动态分配的结构化数据从设备复制到主机。我从 GPU 编程指南中修改了一个简单的代码。编译代码时我没有收到任何错误，但我唯一有问题的是输出错误，即“0”。代码如下：

#include <stdlib.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

typedef struct Point
{
    int2 pt;
};

#define NUMOFBLOCKS 1
#define THREDSPERBLOCK  16

__device__ Point* pnt[NUMOFBLOCKS];
Point dataptr_h[NUMOFBLOCKS][THREDSPERBLOCK];

__global__ void allocmem() 
{   
    if (threadIdx.x == 0)       
        pnt[blockIdx.x] = (Point*)malloc(1*blockDim.x * sizeof(Point)); 
    __syncthreads(); 
} 

__global__ void usemem() 
{ 
    Point* ptr = pnt[blockIdx.x]; 
    if (ptr != NULL) 
    {       
        ptr[threadIdx.x].pt.x = threadIdx.x; 
        ptr[threadIdx.x].pt.y = threadIdx.x;
        printf("Ptr = %d\t", ptr[threadIdx.x].pt.x);
    }
}

__global__ void freemem() 
{ 
    Point* ptr = pnt[blockIdx.x]; 
    if (ptr != NULL) 
        printf("Block %d, Thread %d: final value = %d\n", blockIdx.x, threadIdx.x, ptr[threadIdx.x]); 
    if (threadIdx.x == 0) 
        free(ptr); 
}


int main()
{
    Point* d_pt[NUMOFBLOCKS];
    for (int i = 0 ; i < NUMOFBLOCKS; i++)
        cudaMalloc(&d_pt[i], sizeof(Point)*16);  

    // Allocate memory  
    allocmem<<< NUMOFBLOCKS, THREDSPERBLOCK >>>();  
    // Use memory 
    usemem<<< NUMOFBLOCKS, THREDSPERBLOCK >>>(); 
    cudaMemcpyFromSymbol(d_pt, pnt, sizeof(d_pt));
    cudaMemcpy(dataptr_h, d_pt, sizeof(dataptr_h), cudaMemcpyDeviceToHost);

    for (int j = 0 ; j < 1; j++)
        for (int i = 0 ; i < 16; i++)
        {
            printf("\nPtr_h(%d,%d)->X = %d\t", j, i, dataptr_h[j][i].pt.x);
            printf("Ptr_h(%d,%d)->Y = %d", j, i, dataptr_h[j][i].pt.y);
        }

    freemem<<< NUMOFBLOCKS, THREDSPERBLOCK >>>();
    cudaDeviceSynchronize();
    return 0;
}

代码的输出是：

Ptr_h(0,0)->X = 0       Ptr_h(0,0)->Y = 0
Ptr_h(0,1)->X = 0       Ptr_h(0,1)->Y = 0
Ptr_h(0,2)->X = 0       Ptr_h(0,2)->Y = 0
Ptr_h(0,3)->X = 0       Ptr_h(0,3)->Y = 0
Ptr_h(0,4)->X = 0       Ptr_h(0,4)->Y = 0
Ptr_h(0,5)->X = 0       Ptr_h(0,5)->Y = 0
Ptr_h(0,6)->X = 0       Ptr_h(0,6)->Y = 0
Ptr_h(0,7)->X = 0       Ptr_h(0,7)->Y = 0
Ptr_h(0,8)->X = 0       Ptr_h(0,8)->Y = 0
Ptr_h(0,9)->X = 0       Ptr_h(0,9)->Y = 0
Ptr_h(0,10)->X = 0      Ptr_h(0,10)->Y = 0
Ptr_h(0,11)->X = 0      Ptr_h(0,11)->Y = 0
Ptr_h(0,12)->X = 0      Ptr_h(0,12)->Y = 0
Ptr_h(0,13)->X = 0      Ptr_h(0,13)->Y = 0
Ptr_h(0,14)->X = 0      Ptr_h(0,14)->Y = 0
Ptr_h(0,15)->X = 0      Ptr_h(0,15)->Y = 0

我能做些什么来解决这个问题？

Answer 1

您不能将设备 malloc 操作创建的指针与 CUDA 运行时 API 一起使用（即 cudaMemcpy）

所以这行代码有问题：

cudaMemcpy(dataptr_h, d_pt, sizeof(dataptr_h), cudaMemcpyDeviceToHost);

d_pt 包含从pnt 提取的指针。 pnt 的值由设备 malloc 设置。

相反，您需要创建使用cudaMalloc 正确分配的区域，然后首先将您想要的数据复制到这些区域（从设备上的一个区域到另一个区域），然后使用cudaMemcpy 复制到主机。

在我进一步解释您的下一个反对意见之前，让我们明确一下，以上是您的意图（使用在设备malloc 操作中创建的指针作为cudaMemcpy 的目标之一）。这是不合法的。

“但是我用过 cudaMalloc ??”

d_pt 是一个存在于主机内存中的指针数组。您使用cudaMalloc 获取了这些指针中的每一个并为其分配了一个值（设备内存中的一个指向位置）。

那么这行代码：

cudaMemcpyFromSymbol(d_pt, pnt, sizeof(d_pt));

重写所有你用从设备内存中其他地方获得的指针设置的指针，特别是设备malloc分配的指针。虽然这在技术上是合法的（那行代码不会引发错误），但这些指针在主机上是无用的（无论如何，与运行时 API 一起使用）。