1

所以我之前问过一个关于如何直接在设备上分配对象而不是“正常”的问题：

  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. 现在您可以将该指针传递给另一个内核就好了！

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

干杯

编辑：正如罗伯特所指出的，首先在主机上创建一个指针是没有意义的，所以我从代码中删除了这部分。

4

1 回答 1

3

是的，你可以这么做。

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

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

除了通过主机端 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 的第一个参数并用局部变量替换，从而消除怪异和编译器警告。

于 2013-04-19T20:59:03.017 回答