将一个对象复制到设备？

Question

我能复制一个 C++ 对象到设备上吗？

比如说我有：

class CudaClass
{
public:
int* data;
CudaClass(int x) {
    data = new int[1]; data[0] = x;
}
};

__global__ void useClass(CudaClass cudaClass)
{
    printf("%d" cudaClass.data[0]);
};


int main()
{
    CudaClass c(1);
}

现在我该如何将“c”复制到设备内存并启动“useClass”内核程序？

- SpaceMonkey

Answer 1

是的，您可以将对象复制到设备上以在设备上使用。当对象具有指向动态分配区域的嵌入式指针时，该过程需要一些额外的步骤。

请参见我在这里的答案，了解有关所涉及内容的讨论。该答案还链接了一些示例代码答案。

另外，在类定义中，如果您希望某些函数在设备上可用，您应该适当地装饰这些函数（即可能使用__device__ __host__）；

编辑：针对一个提问（现已删除），以下是我根据提供的代码能想到的最简单的示例代码：

#include <stdio.h>

class CudaClass
{
public:
int* data;
CudaClass(int x) {
    data = new int[1]; data[0] = x;
}
};

__global__ void useClass(CudaClass *cudaClass)
{
    printf("%d\n", cudaClass->data[0]);
};




int main()
{
    CudaClass c(1);
    // create class storage on device and copy top level class
    CudaClass *d_c;
    cudaMalloc((void **)&d_c, sizeof(CudaClass));
    cudaMemcpy(d_c, &c, sizeof(CudaClass), cudaMemcpyHostToDevice);
    // make an allocated region on device for use by pointer in class
    int *hostdata;
    cudaMalloc((void **)&hostdata, sizeof(int));
    cudaMemcpy(hostdata, c.data, sizeof(int), cudaMemcpyHostToDevice);
    // copy pointer to allocated device storage to device class
    cudaMemcpy(&(d_c->data), &hostdata, sizeof(int *), cudaMemcpyHostToDevice);
    useClass<<<1,1>>>(d_c);
    cudaDeviceSynchronize();
    return 0;
}

为了简明扼要，我没有使用通常的cuda错误检查。

回答这个问题，你不能直接使用设备类中的指针从主机分配存储空间。这是因为cudaMalloc需要一个普通的基于主机的指针存储，比如你用以下方式获得的指针：

int *hostdata;

如果指针的存储已经在设备上，则cudaMalloc无法使用该指针。以下代码将无法运行：

cudaMalloc(&(d_c->data), sizeof(int));

因为在主机代码中对设备指针（d_c）进行解引用是不允许的。

- Robert Crovella

---

hostdata 包含指向设备内存的指针。但是，&hostdata 是指向主机内存位置的指针。因此，这个指针并不在设备上，并且它也不指向设备上的位置。&hostdata 指向主机内存中的一个位置，而 hostdata 则指向设备内存中的一个位置。 - Robert Crovella

对于这个问题中的这个非常简单的类，唯一需要做的就是反转复制实际数据的cudaMemcpy操作，因此，在内核之后，类似于：cudaMemcpy(c.data, hostdata, sizeof(int), cudaMemcpyDeviceToHost);。更复杂的对象可能需要几个步骤。由于直接从设备返回主机的对象本身的复制将使对象中的嵌入指针无效，因此您需要一种方法来保留或“修复”这些指针，就像这个答案演示了如何“修复”设备对象中的嵌入指针一样。 - Robert Crovella

设备分配可以在类构造函数中处理吗？ - If_You_Say_So

这对我的应用程序非常有效，谢谢。我想我还需要在d_c->data上使用cudaFree？我正在考虑将其放在CudaClass的析构函数中。假设我的CudaClass还有许多标量整数和双精度浮点数（约为百个），我是否也需要释放它们？ - rinkert

在C++中，我会说一个经验法则是，如果您的应用程序使用动态分配器为变量分配空间，那么最好的做法是确保您的应用程序也包含相应的解除分配器。严格遵守这个规则。动态分配器的例子包括malloc()、new、cudaMalloc、cudaMallocManaged等。相应的解除分配器是free()、delete和cudaFree。如果这还不清楚，请提出一个新问题。我没有试图编写生产就绪的代码。除非问题涉及到它，否则我通常不这样做。 - Robert Crovella