CUDA 异步执行

异步执行 是指 GPU进行的操作（内核启动 或 异步存储器拷贝函数）从主机端启动后，在GPU真正完成这些操作前，在主机端就可以得到这些函数的返回值，CPU线程就可以继续进行下一步操作。也就是说，CPU端的API函数和内核启动的结束 和 GPU端真正完成这些操作是异步的。通过异步函数，CPU可以在GPU端进行 运算 或 数据传输 的同时进行其它操作，更加有效地利用系统中的计算资源。

内核启动 和 显存内的数据拷贝（Device to Device）总是异步的，而 内存和显存间的数据拷贝函数 则有异步和同步两个版本。

下面的代码中进行了一次同步的数据拷贝，只有确实完成CPU和GPU间数据传输后，cudaMemcpy() 函数才会返回，此时CPU才能开始执行后面的 cpuFunction() 函数。

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
cpuFunction();

调用异步的 cudaMemcpyAsync() 函数时，CPU线程通过存储器管理API函数启动了一次数据传输，在GPU传输的同时，cudaMemcpyAsync() 函数就已经返回了。此时，cpuFunction() 函数实质上是与GPU和CPU间的数据传输同时进行的。

cudaMemcpyAsync(a_d, a_h, size, cudaMemcpyHostToDevice);
cpuFunction();

对主机端来说，内核启动是异步的。下面的代码中，cpuFunction() 也会在 kernel函数在GPU上运行结束之前就开始执行：

kernel <<<blocks, threads>>> (a_d);
cpuFunction();

如果主机端需要使用内核函数计算得到的结果，应该加入同步保证操作的顺序一致性：

kernel <<<blocks, threads>>> (a_d);
cudaThreadSynchronsize();  // 重点
cpuFunction();

属于同一个流中的内核启动总是同步的。下面的代码中，没有显式地为 kernel1 和 kernel2 指定所属的流，此时它们都属于默认的stream0。因此，kernel2 总是会在 kernel1 执行完成之后才能开始执行：

kernel1 <<<blocks1, threads1>>> (a_d1);
kernel2 <<<blocks2, threads2>>> (a_d2);

如果几次内核启动分别属于不同的流，它们的执行可能是乱序的，如果它们都要向同一块存储器写数据，就有可能因为竞写而导致错误。因此在使用 stream 时，stream 之间的数据相关性要尽量小。下面的代码中存在两个流，当 stream0 的异步数据传输在GPU上执行完毕时，stream1 的异步数据传输和 stream0 的 kernel 执行可以在GPU上同时执行：

cudaStreamCreate(&stream0);
cudaStreamCreate(&stream1);
cudaMemcpyAsync(a_d, a_h, size, cudaMemcpyHostToDevice, 0);
cudaMemcpyAsync(a_d, a_h, size, cudaMemcpyHostToDevice, 1);
kernel <<<blocks1, threads1, stream0>>> (a_d1);
kernel <<<blocks1, threads1, stream1>>> (a_d1);

下面说明如何使用 流和异步 提高程序性能

方法一
使用 流和异步 使 CPU 和 GPU 同时进行运算：

cudaStreamCreate(&stream1);
cudaMemcpyAsync(a_d, a_h, size, cudaMemcpyHostToDevice, 1);
kernel <<<blocks1, threads1, stream1>>> (a_d1);
cpuFunctionA(CPUresult);
cudaThreadSynchronsize();  // 同步放在这，因为 cpuFunctionA() 和 kernel 无数据关联，可以同时进行运算，而cpuFunctionB() 需要cpuFunctionA() 和 kernel 的结果
cudaMemcpyAsync(a_h, a_d, size, cudaMemcpyHostToDevice, 1);
cpuFunctionB(CPUresult, a_d);

假设 cpuFunctionA() 与 GPU 的操作没有什么冲突，而 cpuFunctionB() 则需要用到 cpuFunctionA() 和 GPU 运算的结果，那么可以利用内核启动的异步执行让 CPU 和GPU 同时运算，再使用同步保证GPU已经完成操作。这时，cpuFunctionB() 就可以安全地消费 CPU 和 GPU 同时生产的数据。

方法二
利用不同流之间的异步执行，使流之间的传输和运算能够同时执行，更好地利用 GPU 资源：

for(int i = 0;i < nStreams; i++)
{
	cudaStreamCreate(&(stream[i]));
}
size = N * sizeof(float) / nStreams;
for(int i = 0;i < nStreams; i++)
{
	offset = i * N / nStreams;
	cudaMemcpyAsync(a_d + offset, a_h + offset, size, cudaMemcpyHostToDevice, stream[i]);
}
for(int i = 0;i < nStreams; i++)
{
	offset = i * N / nStreams;
	kernel <<<N / (threads *nStreams), threads, 0, stream[i]>>> (a_d + offset);  // blocks 数量重点
}

参考书目：GPU高性能运算之CUDA