不能同时使用 cuMemcpyHtoDAsync 和 cuMemcpyDtoHAsync答案

【问题标题】：Cannot use cuMemcpyHtoDAsync and cuMemcpyDtoHAsync at the same time不能同时使用 cuMemcpyHtoDAsync 和 cuMemcpyDtoHAsync
【发布时间】：2020-09-22 00:54:04
【问题描述】：

我对下面的代码 sn-p 有一个相当奇怪的观察。

当我两者都做时 - 将内存复制到设备并将结果复制回主机时，流似乎是同步的 - 即它们按顺序执行内核。一旦我将副本删除到主机并将参数复制到设备，流将并行执行，一旦我删除复制参数并继续复制结果，流也会并行执行。

知道为什么吗？以及如何解决这个问题？

for (int j=0; j<n_streams; j++) {
    cuMemcpyHtoDAsync(gpu_parameters[j], parameters[j].asPointer(), (parameterCount) * Sizeof.FLOAT, stream[j]);
    Pointer kernelParameters1 = Pointer.to(
            Pointer.to(new int[]{0}),
            Pointer.to(new int[] {10000}),
            Pointer.to(gpu_data),
            Pointer.to(gpu_results[j]),
            Pointer.to(gpu_parameters[j])
            );
    cuLaunchKernel(function[j],
            s_grid, 1, 1,      // Grid dimension
            s_block, 1, 1,      // Block dimension
            0, stream[j],               // Shared memory size and stream
            kernelParameters1, null // Kernel- and extra parameters
            );
    cuMemcpyDtoHAsync(results[j].asPointer(), gpu_results[j], (results[j].size()) * Sizeof.FLOAT, stream[j]);
}

【问题讨论】：

除非你有一个带有两个复制引擎的 GPU，否则这是一个硬件限制，不能神奇地编程
哦...即使内核的运行时间是复制所需时间的 1000 倍？
是的：设备查询告诉我：“并发复制和内核执行：是的，有 1 个复制引擎”
最终，两个相反方向的传输将相互阻塞，然后事情将序列化，直到两者都从流中弹出。如果您有两个 dma 引擎，那么两个方向可以重叠并且将消除失速
这对我来说仍然不明显，为什么会发生这种情况。只要复制只在一个方向上完成，使用 1 个流对内核进行相同数量的调用所花费的时间大约是使用 2 个流的两倍。似乎切换内存传输方向会进行设备同步....

标签： performance cuda gpu jcuda

【解决方案1】：

不知道为什么...但是更改顺序消除了问题-并且正在并行执行...。

for (int j=0; j<n_streams; j++) {
    cuMemcpyHtoDAsync(gpu_parameters[j], parameters[j].asPointer(), (parameterCount) * Sizeof.FLOAT, stream[j]);
}
for (int j=0; j<n_streams; j++) {
    Pointer kernelParameters1 = Pointer.to(
            Pointer.to(new int[]{0}),
            Pointer.to(new int[] {getNPrices()}),
            Pointer.to(get_gpu_prices()),
            Pointer.to(gpu_results[j]),
            Pointer.to(gpu_parameters[j])
            //,Pointer.to(new int[]{0})
            );
    cuLaunchKernel(function[j],
            s_grid, 1, 1,      // Grid dimension
            s_block, 1, 1,      // Block dimension
            0, stream[j],               // Shared memory size and stream
            kernelParameters1, null // Kernel- and extra parameters
            );
}
for (int j=0; j<n_streams; j++) {
    cuMemcpyDtoHAsync(results[j].asPointer(), gpu_results[j], (results[j].size()) * Sizeof.FLOAT, stream[j]);
}

【讨论】：

排队系统中的深度优先与广度优先行为。
我仍然想了解为什么内核在第一次尝试中被序列化执行。不幸的是，我的实际用例略有不同——我几乎无法控制将命令添加到流中的顺序，因为每个流都有一个主机线程。因此调度 copyToDevice-executeKernel-copyToHost 总是会一个接一个地快速完成
流数：尝试 2....10 - 总是相同的现象。关于每个流由单独的线程填充的用例：在调度 cuMemcpyDtoHAsync 之前，我为内核执行添加了大约 60% 的预期执行时间的 Thread.sleep() 后，流（主要）并行运行。