复制到 malloc() 分配的全局内存？答案

【问题标题】：Copy to global memory allocated by malloc()?复制到 malloc() 分配的全局内存？
【发布时间】：2012-08-28 08:25:25
【问题描述】：

CUDA 编程指南指出“可以使用运行时复制通过malloc() 分配的内存（即，通过从设备内存调用任何复制内存函数）”，但不知何故我无法重现此功能。代码：

#include <cstdio>
__device__ int* p;

__global__ void allocate_p() {
  p = (int*) malloc(10);
  printf("p = %p  (seen by GPU)\n", p);
}

int main() {
  cudaError_t err;
  int* localp = (int*) malloc(10);

  allocate_p<<<1,1>>>();
  cudaDeviceSynchronize();

  //Getting pointer to device-allocated memory
  int* tmpp = NULL;
  cudaMemcpyFromSymbol(&tmpp, p, 4);
  printf("p = %p  (seen by CPU)\n", tmpp);

  //cudaMalloc((void**)&tmpp, 40);
  err = cudaMemcpy(tmpp, localp, 40, cudaMemcpyHostToDevice);
  cudaDeviceSynchronize();
  printf(" err:%i %s", (int)err, cudaGetErrorString(err));

  delete localp;
  return 0;
}

输出崩溃：

p = 0x601f920  (seen by GPU)
p = 0x601f920  (seen by CPU)
 err:11 invalid argument

我推测，主机在设备上看到了适当的地址，但不知何故不喜欢它来自malloc()。

如果我之前通过cudaMalloc((void**)&np, 40); 分配，然后将指针np 作为参数传递给内核allocate_p，它将被分配给p（而不是malloc()），那么代码运行良好。

我做错了什么/我们如何在主机端函数中使用 malloc() 分配的设备内存？

【问题讨论】：

您引用的部分来自“B.17.2 与主机内存 API 的互操作性”部分。 malloc() 是指主机 malloc() 而不是设备 malloc()。这绝对是令人困惑的，因为它遵循介绍设备 malloc() 的“B.17.1 堆内存分配”部分。
@GregSmith：我无法理解您是如何得出这种解释的。该部分是关于设备动态内存分配的（甚至还有显示内核 malloc 调用的代码示例），并且参考了第 3.2.2 节中描述的主机内存 API 函数，即。 cudaMemcpy。 host malloc 在哪里出现？我不会说它写得很混乱，我会说它完全是错误的（或者在运行时有一个巨大的错误作为替代方案）。
@talonmies 是的，该部分在旧手册中是错误的，并且被复制到不同的位置。该语句应从下一版本的编程指南中删除。

标签： pointers memory-management cuda

【解决方案1】：

据我所知，无法使用主机 API 函数复制运行时堆内存。这在 CUDA 4.x 中肯定是不可能的，而 CUDA 5.0 候选版本并没有改变这一点。我能提供的唯一解决方法是使用内核“收集”最终结果并将它们填充到设备传输缓冲区或零拷贝内存中，可以通过 API 或直接从主机访问。您可以在this answer 和another question 中看到这种方法的示例，其中来自 NVIDIA 的 Mark Harris 确认这是 CUDA 运行时中（当时）当前实现的限制。

【讨论】：

另一种解决方法是根本不使用__device__ malloc()，而是预先分配内存并在内核中对其进行索引，对吗？特别是如果__device__ malloc() 的大小是恒定的，如示例中所示。另外，我认为__device__ malloc() 可能会导致同时进行调用的线程的隐式序列化（以便可以从堆中为每个线程分配自己的块）。
@RogerDahl：在大多数情况下，我会同意。在某些情况下，最好使用“设置”内核使用内核 malloc 在堆中构建数据结构，然后让设备代码纯粹在堆中的该结构上工作，然后最后将最终结果复制回主机“下拉/收集”内核。仅使用主机 API 基本上不可能分配和初始化深度递归数据结构。
谢谢；这让我发疯（鉴于指南所说）。由于我在主机上有设备地址（在tmpp 中），我尝试将从 cudaMalloc() 获得的其他指针移动到指向 malloc() 分配的设备内存。（这很愚蠢，但我认为也许 CUDA 有一些 RTTI 来区分指针。）显然主机端分配器范围检查内存操作到它显式分配的内容。另一方面，设备端分配器非常松散，甚至不会费心报告未分配的s 的free(s)。我正在使用 -arch=sm_21 在 CUDA 5.0 上对此进行测试。
@PMarecki：零拷贝是映射的主机内存，GPU 可以通过 PCI-e 总线直接访问它，而不需要显式的 API 拷贝（因此称为“零拷贝”）。
@harrism：CUDA 4.2 编程指南的第 B.17.2 节，最后一句。我在搜索之前从未注意到它......