我有一个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);
我已经看到与此here,here,here,here和here相似的多个问题,但是我无法理解他们所遇到的问题,或者它似乎不适合这个特殊的问题。
我知道我得到的分段错误是因为
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
$