将重要的类复制到设备答案

【问题标题】：Copy non-trivial classes to Device将重要的类复制到设备
【发布时间】：2022-01-09 22:23:02
【问题描述】：

我有一个这样的课程：

class CudaArray
{
   CudaArray() : Ptr(new double[5]) {}
   double* Ptr;
   int Dimension;
}

然后是另一个类，例如：

class Container
{
   short a;
   CudaArray* ArrayPtr;
   int b;
   int c;
}

现在我正在以这种方式在设备上创建数组：

CudaArray H_Array;
CudaArray* D_Array;

Check(cudaMalloc(&D_Array, sizeof(CudaArray)));
Check(cudaMemcpy(D_Array, &H_Array, sizeof(CudaArray), cudaMemcpyHostToDevice));

double* Tmp;
Check(cudaMalloc(&Tmp, sizeof(double) * 5));
Check(cudaMemcpy(Tmp, H_Array.Ptr, sizeof(double) * 5, cudaMemcpyHostToDevice));
Check(cudaMemcpy(&(D_Array->Ptr), &Tmp, sizeof(double*), cudaMemcpyHostToDevice));

我希望能够在设备代码上使用Container 类型的对象，但我无法从现有数组初始化CudaArray 成员。到目前为止我试过了：

Container* Cont = nullptr;

Check(cudaMalloc(&Cont , sizeof(Container)));
Check(cudaMemcpy(&(Cont->ArrayPtr), &D_Array, sizeof(CudaArray*), cudaMemcpyDeviceToDevice));

但我在最后一个 cudaMemcpy 上得到 GPUassert: invalid argument。

如何初始化包含指向设备内存中现有对象（类）的指针的设备类？

此外，是否有更简单或更优雅的方式在主机和设备之间复制复杂对象？

【问题讨论】：

这是一个常见问题。 SO cuda 标签上已经有很多问题解释了如何在主机和设备之间复制带有嵌入式指针的类。
我阅读了你的几个答案，但遗憾的是它们并没有完全涵盖我在这里寻找的内容......
简单地说，你不能在构造函数中通过 new 使用像 CudaArray 这样的类来分配内存

标签： cuda

【解决方案1】：

Check(cudaMemcpy(&(Cont->ArrayPtr), &D_Array, sizeof(CudaArray*), cudaMemcpyDeviceToDevice));

但我得到 GPUassert: invalid argument on the last cudaMemcpy.

错误是您指定了cudaMemcpyDeviceToDevice，但&D_Array 是主机内存中的一个位置：

CudaArray* D_Array;

你应该使用cudaMempcyHostToDevice。

如何初始化包含指向设备内存中现有对象（类）的指针的设备类？

这种变化似乎为我解决了问题：

$ cat t174.cu
#include <cstdio>
class CudaArray
{
  public:
   CudaArray() : Ptr(new double[5]) {}
   double* Ptr;
   int Dimension;
};

class Container
{
  public:
   short a;
   CudaArray* ArrayPtr;
   int b;
   int c;
};
#define Check(x) x
__global__ void k(Container *c){

  printf("%f\n", c->ArrayPtr->Ptr[0]);
}
int main(){

CudaArray H_Array;
H_Array.Ptr[0] = 1234.0;
CudaArray* D_Array;

Check(cudaMalloc(&D_Array, sizeof(CudaArray)));
Check(cudaMemcpy(D_Array, &H_Array, sizeof(CudaArray), cudaMemcpyHostToDevice));

double* Tmp;
Check(cudaMalloc(&Tmp, sizeof(double) * 5));
Check(cudaMemcpy(Tmp, H_Array.Ptr, sizeof(double) * 5, cudaMemcpyHostToDevice));
Check(cudaMemcpy(&(D_Array->Ptr), &Tmp, sizeof(double*), cudaMemcpyHostToDevice));

Container* Cont = nullptr;

Check(cudaMalloc(&Cont , sizeof(Container)));
Check(cudaMemcpy(&(Cont->ArrayPtr), &D_Array, sizeof(CudaArray*), cudaMemcpyHostToDevice));
k<<<1,1>>>(Cont);
cudaDeviceSynchronize();
}
$ nvcc -o t174 t174.cu
$ cuda-memcheck ./t174
========= CUDA-MEMCHECK
1234.000000
========= ERROR SUMMARY: 0 errors
$

此外，是否有更简单或更优雅的方式在主机和设备之间复制复杂对象？

可能有performance impacts，但从代码复杂性的角度来看，如果您通过托管分配器进行所有分配，事情可能会更简单（简单是旁观者):

$ cat t175.cu
#include <cstdio>
#include <new>
class CudaArray
{
  public:
   CudaArray()  {cudaMallocManaged(&Ptr, 5*sizeof(double)); for (int i = 0; i < 5; i ++) Ptr[i] = 0.0;}
   double* Ptr;
   int Dimension;
};

class Container
{
  public:
   short a;
   CudaArray* ArrayPtr;
   int b;
   int c;
};
#define Check(x) x
__global__ void k(Container *c){

  printf("%f\n", c->ArrayPtr->Ptr[0]);
}
int main(){

CudaArray *my_Array;
cudaMallocManaged(&my_Array, sizeof(CudaArray));
new(my_Array) CudaArray();
my_Array[0].Ptr[0] = 1234.0;

Container* Cont = nullptr;

Check(cudaMallocManaged(&Cont , sizeof(Container)));
Cont[0].ArrayPtr = my_Array;
k<<<1,1>>>(Cont);
cudaDeviceSynchronize();
}
$ nvcc -o t175 t175.cu
$ cuda-memcheck ./t175
========= CUDA-MEMCHECK
1234.000000
========= ERROR SUMMARY: 0 errors
$

【讨论】：