为什么 cudaMemcpyAsync 和内核启动即使使用异步流也会阻塞？答案

【问题标题】：Why do cudaMemcpyAsync and kernel launches block even with an asynchronous stream?为什么 cudaMemcpyAsync 和内核启动即使使用异步流也会阻塞？
【发布时间】：2017-10-26 17:54:56
【问题描述】：

考虑使用以下程序在非阻塞 GPU 流上对一些工作进行排队：

#include <iostream>

using clock_value_t = long long;

__device__ void gpu_sleep(clock_value_t sleep_cycles) {
    clock_value_t start = clock64();
    clock_value_t cycles_elapsed;
    do { cycles_elapsed = clock64() - start; }
    while (cycles_elapsed < sleep_cycles);
}

void callback(cudaStream_t, cudaError_t, void *ptr) { 
    *(reinterpret_cast<bool *>(ptr)) = true; 
}

__global__ void dummy(clock_value_t sleep_cycles) { gpu_sleep(sleep_cycles); }

int main() {
    const clock_value_t duration_in_clocks = 1e6;
    const size_t buffer_size = 1e7;
    bool callback_executed = false;
    cudaStream_t stream;
    auto host_ptr = std::unique_ptr<char[]>(new char[buffer_size]);
    char* device_ptr;
    cudaMalloc(&device_ptr, buffer_size);
    cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
    cudaMemcpyAsync(device_ptr, host_ptr.get(), buffer_size, cudaMemcpyDefault, stream);
    dummy<<<128, 128, 0, stream>>>(duration_in_clocks);
    cudaMemcpyAsync(host_ptr.get(), device_ptr, buffer_size, cudaMemcpyDefault, stream);
    cudaStreamAddCallback(
        stream, callback, &callback_executed, 0 /* fixed and meaningless */);
    snapshot = callback_executed;
    std::cout << "Right after we finished enqueuing work, the stream has "
        << (snapshot ? "" : "not ") << "concluded execution." << std::endl;
    cudaStreamSynchronize(stream);
    snapshot = callback_executed;
    std::cout << "After cudaStreamSynchronize, the stream has "
        << (snapshot ? "" : "not ") << "concluded execution." << std::endl;
}

缓冲区的大小和内核睡眠周期的长度足够高，当它们与 CPU 线程并行执行时，它应该在它们结束之前完成排队（8ms+8ms 用于复制和内核为 20 毫秒）。

然而，看看下面的跟踪，这两个cudaMemcpyAsync() 似乎实际上是同步的，即它们阻塞直到（非阻塞）流实际上结束了复制。这是预期的行为吗？它似乎收缩了CUDA Runtime API documentation 的relevant section。这有什么意义？

跟踪：（编号的行，以微秒为单位的时间）：

      1 "Start"        "Duration"    "Grid X"                             "Grid Y"  "Grid Z"    "Block X"   "Block Y"                       "Block Z"  
    104 14102.830000   59264.347000  "cudaMalloc"
    105 73368.351000   19.886000     "cudaStreamCreateWithFlags"
    106 73388.and 20 ms for the kernel).

然而，看看下面的跟踪，这两个cudaMemcpyAsync() 似乎实际上是同步的，即它们阻塞直到（非阻塞）流实际上结束了复制。这是预期的行为吗？它似乎与 CUDA Runtime API 文档的相关部分相矛盾。这有什么意义？

850000   8330.257000   "cudaMemcpyAsync"
        107 73565.702000   8334.265000   47.683716                            5.587311  "Pageable"  "Device"    "GeForce GTX 650 Ti BOOST (0)"  "1"        
        108 81721.124000   2.394000      "cudaConfigureCall"
        109 81723.865000   3.585000      "cudaSetupArgument"
        110 81729.332000   30.742000     "cudaLaunch (dummy(__int64) [107])"
        111 81760.604000   39589.422000  "cudaMemcpyAsync"
        112 81906.303000   20157.648000  128                                  1         1           128         1                               1          
        113 102073.103000  18736.208000  47.683716                            2.485355  "Device"    "Pageable"  "GeForce GTX 650 Ti BOOST (0)"  "1"        
        114 121351.936000  5.560000      "cudaStreamSynchronize"

【问题讨论】：

