将多分支树复制到 GPU 内存

Question

我有一棵节点树，我正在尝试将其复制到 GPU 内存。节点看起来像这样:

struct Node
{
   char *Key;
   int ChildCount;
   Node *Children;
}

我的复制函数如下所示:

void CopyTreeToDevice(Node* node_s, Node* node_d)
{


     //allocate node on device and copy host node
     cudaMalloc( (void**)&node_d, sizeof(Node));
     cudaMemcpy(node_d, node_s, sizeof(Node), cudaMemcpyHostToDevice);

     //test
     printf("ChildCount of node_s looks to be : %d\n", node_s->ChildCount);
     printf("Key of node_s looks to be : %s\n", node_s->Key);

     Node *temp;
     temp =(Node *) malloc(sizeof(Node));
     cudaMemcpy(temp, node_d, sizeof(Node), cudaMemcpyDeviceToHost);
     printf("ChildCount of node_d on device is actually : %d\n", temp->ChildCount);
     printf("Key of node_d on device is actually : %s\n", temp->Key);
     free(temp);



     //       continue with child nodes
     if(node_s->ChildCount > 0)
     {
         //problem here
         cudaMalloc( (void**)&(node_d->Children), sizeof(Node)*(node_s->ChildCount));

         cudaMemcpy(node_d->Children, node_s->Children, 
                    sizeof(Node)*node_s->ChildCount, cudaMemcpyHostToDevice);

         for(int i=0;i<node_s->ChildCount;i++)
         {
                 CopyTreeToDevice(&(node_s->Children[i]), &(node_d->Children[i]));
         }
     }

}

但是我的线路有问题:

cudaMalloc( (void**)&(node_d->Children), sizeof(Node)*(node_s->ChildCount));

给我访问冲突异常。测试部分工作顺利。初始化字段没有问题。

这是测试部分的输出:

ChildCount of node_s looks to be : 35
Key of node_s looks to be : root
ChildCount of node_d on device is actually : 35
Key of node_d on device is actually : root

这是什么原因？

谢谢。

Answer 1

node_d->Children 是驻留在设备代码中的变量。您不能像使用第二个 cudaMalloc 那样通过主机代码直接使用它。此外，将主机指针复制到设备没有多大意义，因为您无法在设备代码中取消引用它们。

更好更快的方法是:

• 为整棵树预分配一个大数组。
• 使用数组索引而不是指针。索引的有效性将在传输到设备和从设备传输时得到保留。
• 在设备上分配整个数组一次。拥有多个 memAlloc 可能效率低下(尤其是在 Windows 系统中，当监视器连接到该 GPU 时)。此外，由于 memAlloc 返回的地址始终与 512 字节对齐，因此您实际上无法分配更小的内存块。因此，根据您当前的代码，每个子数组将占用至少 512 个字节，即使里面只有 2 个子数组也是如此。
• 将整个数组从主机复制到设备一次。这比使用多个 memCopy 指令要快得多，即使您实际上复制了一些未使用的额外内存区域。