了解内存传输性能 (CUDA)
Understanding memory transfer performance (CUDA)
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);
这是我为自己的教育目的制定的基准测试结果。非常简单，程序的每个 'cycle' 都会启动并行的 t运行 数据传输，并在获取时间戳之前等待这些操作完成。
内核版本添加了一个简单的内核，它对数据的每个字节（也在不同的流上）进行操作。内核执行时间的趋势对我来说很有意义——我的设备只有这么多 SMs/cores，一旦我要求更多，它就会开始花费更长的时间。
我不明白的是为什么内存 t运行sfer only 测试在与核心限制几乎相同的数据大小点开始呈指数增长。我设备的内存带宽标示为 600 GB/s。 T运行在这里传输 10 MB 平均需要大约 1.5 毫秒，这不是餐巾纸数学建议的给定带宽。我的预期是内存 t运行sfer 延迟时间几乎恒定，但情况似乎并非如此。
为了确认这不是我的盗版时间戳方法，我 运行 使用 NSight Compute 的内存版本并确认从 N=1000 KB 到 N=10000 KB 增加了平均异步 t运行将时间从 ~80 us 缩短到 ~800 us。
关于 D/H 内存性能 运行 我错过了什么？获得良好带宽的关键是重叠大量小的 t运行sfer 而不是大的 t运行sfer 还是会因为有限的复制引擎瓶颈而变得更糟？
我 运行 此基准测试是在配备 pcie4 系统的 RTX 3070 Ti 上进行的。
许多 CUDA 操作可以粗略地建模为“开销”和“持续时间”。持续时间通常可以从操作特性中预测 - 例如传输大小除以带宽。 “开销”可以粗略地建模为固定数量 - 例如5 微秒。
您的图表包含多个测量值：
与启动传输或“周期”相关的“开销”。 CUDA 异步操作通常具有大约 5-50 微秒的最短持续时间。这显示在蓝色曲线的“平坦”左侧。这里的“周期”代表两次传输，加上，在“内核”版本的情况下，内核启动开销。这些“开销”数字的组合代表蓝色和橙色曲线的 y-intercept。从蓝色曲线到橙色曲线的距离表示添加了内核操作（您尚未显示）。在曲线的左侧，操作规模非常小，以至于与“开销”贡献相比，“持续时间”部分的贡献很小。这解释了左侧曲线的近似平坦度。
操作的“持续时间”。在曲线的右侧，近似线性区域对应于“持续时间”贡献，因为它变大并使“开销”成本相形见绌。蓝色曲线的斜率应对应于 PCIE 传输带宽。对于 Gen4 系统，每个方向应该约为 20-24GB/s（它与 600GB/s 的 GPU 显存带宽无关——它受 PCIE 总线限制。）橙色曲线的斜率也与 PCIE 有关带宽，因为这是整个操作的主要贡献者。
“内核”的贡献。蓝色和橙色曲线之间的距离代表内核操作的贡献，over/above 只是 PCIE 数据传输。
What I don't understand is why the memory transfer only tests start ramping up exponentially at nearly the same data size point as the core limitations. The memory bandwidth for my device is advertised as 600 GB/s. Transferring 10 MB here takes on average ~1.5 milliseconds which isn't what napkin math would suggest given bandwidth.
这里的显性传输是由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.5ms。
My expectation was that time would be nearly constant around the memory transfer latency, but that doesn't seem to be the case.
我不确定该语句的确切含义，但是对于相当大的传输大小，预计时间不会独立于传输大小而保持不变。时间应该是基于传输大小的乘数（带宽）。
I ran the memory only version with NSight Compute and confirmed that going from N=1000 KB to N=10000 KB increased average async transfer time from ~80 us to around ~800 us.
这是期望。传输更多数据需要更多时间。如果“持续时间”贡献明显大于“开销”贡献，这通常是您会观察到的情况，这在图表的右侧是正确的。
这是一个显示具体示例的电子表格，使用 12GB/s 的 PCIE 带宽和 5 微秒的固定操作开销。 “2 次操作的总数”列非常接近地跟踪您的蓝色曲线：
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);
这是我为自己的教育目的制定的基准测试结果。非常简单，程序的每个 'cycle' 都会启动并行的 t运行 数据传输，并在获取时间戳之前等待这些操作完成。
内核版本添加了一个简单的内核，它对数据的每个字节（也在不同的流上）进行操作。内核执行时间的趋势对我来说很有意义——我的设备只有这么多 SMs/cores，一旦我要求更多，它就会开始花费更长的时间。
我不明白的是为什么内存 t运行sfer only 测试在与核心限制几乎相同的数据大小点开始呈指数增长。我设备的内存带宽标示为 600 GB/s。 T运行在这里传输 10 MB 平均需要大约 1.5 毫秒，这不是餐巾纸数学建议的给定带宽。我的预期是内存 t运行sfer 延迟时间几乎恒定，但情况似乎并非如此。
为了确认这不是我的盗版时间戳方法，我 运行 使用 NSight Compute 的内存版本并确认从 N=1000 KB 到 N=10000 KB 增加了平均异步 t运行将时间从 ~80 us 缩短到 ~800 us。
关于 D/H 内存性能 运行 我错过了什么？获得良好带宽的关键是重叠大量小的 t运行sfer 而不是大的 t运行sfer 还是会因为有限的复制引擎瓶颈而变得更糟？
我 运行 此基准测试是在配备 pcie4 系统的 RTX 3070 Ti 上进行的。
许多 CUDA 操作可以粗略地建模为“开销”和“持续时间”。持续时间通常可以从操作特性中预测 - 例如传输大小除以带宽。 “开销”可以粗略地建模为固定数量 - 例如5 微秒。
您的图表包含多个测量值：
与启动传输或“周期”相关的“开销”。 CUDA 异步操作通常具有大约 5-50 微秒的最短持续时间。这显示在蓝色曲线的“平坦”左侧。这里的“周期”代表两次传输，加上，在“内核”版本的情况下，内核启动开销。这些“开销”数字的组合代表蓝色和橙色曲线的 y-intercept。从蓝色曲线到橙色曲线的距离表示添加了内核操作（您尚未显示）。在曲线的左侧，操作规模非常小，以至于与“开销”贡献相比，“持续时间”部分的贡献很小。这解释了左侧曲线的近似平坦度。
操作的“持续时间”。在曲线的右侧，近似线性区域对应于“持续时间”贡献，因为它变大并使“开销”成本相形见绌。蓝色曲线的斜率应对应于 PCIE 传输带宽。对于 Gen4 系统，每个方向应该约为 20-24GB/s（它与 600GB/s 的 GPU 显存带宽无关——它受 PCIE 总线限制。）橙色曲线的斜率也与 PCIE 有关带宽，因为这是整个操作的主要贡献者。
“内核”的贡献。蓝色和橙色曲线之间的距离代表内核操作的贡献，over/above 只是 PCIE 数据传输。
What I don't understand is why the memory transfer only tests start ramping up exponentially at nearly the same data size point as the core limitations. The memory bandwidth for my device is advertised as 600 GB/s. Transferring 10 MB here takes on average ~1.5 milliseconds which isn't what napkin math would suggest given bandwidth.
这里的显性传输是由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.5ms。
My expectation was that time would be nearly constant around the memory transfer latency, but that doesn't seem to be the case.
我不确定该语句的确切含义，但是对于相当大的传输大小，预计时间不会独立于传输大小而保持不变。时间应该是基于传输大小的乘数（带宽）。
I ran the memory only version with NSight Compute and confirmed that going from N=1000 KB to N=10000 KB increased average async transfer time from ~80 us to around ~800 us.
这是期望。传输更多数据需要更多时间。如果“持续时间”贡献明显大于“开销”贡献，这通常是您会观察到的情况，这在图表的右侧是正确的。
这是一个显示具体示例的电子表格，使用 12GB/s 的 PCIE 带宽和 5 微秒的固定操作开销。 “2 次操作的总数”列非常接近地跟踪您的蓝色曲线：