如何使用 CUDA/Thrust 将阵列的不相交子集从设备异步复制到主机？答案

【问题标题】：How to asynchronously copy a disjoint subset of an array from device to host with CUDA/Thrust?如何使用 CUDA/Thrust 将阵列的不相交子集从设备异步复制到主机？
【发布时间】：2021-04-10 22:53:24
【问题描述】：

我不确定这是否可能以异步方式进行，但我想做的是以下内容。假设我在设备上有以下数组：

d_arr = [0, 1, 2, 3, 4, 5, 6, 7, 8, 9]

还假设我有一个大小为3 的主机数组h_arr。最后，假设我有以下指针数组：

p_arr = [&d_arr[0], &d_arr[4], &d_arr[8]]

我想调用一个虚函数

cudaMemcpyAsyncDisjoint(&d_arr[0], &h_arr[0], &p_arr[0], 3)

然后填充数组h_arr，使其变为

h_arr = [0, 4, 8]

我想异步执行此操作，因为我主要关心的是速度，因为我有一个驱动程序方法，它在循环中运行内核，然后在每轮结束时复制回数据。

【问题讨论】：

thrust::gather 已经这样做了
这种操作一般称为gather（推力、MPI等）。通常，您希望先收集到设备内存的另一段，然后将其复制回主机，因为传输到主机的成本很高。默认情况下，当您将主机向量和设备向量放入收集算法时，我不确定 Thust 会做什么。如果这完全可行，您可能需要明确指定执行策略以明确实际收集发生的位置。

标签： cuda thrust

【解决方案1】：

如果p_arr 的元素之间的步幅是恒定的，那么这在a single operation 和cudaMemcpy2DAsync 中是可能的。

对于不同的步幅，它不可能在单个操作中。此外，单一操作方法（具有恒定步幅）不一定是最快的方法（cudaMemcpy2DAsync 方法不一定接近预期的总线传输速度）。对于最快的方法，以及处理要复制的元素之间不同步幅的能力，通常的建议是将其分为两个步骤。

执行thrust::gather（或带有置换迭代器的thrust::copy）之类的操作，将所有要复制的元素收集到一个连续的临时设备缓冲区中
然后使用cudaMemcpyAsync 将该缓冲区复制到主机。

【讨论】：