我有一个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);


我已经看到与此hereherehereherehere相似的多个问题,但是我无法理解他们所遇到的问题,或者它似乎不适合这个特殊的问题。

我知道我得到的分段错误是因为gpuClass->par在设备上,并且cudaMemCpy不允许设备指针。但是,我看不到将指针“插入” gpuClass对象的其他方法。

我可以看到的解决方案是:

1)整理我的数据结构。但是,我不知道如何使用我想要的继承的类功能来执行此操作。

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

3)我在one of the solutions中看到可以使用cudaMemCpy将动态分配的列表的地址复制到一个对象中,但是再次,我不知道该怎么做(特别是将设备指针复制到该位置)另一个设备指针)。

任何帮助将不胜感激。

最佳答案

在您的first related link中,我为基于对象的深度复制序列提供了5个步骤,但是由于您正在对该链接中给出的示例进行双指针版本操作,因此使这种情况变得复杂。 The complexity associated with a double-pointer deep-copy is such that the usual recommendation is to avoid it(即展平)。

我们需要对您的代码进行的第一个修复是正确处理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
$

07-24 09:51
查看更多