为啥 CUDA 内存复制速度表现得像这样，一些恒定的驱动程序开销？

Posted 2023-03-23

【中文标题】为啥 CUDA 内存复制速度表现得像这样，一些恒定的驱动程序开销？

【英文标题】：Why CUDA memory copy speed behaves like this, some constant driver overhead?为什么 CUDA 内存复制速度表现得像这样，一些恒定的驱动程序开销？

【发布时间】：2012-10-19 07:39:00

【问题描述】：

在我的旧 GeForce 8800GT 上使用 CUDA 中的内存时，我总是有一个奇怪的 0.04 毫秒开销。我需要将 ~1-2K 转移到我设备的常量内存中，处理上面的数据并从设备中只获取一个浮点值。

我有一个使用 GPU 计算的典型代码：

//allocate all the needed memory: pinned, device global
for(int i = 0; i < 1000; i++)

    //Do some heavy cpu logic (~0.005 ms long)        
    cudaMemcpyToSymbolAsync(const_dev_mem, pinned_host_mem, mem_size, 0, cudaMemcpyHostToDevice);
    my_kernel<<<128, 128>>>(output);
    //several other calls of different kernels
    cudaMemcpy((void*)&host_output, output, sizeof(FLOAT_T), cudaMemcpyDeviceToHost);
    // Do some logic with returned value 

我决定用这段代码测量 GPU 内存的工作速度（注释了所有内核调用，添加了cudaDeviceSynchronize 调用）：

//allocate all the needed memory: pinned, device global
for(int i = 0; i < 1000; i++)

    //Do some heavy cpu logic (~0.001 ms long)        
    cudaMemcpyToSymbolAsync(const_dev_mem, pinned_host_mem, mem_size, 0, cudaMemcpyHostToDevice);
    cudaMemcpyAsync((void*)&host_output, output, sizeof(FLOAT_T), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    // Do some logic with returned value 

我测量了循环的执行时间，得到了大约 0.05 秒（因此，每次迭代需要 0.05 毫秒）。奇怪的是，当我尝试做更多的内存工作（添加额外的 cudaMemcpyToSymbolAsync 和 cudaMemcpyAsync 调用）时，每次调用我得到额外的 http://www.cs.virginia.edu/~mwb7w/cuda_support/memory_transfer_overhead.html

他还获得了每次将 1K 块传输到 GPU 的 0.01 毫秒。 那么这 0.04 毫秒（0.05 - 0.01）的开销是从哪里来的呢？有任何想法吗？我应该在新卡上试试这个代码吗？

在我看来，在 cudaDeviceSynchronize 和 CPU 代码之后，我的 GeForce 进入了某种省电模式或类似的模式。

【问题讨论】：

0.05 ms 是每次迭代的平均值。如果我是你，我会一直单独查看它是否是一个常数值。第一次用来有开销。

您如何评估经过的时间？是否在进入循环之前调用了 CUDA 函数，以消除设备的初始化成本，如调用 cudaFree(0)。

@pQB，是的，这是一个平均值，但我认为它显示了我从外部 CPU 代码执行代码时所获得时间的真实情况。例如，当我将迭代次数从 1000 更改为 2000 时，我的时间是以前的两倍（如预期的那样）。当我将迭代次数设置为 500 时也会发生同样的情况 - 时间减少了 2 倍。

@phoad，我只是在循环前后调用 GetSystemTime。关于初始化设备的一些函数调用 - 你能不能给我一些链接，我可以阅读它？我对这里发生的事情的主要版本之一是设备设置为某种“睡眠模式”或类似的东西。这可能发生在执行繁重的 CPU 代码之后。我将运行一些测试，尝试消除 CPU 执行。

***.com/questions/11704681/… 只需检查此链接。它包含有关使用更好的 CUDA 提供的计时器以及如何从计时中消除初始化成本的信息。

【参考方案1】：

我建议您增加正在实现的线程数

    //Use malloc() to allocate memory on CPU. 
    //Change mem_size to the total memory to be tranferred to GPU.        
    cudaMemcpyToSymbolAsync(const_dev_mem, pinned_host_mem, mem_size, 0, cudaMemcpyHostToDevice);
    dim3 dimBlock(128,2);
    dim3 dimGrid(64000,1);
    my_kernel<<<dimGrid, dimBlock>>>(output);
    //several other calls of different kernels
    //change size field to 1000*sizeof(FLOAT_T)
    cudaMemcpy((void*)&host_output, output, sizeof(FLOAT_T), cudaMemcpyDeviceToHost);
    // Do some logic with returned value 

如果代码崩溃（因为更多线程或更多 GPU 内存），请使用循环。但是，减少它们。

【讨论】：

很遗憾，我无法按照您描述的方式进行操作。在这个周期（从 0 到 1000）中，我尝试模拟我的真实项目行为：它从互联网接收数据，每秒更新 1000-10000 次（某些股票价格）。在每次更新时，我都应该执行一些繁重的逻辑（大约 10000-100000 次迭代），CUDA 卡比任何 CPU 都能更好地应对这些逻辑。我现在的问题是尽量减少与 GPU 交互时可能获得的任何开销。

看来您的传输是如此之小，以至于它们完全被硬件和驱动程序堆栈的基本延迟所支配。在我的 64 位 Linux 工作站（Xeon 5272 + C2050；PCIe gen2）上，对于 1 字节和 1 KB 之间的任何大小的传输，设备/主机和主机/设备传输的延迟约为 25 us，这与您的 50 微秒往返行程相匹配看。异步副本的想法是创建多个副本和内核流，以便副本可以在管道中与内核重叠。人们通常可以获得（接近）完美的重叠。这可以优化吞吐量，但不能改善往返延迟。

@njuffa，非常感谢你给我你的真实数字。这让我脚踏实地，让我觉得真的没有办法让它更快地工作。我只是想实现this results - 10us 的传输。正式地我得到了它们：当添加任何其他异步传输时，测量的时间增加了约 10 微秒。但是，如果我只留下一次转移 - 时间不是 10 我们，而是 50 我们。这 40 我们 (50-10) 对我来说非常重要。在此期间可以在我的内核中进行很多计算。

我不知道是什么底层机制决定了最短周转时间，但根据经验观察，在我的测试中，主板和 GPU 的各种组合具有从 14 微秒到 25 微秒的一次性 1KByte 复制延迟框架。这似乎表明延迟可能主要是由硬件引起的，您可以通过尝试不同的硬件组合来减少它。

按照我的建议尝试不同的硬件组合，我能找到的最低主机/设备传输时间是在 64 位 Linux 系统上，配备 Xeon X5550 CPU 和 M2090，运行 CUDA 5.0。我使用主机上的固定内存在任一方向（主机->设备、设备->主机）传输 1KB 的时间为 10 微秒。