zero copy（零复制）是一种特殊形式的内存映射，它允许你将host内存直接映射到设备内存空间上。其实就是设备可以通过直接内存访问（direct memory access，DMA）方式来访问主机的锁页内存。

锁页主机内存

现代操作系统都支持虚拟内存，操作系统实现虚拟内存的主要方法就是通过分页机制。操作系统将内存中暂时不使用的内容换出到外存（硬盘等大容量存储）上，从而腾出空间存放将要调入内存的信息。这样，系统好像为用户提供了一个比实际内存大得多的存储器，称为虚拟存储器。
锁页就是将内存页面标记为不可被操作系统换出的内存。所以设备驱动程序给这些外设编程时，可以使用页面的物理地址直接访问内存（DMA），从而避免从外存到内存的复制操作。CPU 仍然可以访问上述锁页内存，但是此内存是不能移动或换页到磁盘上的。CUDA 中把锁页内存称为pinned host memory 或者page-locked host memory。

锁页主机内存的优势

使用锁页内存（page-locked host memory）有一些优势：

• 锁页内存和GPU内存之间的拷贝可以和内核程序同时执行，也就是异步并发执行。
• 在一些设备上锁页内存的地址可以从主机地址空间映射到CUDA 地址空间，免去了拷贝开销。
• 在拥有前线总端的系统上，如果主机内存被分配为锁页内存，主机内存和GPU 内存带宽可以达到更高，如果主机内存被分配为Write-Combining Memory，带宽会进一步提升。

然而锁页主机存储器是稀缺资源，所以锁页内存分配得多的话，分配会失败。另外由于减少了系统可分页的物理存储器数量，分配太多的分页锁定内存会降低系统的整体性能

使用锁页主机内存

在GPU 上分配的内存默认都是锁页内存，这只是因为GPU 不支持将内存交换到磁盘上。在主机上分配的内存默认都是可分页，如果需要分配锁页内存，则需要使用cudaMallocHost() 或者cudaHostAlloc()。释放时需要使用cudaFreeHost() 释放这一块内存。调用常规的C函数释放，可能会崩溃或者出现一些不常见的错误。也可以通过函数cudaHostRegister() 把可分页内存标记为锁页内存。

__host__ ​cudaError_t cudaMallocHost ( void** ptr, size_t size )
__host__ ​cudaError_t cudaHostAlloc ( void** pHost, size_t size, unsigned int  flags )
__host__ ​cudaError_t cudaFreeHost ( void* ptr )

cudaHostAlloc() 多了一个可选形参flags ，功能更强大。flags 的值可以取如下值。

#define cudaHostAllocDefault 0x00
Default page-locked allocation flag
#define cudaHostAllocMapped 0x02
Map allocation into device space
#define cudaHostAllocPortable 0x01
Pinned memory accessible by all CUDA contexts
#define cudaHostAllocWriteCombined 0x04
Write-combined memory

cudaHostRegister() 函数用于把已经的存在的可分页内存注册为分页锁定的。

__host__ ​cudaError_t cudaHostRegister ( void* ptr, size_t size, unsigned int  flags )

flags 是一个可选形参，可以取如下值。

#define cudaHostRegisterDefault 0x00
Default host memory registration flag
#define cudaHostRegisterIoMemory 0x04
Memory-mapped I/O space
#define cudaHostRegisterMapped 0x02
Map registered memory into device space
#define cudaHostRegisterPortable 0x01
Pinned memory accessible by all CUDA contexts

下面分别介绍这些flags 的作用。

Portable Memory

一块锁页内存可被系统中的所有设备使用（一个系统中有多个CUDA设备时）。 启用这个特性需要在调用cudaHostAlloc() 时使用cudaHostAllocPortable 选项，或者在调用cudaHostRegister() 使用cudaHostRegisterPortable 选项。

Write-Combining Memory

默认情况下，锁页主机存储是可缓存的。可以在调用cudaHostAlloc() 时传入cudaHostAllocWriteCombined 标签使其被分配为写结合的（Write-Combining Memory）。写结合存储不使用L1 和L2 cache，所以程序的其它部分就有更多的缓存可用。此外，写结合内存通过PCI-E 传输数据时不会被监视（snoop），这能够获得高达40%的传输加速。 从主机读取写结合存储非常慢（因为没有使用L1、L2cache），所以写结合存储应当只用于那些主机只写的存储。

Mapped Memory

一块锁页内存可以在调用cudaHostAlloc() 分配时传入cudaHostAllocMapped 标签或者在使用cudaHostRegister() 注册时使用cudaHostRegisterMapped 标签，把锁页内存地址映射到设备地址空间。这样，这块存储会有两个地址：一个是从cudaHostAlloc() 或malloc() 返回的在主机内存地址空间上；另一个在设备存储器上，可以通过cudaHostGetDevicePointer() 取得。内核函数可以使用这个指针访问这块存储。 cudaHostAlloc() 返回的地址指针一个的例外情况是，主机和设备使用统一地址空间（Unified Virtual Address Space）。
内核直接存取主机内存有很多优势：

• 无需在设备上分配内存，也无需在主机内存和设备内存之间拷贝数据。数据传输是在内核需要的时候隐式进行的。
• 无须使用流（cuda stream）就可以并发数据传输和内核执行；数据传输和内核执行自动并发执行。

因为映射的锁页主机内存是主机和设备之间共享的，所以在使用cuda stream 或者cuda event 时必须对内存读写同步；避免潜在的写后读，读后写或者写后写等多线程同步问题。
为了能够对任何映射的锁页主机内存解引用设备指针，必须在调用任何cuda 运行时函数前调用cudaSetDeviceFlags()，并传入cudaDeviceMapHost 标签。否则，cudaHostGetDevicePointer() 将会返回错误。
如果设备不支持被映射分页锁定存储，cudaHostGetDevicePointer() 将会返回错误。程序员可以检查canMapHostMemory 属性，如果设备支持映射锁页主机内存，将会返回1。

注意：使用映射锁页主机内存看，原子操作将不再保证原子性。cudaHostRegisterIoMemory 是cudaHostRegister() 特有的选项，可以把主机内存映射到IO 地址空间。

[1] https://en.wikipedia.org/wiki/CUDA_Pinned_memory
[2] Cook, Shane (2013). CUDA Programming: A Developer’s Guide to Parallel Computing with GPUs (1st ed.). Morgan Kaufmann Publishers Inc. pp. 334–335. ISBN 9780124159334.

零复制　　zero copy（零复制）是一种特殊形式的内存映射，它允许你将host内存直接映射到设备内存空间上。其实就是设备可以通过直接内存访问（direct memory access，DMA）方式来访问主机的锁页内存。 　　锁页主机内存　　现代操作系统都支持虚拟内存，操作系统实现虚拟内存的主要方法就是通过分页机制。操作系统将内存中暂时不使用的内容换出到外存（硬盘等大容量存储）上，从而腾出空间存 运行时提供的函数允许使用锁页（也称为固定）主机内存（与 malloc() 分配的常规可分页主机内存相反）： cuda HostAlloc() 和 cuda FreeHost() 分配和释放锁页主机内存； cuda HostRegister() 将 malloc() 分配的内存范围变为锁页内存（有关限制，请参阅参考手册）。 使用页面锁定的主机内存有几个好处： 锁页主机内存和设备内存之间的复制可以与异步并发执行中提到的某些设备的内核执行同时执行。 在某些设备上，锁页主机内存可以映射