了解内存传输性能 (CUDA)
Tre*_*man 2 memory cuda bandwidth
auto ts = std::chrono::system_clock::now();
cudaMemcpyAsync((void**)in_dev, in_host, 1000 * size, cudaMemcpyHostToDevice, stream_in);
cudaMemcpyAsync((void**)out_host, out_dev, 1000 * size, cudaMemcpyDeviceToHost, stream_out);
cudaStreamSynchronize(stream_in);
cudaStreamSynchronize(stream_out);
time_data.push_back(std::chrono::system_clock::now() - ts);
这是我为自己的教育目的制定的基准的结果。非常简单，程序的每个“周期”都会启动并行数据传输，并在获取时间戳之前等待这些操作完成。
内核版本添加了一个简单的内核，该内核对每个数据字节（也在不同的流上）进行操作。内核执行时间的趋势对我来说很有意义 - 我的设备只有这么多 SM/核心，一旦我要求更多，它就会开始花费更长的时间。
我不明白的是，为什么仅内存传输测试在与核心限制几乎相同的数据大小点上开始呈指数级增长。我的设备的内存带宽标榜为 600 GB/s。此处传输 10 MB 平均需要约 1.5 毫秒，这并不是给定带宽的餐巾纸数学建议的结果。我的预期是内存传输延迟周围的时间几乎是恒定的，但情况似乎并非如此。
为了确认这不是我的盗版时间戳方法，我使用 NSight Compute 运行了仅内存版本，并确认从 N=1000 KB 到 N=10000 KB 将平均异步传输时间从约 80 us 增加到约 800 us。
我对 D/H 内存传输性能缺少什么？获得良好带宽的关键是重叠大量小型传输而不是大型传输，还是会因为有限的复制引擎瓶颈而变得更糟？
我在配备 PCIe4 系统的 RTX 3070 Ti 上运行了此基准测试。
许多 CUDA 操作可以粗略地建模为“开销”和“持续时间”。持续时间通常可以根据操作特性来预测，例如传输大小除以带宽。“开销”可以粗略地建模为固定量——例如5微秒。
您的图表由多个测量值组成：
与启动传输或“周期”相关的“开销”。CUDA 异步操作的最短持续时间通常为 5-50 微秒。这在蓝色曲线的“平坦”左侧表示。这里的“周期”代表两次传输，加上“内核”版本的内核启动开销。这些“开销”数字的组合代表蓝色和橙色曲线的 y 截距。从蓝色曲线到橙色曲线的距离表示内核操作的添加（您尚未显示）。在曲线的左侧，操作规模非常小，因此与“开销”贡献相比，“持续时间”部分的贡献很小。这解释了左侧曲线的近似平坦度。
操作的“持续时间”。在曲线的右侧，近似线性区域对应于“持续时间”贡献，因为它变大并且使“间接费用”成本相形见绌。蓝色曲线的斜率应与 PCIE 传输带宽相对应。对于 Gen4 系统，每个方向应约为 20-24GB/s（它与 600GB/s 的 GPU 内存带宽没有联系 - 它受到 PCIE 总线的限制。）橙色曲线的斜率也与 PCIE 有关带宽，因为这是整体操作的主要贡献者。
“内核”的贡献。蓝色和橙色曲线之间的距离代表内核操作对 PCIE 数据传输的贡献。
我不明白的是，为什么仅内存传输测试在与核心限制几乎相同的数据大小点上开始呈指数级增长。我的设备的内存带宽标榜为 600 GB/s。此处传输 10 MB 平均需要约 1.5 毫秒，这并不是给定带宽的餐巾纸数学建议的结果。
这里的主导传输由 PCIE 总线控制。该带宽不是 600GB/s，而是每个方向 20-24GB/s。此外，除非您使用固定内存作为传输的主机内存，否则实际带宽将约为可实现的最大带宽的一半。这与您的测量结果非常吻合：10MB/1.5ms = 6.6GB/s。为什么这是有道理的？第一次传输时，您将以约 10GB/s 的速率传输 10MB。除非您使用固定内存，否则该操作将阻塞并且不会与第二次传输同时执行。然后，您在第二次传输时以约 10GB/s 的速率传输 10MB。这是 10GB/s 时的 20MB，因此我们预计传输时间约为 2ms。您的实际传输速度可能接近 12GB/s，这将使预期非常接近 1.5 毫秒。
我的预期是内存传输延迟周围的时间几乎是恒定的，但情况似乎并非如此。
我不确定该声明的确切含义，但对于相当大的传输大小，预计时间不会独立于传输大小而恒定。时间应该是基于传输大小的乘数（带宽）。
我使用 NSight Compute 运行了仅内存版本，并确认从 N=1000 KB 到 N=10000 KB 将平均异步传输时间从约 80 us 增加到约 800 us。
这就是期望。传输更多数据需要更多时间。如果“持续时间”贡献明显大于“开销”贡献，您通常会观察到这种情况，图表右侧就是如此。
下面是一个电子表格，显示了一个具体示例，使用 12GB/s 的 PCIE 带宽和 5 微秒的固定操作开销。“2 次操作总计”列非常接近地跟踪您的蓝色曲线：
