CUDA：为什么大于 64KB 的内存传输会阻塞调用？答案

【问题标题】：CUDA: Why memory transfers greater than 64KB are blocking calls?CUDA：为什么大于 64KB 的内存传输会阻塞调用？
【发布时间】：2022-01-19 15:27:07
【问题描述】：

我是一名学习 CUDA 的学生，我正在尝试了解内存传输的工作原理。我在 Internet 上读到，大于 64KB 的内存传输被视为阻塞调用，而 64KB 以下的内存传输是非阻塞的。我试图用我的课堂笔记来解释它，但我不确定我的推理。我想这是因为传输大量数据可能会导致空闲时间，GPU、CPU 和内存总线都没有工作，所以最好在计算和内存传输之间有重叠。不过我不明白为什么限制真的是 64KB，我什至不确定我刚才说的是否正确。

谁能帮助我或提供更好的解释？非常感谢。

编辑：我首先在以下参考资料中找到了这些信息：https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf

幻灯片编号为 7，标题为“Default Stream (aka Stream '0')”

然后我加深了我的研究，我发现了这个：

在以下参考中：http://gpu.di.unimi.it/slides/lezione7.pdf

但还有很多其他类似的参考资料。

【问题讨论】：

这样的互联网充满了错误信息。请提供参考资料，最好来自 CUDA 文档，以获取此处查询的信息。
我发现的唯一 CUDA 参考如下：developer.download.nvidia.com/CUDA/training/… at slide n.7
我没有看到该幻灯片中提到的任何似乎与问题所声称的事实有关的内容。你能指出与问题相关的幻灯片编号吗？此外，为了提高问题的清晰度，我建议首先指出提示您问题的在线资源。现在在我看来，这个问题是基于一个不正确的前提。
我编辑了我的问题，希望更准确，很抱歉不清楚。
参考似乎是幻灯片 4。措辞有点糟糕（恕我直言）但没有错。它告诉我们：cudaMemcpyAsync() 在主机执行方面是非阻塞的。 cudaMemcpy() 主机设备在主机执行方面通常是阻塞的，除了对于大概（推测！）源于对命令长度的限制。总是喜欢权威来源。此处：CUDA 编程指南第 3.2.6.1 节。

标签： performance memory-management cuda blocking memory-access

【解决方案1】：

CUDA API 中描述了对 cudaMemcpy 或 cudaMemcpyAsync 的调用是否阻塞：https://docs.nvidia.com/cuda/cuda-runtime-api/api-sync-behavior.html#api-sync-behavior__memcpy-async

对于 cudaMemcpy，它表示 cuda 11.6：

对于从可分页主机内存到设备内存的传输，在启动复制之前会执行流同步。一旦可分页缓冲区被复制到暂存内存以便 DMA 传输到设备内存，该函数将返回，但到最终目的地的 DMA 可能尚未完成。

这可以解释为“有点异步”，因为调用可能会在传输完成之前返回，但不会立即返回，具体取决于总传输大小。这完全取决于暂存内存的大小。幻灯片似乎很旧，因为它提到了费米架构。

我尝试使用以下使用nvcc -O3 test1.cu -o test1 编译的代码来找出我的 linux 机器上 CUDA 11.5 和驱动程序 495.46 的暂存缓冲区的当前大小

#include <iostream>

int main(){
    size_t maxsize = 1024ull*1024ull*1024ull;
    void* d_buffer;
    cudaMalloc(&d_buffer, maxsize);
    void* buffer = malloc(maxsize);

    for(size_t bytes = 1; bytes <= maxsize; bytes *= 2){
        std::cerr << bytes << "\n";
        cudaMemcpy(d_buffer, buffer, bytes, cudaMemcpyHostToDevice);
    }
    cudaDeviceSynchronize();
}

使用 nsight-systems 进行分析可以看到，在 API 调用返回之前，最多 1MB 的传输大小不会传输任何数据，并且传输的预期吞吐量为 12GB/s。我的解释是，在 API 返回之前，数据已完全复制到暂存缓冲区，这意味着该缓冲区的大小至少为 1MB，而不是 64KB。

对于大于 1MB 的传输，比如 2MB，使用不同的方法，因为到 GPU 的数据传输在 API 调用返回之前开始。对于 2MB，API 端和传输端之间的时间差约为 86µs。显示的吞吐量为 6.5 GB/s，这对应于 ~559 KB（实际上可能是 512 KB）。因此，对于大于 1MB 的大小，暂存缓冲区似乎被分成两半以允许流水线操作。虽然当前正在将 512 KB 传输到 GPU，但可以将下一个数据块暂存到剩余的 512 KB。

【讨论】：