如果我理解得很好，我想你可能错过了这个：docs.nvidia.com/cuda/cuda-runtime-api/…
@RobinThoni：所以基本上你是说如果主机端内存被固定，我会得到异步行为？
是的，这将是预期的行为
@RobinThoni：但是主机到设备的传输呢？这似乎符合异步链接的标准，但似乎是同步完成的。
@RobinThoni：你是说我应该理解“from”和“to”具有它们的实际含义，但含义不兼容:-( ...无论如何，谢谢，也许我会感兴趣你在我的followup question。

标签： asynchronous cuda cuda-streams

【解决方案1】：

这看起来很奇怪，所以我联系了 CUDA 驱动团队的人，他确认文档是正确的。我也能够确认：

#include <iostream>
#include <memory>

using clock_value_t = long long;

__device__ void gpu_sleep(clock_value_t sleep_cycles) {
    clock_value_t start = clock64();
    clock_value_t cycles_elapsed;
    do { cycles_elapsed = clock64() - start; }
    while (cycles_elapsed < sleep_cycles);
}

void callback(cudaStream_t, cudaError_t, void *ptr) { 
    *(reinterpret_cast<bool *>(ptr)) = true; 
}

__global__ void dummy(clock_value_t sleep_cycles) { gpu_sleep(sleep_cycles); }

int main(int argc, char* argv[]) {
  cudaFree(0);
  struct timespec start, stop;
    const clock_value_t duration_in_clocks = 1e6;
    const size_t buffer_size = 2 * 1024 * 1024 * (size_t)1024;
    bool callback_executed = false;
    cudaStream_t stream;
    void* host_ptr;
    if (argc == 1){
      host_ptr = malloc(buffer_size);
    }
    else {
      cudaMallocHost(&host_ptr, buffer_size, 0);
    }
    char* device_ptr;
    cudaMalloc(&device_ptr, buffer_size);
    cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);

    clock_gettime(CLOCK_PROCESS_CPUTIME_ID, &start);
    cudaMemcpyAsync(device_ptr, host_ptr, buffer_size, cudaMemcpyDefault, stream);
    clock_gettime(CLOCK_PROCESS_CPUTIME_ID, &stop);
    double result = (stop.tv_sec - start.tv_sec) * 1e6 + (stop.tv_nsec - start.tv_nsec) / 1e3;
    std::cout << "Elapsed: " << result / 1000 / 1000<< std::endl;

    dummy<<<128, 128, 0, stream>>>(duration_in_clocks);

    clock_gettime(CLOCK_PROCESS_CPUTIME_ID, &start);
    cudaMemcpyAsync(host_ptr, device_ptr, buffer_size, cudaMemcpyDefault, stream);
    clock_gettime(CLOCK_PROCESS_CPUTIME_ID, &stop);
    result = (stop.tv_sec - start.tv_sec) * 1e6 + (stop.tv_nsec - start.tv_nsec) / 1e3;
    std::cout << "Elapsed: " << result / 1000 / 1000 << std::endl;

    cudaStreamAddCallback(
        stream, callback, &callback_executed, 0 /* fixed and meaningless */);
    auto snapshot = callback_executed;
    std::cout << "Right after we finished enqueuing work, the stream has "
        << (snapshot ? "" : "not ") << "concluded execution." << std::endl;
    cudaStreamSynchronize(stream);
    snapshot = callback_executed;
    std::cout << "After cudaStreamSynchronize, the stream has "
        << (snapshot ? "" : "not ") << "concluded execution." << std::endl;
}

这基本上是你的代码，有一些修改：

时间测量
从可分页或固定内存分配的开关
缓冲区大小为 2 GiB，以确保可测量的复制时间
cudaFree(0) 强制 CUDA 延迟初始化。

结果如下：

$ nvcc -std=c++11 main.cu -lrt

$ ./a.out # using pageable memory
Elapsed: 0.360828 # (memcpyDtoH pageable -> device, fully async)
Elapsed: 5.20288 # (memcpyHtoD device -> pageable, sync)

$ ./a.out 1 # using pinned memory
Elapsed: 4.412e-06 # (memcpyDtoH pinned -> device, fully async)
Elapsed: 7.127e-06 # (memcpyDtoH device -> pinned, fully async)

从可分页复制到设备时速度较慢，但它确实是异步的。

我很抱歉我的错误。我删除了我以前的 cmets 以避免混淆。

【讨论】：

但是你的时间真的能正确显示异步行为吗？ 0.3 秒不是排队传输所需的时间。也许只是季后赛部分是异步的？
平。另外 - 是否保证 cudaFree(nullptr) 强制 CUDA 的延迟初始化？
是的，但是要复制到可分页内存，CUDA 从设备复制到固定然后再到可分页。这就是为什么会有一些开销。是的，这是有保证的。更准确地说，cudaFree(0) 将初始化 CUDA 并创建一个 CUDA 上下文（相当于调用 cuInit() 和 cuCtxCreate()）。

【解决方案2】：

正如@RobinThoni 所指出的那样，CUDA 内存副本仅在严格的条件下才异步。对于有问题的代码，问题主要是使用未固定（即分页）的主机内存。

引用运行时 API 文档的单独部分（强调我的）：

2。 API 同步行为

API 提供了同步和同步的 memcpy/memset 函数异步形式，后者具有“Async”后缀。 这是一个用词不当，因为每个函数都可能表现出同步或异步行为取决于传递给函数的参数。

...

异步

对于从设备内存到可分页主机内存的传输，该函数仅在复制完成后返回。

这只是它的一半！确实是这样的

对于从可分页主机内存到设备内存的传输，数据将首先暂存于固定主机内存，然后复制到设备；并且该函数只有在暂存发生后才会返回。

【讨论】：