11.4 映射锁页内存

对于传送受限型流式负载，例如SAXPY，重新组织程序把映射锁页内存用于输入和输出，有许多好处。

·从节选的stream4Mapped.cu（代码清单11-7）可以看出，它不需要调用复制函数CUDAMemcpy()。

• 它不需要分配设备内存。
  ·对于独立GPU，映射锁页内存执行总线传输，但最大限度地减少了上一节中提到的“突出部分”的数量。无须等待主机到设备的内存复制操作完成，而是在输入数据到达的当时被SM处理。无须等待内核完成后再发起设备到主机的传输，而是在SM处理结束就把数据发布到总线。
  ·对于集成GPU，主机和设备内存共存于同一个内存池，因此映射锁页内存支持“零复制”特性，并消除了任何总线数据传输需要。

代码清单11-7 stream4Mapped函数节选

chTimerGetTime( &chStart );  
cudaEventRecord( evStart, 0 );  
saxpyGPU<<nBlocks, nThreads>>>(dptrOut, dptrX, dptrY, N, alpha);  
cudaEventRecord( evStop, 0 );  
cudaDeviceSynchronize();

映射锁页内存在需要写入主机内存时效果尤为突出（例如，把归约结果传给主机），因为这时与读取操作不同，没有必要等到写操作结束再继续执行。[1]读取映射锁页内存的负载更容易出问题。如果在读取映射锁页内存时，GPU无法维持全速的总线性能，较小的传输性能可能压倒较小“突出部分”带来的益处。同样，对于某些负载，让SM忙于更多的事情比保持（等待）PCIe总线传输更有意义。

就我们的应用程序而言，在我们的测试系统上，使用映射锁页内存的性能优势明显。

Measuring times with 128M floats (use --N to specify number of Mfloats) Total time: 204.54 ms (7874.45 MB/s)

它在204.54毫秒内完成计算任务，显著快于次佳的273毫秒。7.9GB/s的高效带宽表明GPU充分利用了PCIe两个方向的传输潜能。

并不是所有组合使用系统和GPU的方式，都可以借助映射锁页内存维持如此高水平的性能。如果有任何异常的端倪，请把数据保持在设备内存并使用异步内存复制策略，类似于stream2Async.cu。

[1] 硬件设计人员称之为“延迟隐藏”。