在CUDA中，如何将设备指针的数组复制到设备内存？

Question

例如，我分配以下指针：

float *data_1, *data_2, *data_3, *data_4;

//Use malloc to allocate memory and fill out some data to these pointers
......
//Filling complete

float *data_d1,*data_d2,*data_d3,*data_d4;

cudaMalloc((void **)&data_d1,size1);
cudaMalloc((void **)&data_d2,size2);
cudaMalloc((void **)&data_d3,size3);
cudaMalloc((void **)&data_d4,size4);

cudaMemcpy(data_d1,data_1,size1,cudaMemcpyHostToDevice);
cudaMemcpy(data_d2,data_2,size2,cudaMemcpyHostToDevice);
cudaMemcpy(data_d3,data_3,size3,cudaMemcpyHostToDevice);
cudaMemcpy(data_d4,data_4,size4,cudaMemcpyHostToDevice);

在此之后，我应该已经获得了4个设备指针，其中包含与主机指针完全相同的数据。现在我想将这些指针存储到指针数组中，如下所示，

float *ptrs[4];

ptrs[0] = data_d1;
ptrs[1] = data_d2;
ptrs[2] = data_d3;
ptrs[3] = data_d4;

现在我想把这个指针数组传递给CUDA内核。但是，我知道由于ptrs[4]实际上位于主机内存中，我需要在设备上分配一个新的指针。所以我这样做了：

float **ptrs_d;
size_t size = 4 * sizeof(float*);
cudaMalloc((void ***)&ptrs_d,size);
cudaMemcpy(ptrs_d,ptrs,size,cudaMemcpyHostToDevice);

然后调用内核：

kernel_test<<<dimGrid,dimBlock>>>(ptrs_d, ...);
//Declaration should be 
//__global__ void kernel_test(float **ptrs_d, ...);

在 kernel_test 中，使用以下语法加载数据：

if (threadIdx.x < length_of_data_1d)
{
    float element0 = (ptrs[0])[threadIdx.x];
}

编译没有问题，但在调试时，出现了访问冲突的错误。

也许我的代码中有很多错误。但是我只想知道为什么不能以这种方式传递设备指针，如果在CUDA中允许将设备指针数组传递给内核函数，正确的访问方式是什么。

那么我该如何解决这个问题？欢迎提供任何建议。提前感谢。

- Coding_new_bird

---

我没有看到你的方法有任何明显的问题。我围绕你展示的内容构建了一个简单的代码，并且对我来说似乎工作得正确，它在这里（http://pastebin.com/n1S63xLb）。你的访问违规可能只是基于数据长度和一些你没有展示的代码而导致的数组越界。它可能与你复制设备指针数组的基本方法无关。我建议你提供一个完整的代码来重现问题，而不是一系列片段。问题出在你这里没有展示的东西。 - Robert Crovella

非常感谢您的建议。我已经试过几次调试，最终学会了这种方法是可行的。正如您所建议的那样，真正的问题似乎很可能是超出界限的问题，现在我正在努力解决。再次感谢您的快速帮助。 - Coding_new_bird

当然，使用调试器运行代码应该可以让您了解出现了什么问题。使用 cuda-memcheck 运行代码也可能有助于查明问题所在。 - Robert Crovella

Answer 1

一个可能的解决方案是分配一个空指针，就像CUDA也期望这样。当将其传递到您的内核中时，您可以将其转换为float **。我是这样做的：

void* ptrs_d = 0;
cudaMalloc(&ptrs_d, 4*sizeof(float*));
cudaMemcpy(ptrs_d, ptrs, 4*sizeof(float*), cudaMemcpyHostToDevice);
kernel_test<<<dimGrid, dimBlock>>>((float**)ptrs_d);

- hubs

---

这个解决方案解决了什么问题？ - Robert Crovella

@robert 可以将指针数组传递到 CUDA 内核中。 - hubs

不需要将浮点数转换为void，然后再转换回浮点数，以完成所需的操作。我在我的评论链接中提供的代码演示了这一点。 - Robert Crovella

谢谢您的回答。结果表明该方法是正确的，尽管似乎还存在其他问题，我仍在努力解决。致意。 - Coding_new_bird