CUDA 将继承的类对象复制到设备

Question

我有一个Parent类和一个继承Child类：

class Parent {};
class Child : public Parent {};

有几个子类继承自Parent，但为了简单起见，我只包含了一个。这些继承的类对于我正在从事的项目是必需的。我还有另一个类的对象，我希望将其复制到设备上：

class CopyClass {
  public:
    Parent ** par;
};

请注意，之所以Parent ** par;存在，是因为我需要一个Child对象列表，但它将使用哪个子对象（以及列表的长度）在编译时未知。这是我将CopyClass对象复制到设备上的尝试：

int length = 5;

//Instantiate object on the CPU
CopyClass cpuClass;
cpuClass.par = new Parent*[length];
for(int i = 0; i < length; ++i) cpuClass.par[i] = new Child;

//Copy object onto GPU
CopyClass * gpuClass;
cudaMalloc(&gpuClass,sizeof(CopyClass));
cudaMemcpy(gpuClass,&cpuClass,sizeof(CopyClass),cudaMemcpyHostToDevice);

//Copy dynamically allocated variables to GPU
Parent ** d_par;
d_par = new Parent*[length];
for(int i = 0; i < length; ++i) {
    cudaMalloc(&d_par[i],sizeof(Child));
    printf("\tCopying data\n");
    cudaMemcpy(d_par[i],cpuClass.par[i],sizeof(Child),cudaMemcpyHostToDevice);
}

//SIGSEGV returned during following operation
cudaMemcpy(gpuClass->par,d_par,length*sizeof(void*),cudaMemcpyHostToDevice);

我在这里、这里、这里、这里和这里看到了多个类似的问题，但要么我无法理解他们遇到的问题，要么它似乎与这个特定问题不相符。

我知道我遇到的分段错误是因为gpuClass->par它位于设备上，并且 cudaMemCpy 不允许设备指针。但是，我认为没有其他方法可以将指针“插入”到对象中gpuClass。

我可以看到解决方案的方法是：

1）扁平化我的数据结构。但是，我不知道如何使用我想要的继承类功能来做到这一点。

2）gpuClass最初在gpu上实例化，我不知道该怎么做，或者

3）我在其中一个解决方案中看到，您可以使用 cudaMemCpy 将动态分配列表的地址复制到对象中，但我再次不知道该怎么做（特别是将设备指针复制到另一个设备指针的位置）。

任何帮助将不胜感激。

Answer 1

在您的第一个相关链接中，我为基于对象的深度复制序列提供了 5 个步骤，但由于您正在执行该链接中给出的示例的双指针版本，因此这种情况变得复杂。 与双指针深拷贝相关的复杂性使得通常的建议是避免它（即展平）。

我们需要对代码进行的第一个修复是正确处理数组d_par。您需要在设备上进行相应的分配来保存与 关联的数组d_par。关联的数组d_par可存储 5 个对象指针。您已经为其分配了主机端存储（使用new），但您没有为其进行设备端分配。（我不是在谈论d_par指针本身，我是在谈论它指向的内容，即一个由 5 个指针组成的数组）。

par我们需要进行的第二个修复是在顶级设备端对象中调整指针本身的修复（而不是它指向的内容）。您尝试将这两个步骤合并到一个步骤中，但这行不通。

这是代码的修改版本，似乎可以在上述更改的情况下正常工作：

$ cat t29.cu
#include <stdio.h>

class Parent {public: int my_id;};
class Child : public Parent {};

class CopyClass {
  public:
    Parent ** par;
};

const int length = 5;

__global__ void test_kernel(CopyClass *my_class){

  for (int i = 0; i < length; i++)
    printf("object: %d, id: %d\n", i, my_class->par[i]->my_id);
}

int main(){


//Instantiate object on the CPU
  CopyClass cpuClass;
  cpuClass.par = new Parent*[length];
  for(int i = 0; i < length; ++i) {
    cpuClass.par[i] = new Child;
    cpuClass.par[i]->my_id = i+1;} // so we can prove that things are working

//Allocate storage for object onto GPU and copy host object to device
  CopyClass * gpuClass;
  cudaMalloc(&gpuClass,sizeof(CopyClass));
  cudaMemcpy(gpuClass,&cpuClass,sizeof(CopyClass),cudaMemcpyHostToDevice);

//Copy dynamically allocated child objects to GPU
  Parent ** d_par;
  d_par = new Parent*[length];
  for(int i = 0; i < length; ++i) {
    cudaMalloc(&d_par[i],sizeof(Child));
    printf("\tCopying data\n");
    cudaMemcpy(d_par[i],cpuClass.par[i],sizeof(Child),cudaMemcpyHostToDevice);
  }

//Copy the d_par array itself to the device

  Parent ** td_par;
  cudaMalloc(&td_par, length * sizeof(Parent *));
  cudaMemcpy(td_par, d_par, length * sizeof(Parent *), cudaMemcpyHostToDevice);

//copy *pointer value* of td_par to appropriate location in top level object
  cudaMemcpy(&(gpuClass->par),&(td_par),sizeof(Parent **),cudaMemcpyHostToDevice);

  test_kernel<<<1,1>>>(gpuClass);
  cudaDeviceSynchronize();
  return 0;


}
$ nvcc -arch=sm_61 -o t29 t29.cu
$ cuda-memcheck ./t29
========= CUDA-MEMCHECK
        Copying data
        Copying data
        Copying data
        Copying data
        Copying data
object: 0, id: 1
object: 1, id: 2
object: 2, id: 3
object: 3, id: 4
object: 4, id: 5
========= ERROR SUMMARY: 0 errors
$