这种分配设备对象的方式“正确”吗？

Question

SO I asked a question 之前关于如何直接在设备上分配对象而不是“normal”：

1. 在主机上分配
2. 复制到设备
3. 将动态分配的字段一一复制到设备

我希望它直接在设备上分配的主要原因是我不想手动将每个动态分配的字段一一复制。

无论如何，我认为我实际上已经找到了一种方法，我希望看到更有经验的 CUDA 程序员（如 Robert Crovella）的一些意见。

我们先看代码：

class Particle
{
    public:
    int *data;

    __device__ Particle()
    {
        data = new int[10];
        for (int i=0; i<10; i++)
        {
            data[i] = i*2;
        }
    }
};


__global__ void test(Particle **result)
{
    Particle *p = new Particle();

    result[0] = p; // store memory location
}

__global__ void test2(Particle *p)
{
    for (int i=0; i<10; i++)
        printf("%d\n", p->data[i]);

}

int main() {
    // initialise and allocate an object on device
    Particle **d_p_addr;
    cudaMalloc((void**)&d_p_addr, sizeof(Particle*));
    test<<<1,1>>>(d_p_addr);

    // copy pointer to host memory
    Particle  **p_addr = new Particle*[1];
    cudaMemcpy(p_addr, d_p_addr, sizeof(Particle*), cudaMemcpyDeviceToHost);

    // test:
    test2<<<1,1>>>(p_addr[0]);

    cudaDeviceSynchronize();

    printf("Done!\n");

}

如你所见，我所做的是：

1. 调用内核初始化设备上的对象并将其指针存储为输出参数
2. 将指向分配对象的指针从设备内存复制到主机内存
3. 现在您可以将该指针传递给另一个内核！

这段代码确实有效，但我不确定是否有缺点。

干杯

编辑：正如 Robert 所指出的，首先在主机上创建指针没有意义，所以我从代码中删除了该部分。

Answer 1

是的，你可以这样做。

您正在设备上分配一个对象，并将指向它的指针从一个内核传递到下一个内核。由于device malloc/new 的一个特点是分配在上下文的生命周期内（不仅仅是内核）持续存在，因此分配不会在内核结束时消失。这基本上是标准的 C++ 行为，但我认为它可能值得重复。因此，您从一个内核传递到下一个内核的指针在您的程序上下文中的任何后续设备代码中都是有效的。

不过，您可能需要注意一个问题。在设备上完成的动态分配返回的指针（例如通过设备代码中的new 或malloc）不可用于将数据从设备传输到主机，至少在当前的 cuda 化身中（cuda 5.0 及更早版本）。造成这种情况的原因有些晦涩（翻译：我无法充分解释），但考虑动态分配来自设备堆这一事实是有启发性的，设备堆是一个逻辑上与运行时全局内存区域分开的区域像 cudaMalloc 和 cudaMemcpy 这样的 API 函数使用。对此的间接指示是here:

为设备堆保留的内存是通过主机端 CUDA API 调用（例如 cudaMalloc()）分配的内存之外的。

如果您想向自己证明这一点，请尝试在您的第二次内核调用之后添加以下看似无害的代码：

Particle *q;
q = (Particle *)malloc(sizeof(Particle));
cudaMemcpy(q, p_addr[0], sizeof(Particle), cudaMemcpyDeviceToHost);

如果您随后检查从该 cudaMemcpy 操作返回的 API 错误值，您将观察到错误。

作为一个不相关的评论，在我的书中，您对指针 *p 的使用有点怪异，并且给出的编译器警告表明它很奇怪。这在技术上并不是非法的，因为你实际上并没有用那个指针做任何有意义的事情（你立即在你的内核 1 中替换它）但是它很奇怪，因为你正在传递一个指向你没有正确 cudaMalloc'ed 的内核的指针.在您演示的上下文中，这是完全没有必要的，您可以消除内核 1 的第一个参数并用局部变量替换，从而消除怪异和编译器警告。