CUDA C++ 编程指引：异步并发执行 | CUDA

本系列参考自 CUDA C++ Programming Guide。

并发行为

CUDA 允许以下操作彼此之间并发执行：

• 主机端计算
• 设备端计算（内核函数执行）
• 主机端到设备端传输数据
• 设备端到主机端传输数据
• 设备端内部传输数据
• 设备间传数据

主机端/设备端并发执行

为了易于使用主机和设备间的异步执行，一些函数是异步的：在设备完全完成任务前，控制权已经返回给主机线程了。包括：

• 内核函数执行（可以通过将 CUDA_LAUNCH_BLOCKING 设为 1，来禁用内核并发执行，调试时使用）；
• 设备内两个不同地址间的内存拷贝函数；
• 主机和设备内拷贝小于 64KB 的内存块；
• 内存拷贝函数中带有 Async 后缀的；
• 设置设备内存的函数调用（cudaMemset()）。

内核函数并发执行

计算能力 ≥2.0 及以上的设备，支持多个内核函数并发执行。可以通过检查 concurrentKernels 来确定，如果返回 1，说明支持。

多个并发的内核函数必须位于同一个 CUDA 上下文，不同 CUDA 上下文上的内核函数不能并行。

数据传输和内核并发执行（需要使用锁页内存）

一些设备支持数据传输（主机端到设备端、设备端到主机端）和内核执行并行，可通过检查 asyncEngineCount 来确认，大于 0 即可使用，前提是主机端使用锁页内存。一些设备支持设备端内部数据传输和内核执行/数据传输并行，可通过检查 concurrentKernels 来确认。

并发数据传输（需要使用锁页内存）

在计算能力 ≥2.0 的设备上，从主机锁页内存复制数据到设备内存和从设备内存复制数据到主机锁页内存，这两个操作可并发执行。可以通过检查 asyncEngineCount 属性查询这种能力，如果等于 2，说明支持。

流

创建和销毁

cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
    cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);

for (int i = 0; i < 2; ++i) {
    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel <<<100, 512, 0, stream[i]>>>(outputDevPtr + i * size, inputDevPtr + i * size, size);
    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
}

for (int i = 0; i < 2; ++i)
    cudaStreamDestroy(stream[i]);

上述代码创建了两个流且在锁页内存中分配了一个名为 hostPtr 的浮点数组。每个流都是一个由一次主机到设备的传输、一次内核发射、一次设备到主机的传输组成的系列。每个流将它的 hostPtr 输入数组的部分拷贝到设备内存数组 inputdevPtr，调用 MyKernel() 内核处理 inputDevPtr，然后将结果 outputDevPtr 传输回 hostPtr 同样的部分。

当设备还在执行流中的任务，而用户调用 cudaStreamDestroy() 函数时，函数会立刻执行（不会阻塞）。之后，当流中的任务完成后，与流相关的资源会自动释放。

另外需要注意的是，上例中主机端线程、数据拷贝和内核执行完全异步，因此在“拷贝回主机端”这一操作完成之前，主机端的内存数据是不正确的。必须在数据返回的一步做同步操作，方能保证数据是正确的。

默认流

在调用内核函数时，不指定流或者将流指定为 0，则代表使用了默认流。关于默认流的具体内容，可以查看本系列另一篇文章。

显式同步

有很多方法显式的在流之间同步：

• cudaDeviceSynchronize() 直到所有线程向设备端的所有流的所有已送入指令完成。
• cudaStreamSynchronize() 直到指定流的之前所有已送入指令完成。
• cudaStreamWaitEvent() 以一个流和一个事件为参数，使得在调用 cudaStreamWaitEvent() 后加入到指定流的所有命令暂缓执行直到事件完成。流可以是 0，此时在调用 cudaStreamWaitEvent() 后加入到所有流的所有命令等待事件完成。
• cudaStreamQuery() 用于查询流中的所有之前的命令是否已经完成。

为了避免不必要的性能损失，这些函数最好用于计时或隔离失败的发射或内存拷贝。

隐式同步

一般来讲，不同流内的命令可以并行。但是当任何一个流执行如下的命令时，情况例外，不能并发：

• 锁页内存的分配
• 设备端内存分配
• 设备端内存设置（cudaMemset）
• 设备内部拷贝
• NULL stream 内的命令
• L1 /共享内存配置之间的切换

重叠行为

操作的重叠程度，一方面取决于各个操作的顺序，另一方面取决于设备支持重叠的程度：

• 是否支持内核执行并发
• 数据传输与内核执行并发
• 数据传输并发

主机函数（回调）

运行时库通过 cudaStreamAddCallback() 提供了一种在任何执行点向流插入 CPU 回调函数的方式。一旦在插入点之前发射到流的所有命令执行完成，回调函数就会在主机上执行。如果是 NULL 流，则只能在插入点之前其它流的所有命令都完成后才能执行。

在下面的代码中，两个流在将数据复制到主机端时，会调用回调函数 MyCallbak。

void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void *data){
    printf("Inside callback %d\n", (size_t)data);
}
...
for (size_t i = 0; i < 2; ++i) {
    cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel<<<100, 512, 0, stream[i]>>>(devPtrOut[i], devPtrIn[i], size);
    cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost, stream[i]);
    cudaStreamAddCallback(stream[i], MyCallback, (void*)i, 0);
}

需要注意的是，回调函数中不能直接或间接的执行 CUDA 函数，否则会因为等待自己完成而造成死锁。

流的优先级

可以通过 cudaStreamCreateWithPriority() 来在创建流时指定流的优先级。设备允许的优先级范围可由 cudaDeviceGetStreamPriorityRange() 来获得。

运行时，高优先级流中的线程块不能打断正在执行的低优先级中的线程块。但是当低优先级流的线程块退出 SM 时，高优先级流中的线程块会被优先调度进 SM。

下面的代码示例获取当前设备的允许的优先级范围，并创建具有最高和最低可用优先级的流。

// get the range of stream priorities for this device
int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
// create streams with highest and lowest available priorities
cudaStream_t st_high, st_low;
cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low);

事件

事件可以被压入流中以监视流的运行情况，或者用于精确计时。当事件记载点前面，事件指定的流中的所有任务全部完成时，事件被触发。如果向 NULL 流压入事件，则当压入事件前的向所有流压入的任务完成后，事件才被触发。

创建与销毁

下面的代码创建并销毁了两个事件：

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
...
cudaEventDestroy(start);
cudaEventDestroy(stop);

计算时间

下例是一个使用事件计算时间的例子：

cudaEventRecord(start, 0);  // 记录事件（将事件压入流），流 0 则代表所有流完成任务后事件才会被触发
for (int i = 0; i < 2; ++i) {
    cudaMemcpyAsync(inputDev + i * size, inputHost + i * size, size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel<<<100, 512, 0, stream[i]>>>(outputDev + i * size, inputDev + i * size, size);
    cudaMemcpyAsync(outputHost + i * size, outputDev + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);    // 获取两个事件发生的时间差（ms）

同步调用

调用同步函数时，在设备完成请求的任务之前，控制权不会返回给主机线程。在主机线程执行任何其他 CUDA 调用之前，可以通过调用带有某些特定标志的 cudaSetDeviceFlag() 来指定主机线程的让步、阻塞或自旋状态。

参考

CUDA C++ Programming Guide
《CUDA C Programming Guide》(《CUDA C 编程指南》)导读
《CUDA 并行程序设计：GPU 编程指南》