CUDA - 复制整个托管内存类的性能损失有多严重

Posted 2023-02-17

【中文标题】CUDA - 复制整个托管内存类的性能损失有多严重

【英文标题】：CUDA - how bad is the performance penalty for copying entire class of managed memory

【发布时间】：2017-12-13 13:14:28

【问题描述】：

我有一些从托管内存分配器派生的类，例如：

/* --------- managed 来自https://devtalk.nvidia.com/default/topic/987577/-thrust-is-there-a-managed_vector-with-unified-memory-do-we-still-only-have-device_vector-cuda-thrust-managed-vectors-/

------------ 它覆盖新的操作符，执行 cudaMallocManaged 然后强制转换 */

class Cell : public Managed 
    int a;float b;char c; // say ~50 fields
 

现在，假设我有一个包含 100,000 个 Cell 对象的数组，并且想要发送到某个全局函数，该函数仅使用一小部分（例如 5-10 个）字段来进行一些计算。

最简单的方法是发送整个单元格对象数组。但是，它确实会复制大量未使用的数据。

一种更严格的方法是只分配所需的 5-10 个字段的设备数组，复制值并将它们发送到全局函数。这有点烦人，因为如果全局函数体需要单元类中的一些其他字段，则必须重写其签名以接受新数组。

我的问题 - 一般来说，使用最简单方法的性能损失有多严重？

谢谢！

【问题讨论】：

“发送整个数组”是什么意思？如果您使用托管内存，则不会“发送”任何内容。驱动程序和设备正在协调设备通过 PCI-E 接口进行的 ad-hoc 访问。

"sending" - 将“Cell*”对象的数组传递给全局函数。我不知道（我猜它的实现取决于？）引擎盖下发生了什么。

同样，将指向内核的指针作为参数传递是按值传递，没有数据复制或底层结构数组的传输。您只传递了一个 64 位地址。

【参考方案1】：

托管内存的处理方式非常多depends on the compute capability of you device。 Pascal (6.x) 及更高版本将仅在那些被访问的页面中请求分页。

计算能力较低的设备通常会传输整个托管内存集，无论它有多少被访问，甚至是否被访问。 但是你可以explicitly declare the memory regions to transfer on a per-stream basis using cudaStreamAttachMemAsync()。这使您可以限制传输的数据量，而无需更改分配或数据结构。