在 CUDA 的 __device__ 函数中使用动态分配时出现“未知错误”

Posted 2023-02-21

【中文标题】在 CUDA 的 __device__ 函数中使用动态分配时出现“未知错误”

【英文标题】："unknown error" while using dynamic allocation inside __device__ function in CUDA

【发布时间】：2014-07-17 22:34:42

【问题描述】：

我正在尝试在 CUDA 应用程序中实现一个链表来为不断增长的网络建模。为此，我在__device__ 函数中使用malloc，旨在在全局内存中分配内存。 代码是：

void __device__ insereviz(Vizinhos **lista, Nodo *novizinho, int *Gteste)

   Vizinhos *vizinho;

   vizinho=(Vizinhos *)malloc(sizeof(Vizinhos));

   vizinho->viz=novizinho;

   vizinho->proxviz=*lista;

   *lista=vizinho;

   novizinho->k=novizinho->k+1;

在分配一定数量的元素（大约 90000 个）后，我的程序返回“未知错误”。起初我虽然这是一个内存限制，但我检查了nvidia-smi，我得到了

+------------------------------------------------------+                       
| NVIDIA-SMI 331.38     Driver Version: 331.38         |                       
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|===============================+======================+======================|
|   0  GeForce GTX 770     Off  | 0000:01:00.0     N/A |                  N/A |
| 41%   38C  N/A     N/A /  N/A |    159MiB /  2047MiB |     N/A      Default |
+-------------------------------+----------------------+----------------------+

所以这似乎不是内存问题，除非malloc 在共享内存中分配。为了测试这一点，我尝试在单独的块中运行两个网络，但我能够分配的结构数量仍然有限。但是，当我尝试使用较少数量的结构运行同一程序的两个实例时，它们都可以顺利完成。

我也试过cuda-memcheck得到了

========= CUDA-MEMCHECK
========= Invalid __global__ write of size 8
=========     at 0x000001b0 in     /work/home/melo/proj_cuda/testalloc/cuda_testamalloc.cu:164:insereviz(neighbor**, node*, int*)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x00000000 is out of bounds
=========     Device Frame:/work/home/melo/proj_cuda/testalloc/cuda_testamalloc.cu:142:insereno(int, int, node**, node**, int*) (insereno(int, int, node**, node**, int*) : 0x648)
=========     Device Frame:/work/home/melo/proj_cuda/testalloc/cuda_testamalloc.cu:111:fazrede(node**, int, int, int, int*) (fazrede(node**, int, int, int, int*) : 0x4b8)
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/libcuda.so.1 (cuLaunchKernel + 0x331) [0x138281]
=========     Host Frame:gpu_testamalloc5 [0x1bd48]
=========     Host Frame:gpu_testamalloc5 [0x3b213]
=========     Host Frame:gpu_testamalloc5 [0x2fe3]
=========     Host Frame:gpu_testamalloc5 [0x2e39]
=========     Host Frame:gpu_testamalloc5 [0x2e7f]
=========     Host Frame:gpu_testamalloc5 [0x2c2f]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xfd) [0x1eead]
=========     Host Frame:gpu_testamalloc5 [0x2829]

内核启动是否有任何限制或我缺少什么？如何查看？

谢谢，

里卡多

【问题讨论】：

为什么不检查 malloc 返回的值的有效性？

@talonmies 好问题 :-) 吸取教训，谢谢！

【参考方案1】：

最可能的原因是“设备堆”上的空间不足。这最初默认为 8MB，但您可以更改它。

参考documentation，我们看到设备malloc在设备堆外分配。

如果发生错误，malloc 将返回一个 NULL 指针。最好在设备代码中测试这个 NULL 指针（和主机代码——在这方面它与主机 malloc 没有什么不同）。如果你得到一个 NULL 指针，你已经用完了设备堆空间。

如文档中所述，设备堆的大小可以在内核调用之前通过以下方式调整：

cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)

运行时 API 函数。

如果您忽略所有这些并尝试使用返回的 NULL 指针，您将在设备代码中获得无效访问，如下所示：

=========     Address 0x00000000 is out of bounds

【讨论】：

非常感谢！我完全解决了我的问题！