11.2 异步内存复制

除非输入和输出数据可以驻留在GPU上，让数据流过GPU（即把输入和输出数据传入和传出设备内存）的策略变得尤为重要。最适合用来提高传输性能的2个工具是锁页内存和异步内存复制（它仅可以运行在锁页内存）。

stream2Async.cu应用程序演示了把stream1Device.cu的分页内存替换成锁页内存，并启用异步内存复制的效果。

Measuring times with 128M floats (use --N to specify number of Mfloats)  
Memcpy(host->device): 181.03 ms (5931.33 MB/s)  
Kernel processing: 13.87 ms (116152.99 MB/s)  
Memcpy (device->host): 90.07 ms (5960.35 MB/s)  
Total time (wall clock): 288.68 ms (5579.29 MB/s)

代码清单11-5对比了stream1Device.cu（它执行同步传输）和stream2Async.cu（它执行异步传输）的计时部分的差异。[1]在这两种情况下，4个CUDA事件用于分别记录开始时刻、主机到设备传输完成的时刻、内核执行完成的时刻以及结束时刻。对于stream2Async.cu，所有这些操作在短时间内相继请求了GPU，而GPU在执行它们的时候记录下事件的时间。对于stream1Device.cu，GPU的基于事件的时间是有点不够准确，因为任何对CUDAMemcpy()的调用必须等待GPU完成之后再继续下一操作，会在以evHtoD和evDtoH参数执行CUDAEventRecord()之前导致流水线出现空隙。

值得注意的是，尽管使用了较慢的、简单实现的saxpyGPU函数（来自代码清单11-2），该应用程序消耗的系统时钟时间显示，它的计算速度仍然快了将近一倍：289毫秒对比之前的570.22毫秒。结合使用更快的数据传输和异步执行会获得更好的性能。

尽管性能提高了，应用程序的输出指示了另一个提升性能的机会：一些内核的处理可以与数据传输并发进行。接下来的两节将介绍2种实现“内核执行和传输”重叠进行的方法。

代码清单11-5 同步内核（stream1Device.cu）与异步内核

stream2Async.cu) 对比

//   
// from stream1Device.cu   
//   
cudaEventRecord(evStart, 0);   
cudaMemcpy(dptrX, hptrX, ...,udaMemcpyHostToDevice);   
cudaMemcpy(dptrY, hptrY, ...,udaMemcpyHostToDevice);   
cudaEventRecord(evHtoD, 0);   
saxpyGPU<<nBlocks, nThreads>>>(dptrOut, dptrX, dptrY, N, alpha;   
cudaEventRecord(evKernel, 0);   
cudaMemcy(hptrOut, dptrOut, N*sizeOf(float),udaMemcpyDeviceToHost);   
cudaEventRecord(evDtoH, 0);   
cudaDeviceSynchronize();   
//   
// from stream2Async.cu   
//   
cudaEventRecord(evStart, 0);   
cudaMemcpyAsync(dptrX, hptrX, ...,udaMemcpyHostToDevice, NULL);   
cudaMemcpyAsync(dptrY, hptrY, ...,udaMemcpyHostToDevice, NULL);   
cudaEventRecord(evHtoD, 0);   
saxpyGPU<<nBlocks, nThreads>>>(dptrOut, dptrX, dptrY, N, alpha;   
cudaEventRecord(evKernel, 0);   
cudaMemcpyAsync(hptrOut, dptrOut, N*sizeOf(float), ..., NULL);   
cudaEventRecord(evDtoH, 0);   
cudaDeviceSynchronize();

[1] 为清晰起见，删除了进行错误检查的代码。