cuda - 关于 cudaMemcpyAsync 函数

Question

我有一些问题。

最近我正在使用CUDA制作一个程序。

在我的程序中，主机上有一个用 std::map(string, vector(int)) 编程的大数据。

通过使用这些数据，一些向量(int)被复制到 GPU 全局内存并在 GPU 上处理

处理后，在 GPU 上生成一些结果，并将这些结果复制到 CPU。

这些都是我的节目表。

cudaMemcpy( ... , cudaMemcpyHostToDevice)

内核函数(内核函数只有在将必要的数据复制到 GPU 全局内存时才能执行)

cudaMemcpy( ... , cudaMemcpyDeviceToHost)

重复 1~3 步 1000 次(对于另一个数据(向量))

但我想减少处理时间。

所以我决定在我的程序中使用 cudaMemcpyAsync 函数。

在搜索了一些文档和网页后，我意识到要使用 cudaMemcpyAsync 函数主机内存，必须将数据复制到 GPU 全局内存作为固定内存分配。

但是我的程序正在使用 std::map，所以我无法将此 std::map 数据制作到固定内存。

因此，我没有使用它，而是制作了一个缓冲区数组类型的固定内存，这个缓冲区总是可以处理复制向量的所有情况。

最后，我的程序是这样工作的。

Memcpy(使用循环将数据从 std::map 复制到缓冲区，直到将整个数据复制到缓冲区)

cudaMemcpyAsync( ... , cudaMemcpyHostToDevice)

内核(内核函数只有在将整个数据复制到 GPU 全局内存时才能执行)

cudaMemcpyAsync( ... , cudaMemcpyDeviceToHost)

重复 1~4 步 1000 次(对于另一个数据(向量))

我的程序变得比以前的情况快得多。

但问题(我的好奇心)就在这一点上。

我试图以类似的方式制作另一个程序。

Memcpy(仅针对一个向量将数据从 std::map 复制到缓冲区)

cudaMemcpyAsync( ... , cudaMemcpyHostToDevice)

循环 1~2，直到将整个数据复制到 GPU 全局内存

内核(内核函数只有在将必要的数据复制到 GPU 全局内存时才能执行)

cudaMemcpyAsync( ... , cudaMemcpyDeviceToHost)

重复 1~5 步 1000 次(对于另一个数据(向量))

这种方法比上面讨论的方法快了大约 10%。

但我不知道为什么。

我认为 cudaMemcpyAsync 只能与内核函数重叠。

但我的情况我认为不是。而不是看起来可以在 cudaMemcpyAsync 函数之间重叠。

对不起，我的问题很长，但我真的很想知道为什么。

有人可以教我或向我解释什么是“cudaMemcpyAsync”的确切设施以及哪些功能可以与“cudaMemcpyAsync”重叠？

Answer 1

cudaMemcpyAsync 的复制事件(以及内核事件)可以与任何主机代码重叠。此外，与设备之间的数据复制(通过 cudaMemcpyAsync)可以与内核事件重叠。所有 3 个事件:主机事件、数据复制事件和内核事件，可以相互异步完成，并且可以相互重叠。

正如您所看到和演示的，主机事件和数据复制或内核事件可以以一种相对简单的方式相互重叠:内核启动立即返回到主机，cudaMemcpyAsync 也是如此。然而，为了在数据复制和内核事件之间获得最佳的重叠机会，有必要使用一些额外的概念。为了获得最佳重叠机会，我们需要:

固定的主机内存缓冲区，例如通过 cudaHostAlloc()

使用 cuda 流来分离各种类型的事件(数据复制和内核计算)

使用 cudaMemcpyAsync(而不是 cudaMemcpy)

当然，您的工作也需要以可分离的方式分解。这通常意味着，如果您的内核正在执行特定功能，您可能需要多次调用该内核，以便每次调用都可以处理单独的数据片段。例如，这允许我们在第一次内核调用正在处理数据块 A 时将数据块 B 复制到设备。这样做我们有机会将数据块 B 的副本与数据块 A 的内核处理重叠。

与 cudaMemcpyAsync(与 cudaMemcpy 相比)的主要区别在于:

它可以在任何流中发出(它需要一个流参数)

通常，它会立即将控制权返回给主机(就像内核调用那样)，而不是等待数据复制完成。

第 1 项是必要的功能，因此数据复制可以与内核计算重叠。第 2 项是必要的功能，以便数据复制可以与主机事件重叠。

尽管复制/计算重叠的概念非常简单，但在实践中实现需要一些工作。如需其他引用资料，请参阅:

Overlap copy/compute section CUDA 最佳实践指南。

显示 basic implementation of copy/compute overlap 的示例代码.

显示完整 multi/concurrent kernel copy/compute overlap scenario 的示例代码.

请注意，上面的一些讨论是基于具有计算能力 2.0 或更高版本的设备(例如并发内核)。此外，不同的设备可能有一个或两个复制引擎，这意味着只能在某些设备上同时复制到设备和从设备复制。