CUDA 内核代码的设备内存：它是不是明确可管理？

Posted 2023-04-15

【中文标题】CUDA 内核代码的设备内存：它是不是明确可管理？

【英文标题】：Device memory for CUDA kernel code: Is it explicitly manageable?CUDA 内核代码的设备内存：它是否明确可管理？

【发布时间】：2011-11-01 01:16:00

【问题描述】：

上下文： CUDA 4.0，Linux 64 位，NVIDIA UNIX x86_64 内核模块 270.41.19，在 GeForce GTX 480 上。

我尝试在我的程序中查找（设备）内存泄漏。我使用运行时 API 和 cudaGetMemInfo(free,total) 来测量设备内存使用情况。我注意到内核执行后有重大损失（在本例中为 31M）。内核代码本身不分配任何设备内存。所以我猜它是保留在设备内存中的内核代码。即使我会认为内核没有那么大。 （有没有办法确定内核的大小？）

内核代码何时加载到设备内存中？我猜在执行主机代码行时：

kernel<<<geom>>>(params);

对吗？ 调用后代码是否保留在设备内存中？如果是这样，我可以显式卸载代码吗？

我担心的是设备内存碎片。想想大量交替的设备内存分配和内核执行（不同的内核）。然后过了一段时间设备内存变得非常稀缺。即使你释放了一些内存，内核代码仍然只留下内核之间的空间用于新的分配。这会在一段时间后导致巨大的内存碎片。这就是 CUDA 的设计方式吗？

【问题讨论】：

你在使用全局内存吗？全局内存将保持分配状态，直到您关闭设备或 cuda 释放它（如果它是 cudamalloc'ed）

内核被定义为“global”，但没有其他的。

这里的全局并不意味着__global__ 关键字，它只是意味着在.cu 文件中的全局范围内声明的任何内容或通过cudamalloc 分配的任何内存。因此，例如，如果您在 .cu 文件中的全局范围内有 int8[1024]，或者如果您在没有匹配 cudaFree 的情况下调用 cudaMalloc(8 * 1024)，您将拥有 1KB 的显着内存占用。

【参考方案1】：

您正在观察的内存分配由 CUDA 上下文使用。它不仅包含内核代码，它还包含任何其他静态范围设备符号、纹理、本地内存的每线程暂存空间、printf 和堆、常量内存，以及驱动程序和 CUDA 运行时本身所需的 gpu 内存。当加载二进制模块或 PTX 代码由驱动程序 JIT 编译时，大部分内存只分配一次。最好将其视为固定开销，而不是泄漏。 PTX 代码中有 200 万条指令限制，而当前硬件使用 32 位字来表示指令，因此即使是最大允许内核代码的内存占用量与它所需的其他全局内存开销相比也很小。

在最新版本的 CUDA 中，有一个运行时 API 调用 cudaDeviceSetLimit，它允许对给定上下文可以消耗的暂存空间量进行一些控制。请注意，可以将限制设置为低于设备代码要求的值，在这种情况下，可能会导致运行时执行失败。

【讨论】：

感谢您的回答！您知道内核代码是否保留在该内存空间中直到上下文重置？还是它就像内核代码的缓存一样？如果一个程序有大量不同的内核，它们都留在内存中怎么办？

没有记录。当一个模块被加载到一个上下文中时，所有外部可见的符号和架构兼容的二进制有效载荷都会被加载。但我不知道他们是否在加载时直接进入设备。