Unified Memory：
CUDA 6.0开始支持Unified Memory，在CUDA 6.0之前，需要手动的在Host和GPU之间分配内存，并在两者之间不断地进行拷贝（cudaMemcpy），即需要自己进行CPU和GPU之间地内存管理。 采用UM后，从程序员的视角，可通过一个统一的指针进行内存管理，由系统自动的迁移内存。
使用UM机制，不能再用cudaMalloc分配GPU memory上的内存，应该使用cudaMallocManaged
cudaMallocManaged ( void** devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal )
UM机制极大地减少了代码量，可以很大程度上减少程序员的工作量。在UM出现之前，由于CPU和GPU之间的地址空间是各自独立的，需要进行多次的手动分配和频繁地使用cudaMemcpy在CPU和GPU的memory之间来回拷贝内存。
当实际数据结构更加复杂时，会使得内存管理变得很复杂。
UM的例子：
如上图程序示例图和运行结果图所示：在CPU端可以直接访问GPU内存中的ret[i]的值
一个UM的示例程序，直接使用cudaMallocManaged分配GPU内存，指针ret指向GPU的内存，在CPU端同样可以直接访问ret指针。
UM机制使得程序员不用频繁地使用cudaMemcpy在GPU和CPU之间来回拷贝数据。
需要注意的是，UM不会消除CPU Memory和GPU Memory之间的拷贝，这部分copy工作交给CUDA执行，程序员不感知CPU Memory和GPU memory之间的数据拷贝，但copy依然是存在的。
Unified Memory page_migration机制：
GPU和CPU一样，有自己的页表和TLB，当CPU或GPU采用UM机制访问一块内存时，CPU和GPU均可能发生page falut，从而导致memory在CPU和GPU之间来回拷贝。
上图是一个使用UM机制的例子，我们来看一下这个例子的具体流程：
1 首先使用cudaMallocManaged分配GPU内存
2 在fill_data中，CPU会写这块内存，此时CPU触发page faｕｌｔ流程，GPU memory中的内容通过PCIE migrate到GPU中。
同时为了数据保证一致性，CUDA会invalidateGPU中对应的页表项 ，这样下一次GPU访问的时候就会触发GPU端的页表项。
３qsort<<<...>>>中，GPU端执行kernel函数，由于之前GPU的页表已经无效了，此时触发GPU端的page falut，数据通过PCIE从CPU migrate至GPU中。然后CPU短的页表项被invalidate。
4 use_data中，CPU处理数据，由于CPU的页表项已经在上一步被invalidate了，于是CPU端触发page falut, 数据再从CPUmigrate至GPU。
由此我们可以看到，其实UM机制中，数据就是来回在CPU和GPU的memory中来回migrate，只是程序员在上层并不感知。
为什么不采用zero-copy机制？
如果数据内容只有一份，那么就无法利用data locality，这样会影响性能。
GDR(GPU Direct RDMA)实现：
GDR，即GPU direct RDMA，可以在GPU中直接分配memory。然后直接在GPU和device上进行数据传输。 Device可以是NIC，storage adapters，video acquisition devices这些PCIE设备。
想要使用GDR功能，device和GPU需要在同一个RC下。
GDR Example：
传统RDMA程序注册内存： GDR注册内存：
char *addr; char *addr;
struct ibv_mr *mr; struct ibv_mr *mr;
size_t length; size_t length;
addr = malloc(length); cudaMalloc(&addr, length);
mr = ibv_reg_mr(pd, addr, length, access); mr = ibv_reg_mr(pd, addr, length, access);
... ...
从上面的程序例子可以看到，其实GDR和传统的RDMA程序几乎没有什么不同，最主要的不同就在于，传统的RDMA程序采用malloc分配CPU memory上面的内存，而GDR程序中采用cudaMalloc分配位于GPU memory上面的内存。
GDR的使用：
使能GDR功能需要插入nv_peer_mem内核模块： insmod nv_peer_mem 插入nv_peer_mem内核模块后，调用ibv_reg_mr后将会在GPU memory中注册内存而不是在CPU DDR中注册memory。
GDR的实现 ：
B提供了一种叫做peer_memory_client的机制，第三方设备在自己的_init函数中调用ib_register_peer_memory_client在ib_core中注册了一个peer_memory_client。
insmod nv_peer_mem.ko时会执行nv_mem_client_init函数，其中就调用了ib_register_peer_memory_client向ib_core注册了一个peer_memory_client 。
peer_memory.h中定义了各种接口，使得注册为peer_memory_client的第三方device能够实现自己的get_pages，map等相关的接口。peer_memory_client的实现如下：
peer_memory_client在peer_mem.h定义，位于mlnx-ofed-kernel-4.7\include\rdma
每个第三方设备都可以实现自己的peer_memory_client，然后调用ib_register_peer_memory_client往ib_core里面注册为一个peer_memory_client。 每个peer_memory_client都是链表peer_memory_list中的一个元素。 这样任何实现了peer_memory_client接口的设备都可以实现Direct RDMA的功能。
ibv_reg_mr的流程如上，最终会调用到ib_client_umem_get函数，其中就会调用到peer_memory_client自己注册的acquire，get_pages，dma_map等函数。
如上图所示，nv_mem_client_ex中注册了自己的acuqire，get_pages，dma_map等函数。当insmod nv_peer_mem.ko时，最终执行的
peer_mem->acquire( )就是nv_mem_acquire，peer_mem->get_pages( )就是nv_mem_get_pages，peer_mem->dma_map( )就是
nv_dma_map。
acquire函数返回1，表示ib_core找到了对应的peer_memory_client，之后就由这个peer_memory_client来handle这块注册的内存。其中nv_peer_mem中的acquire最终会调用到nvidia_p2p_get_pages通过虚拟页面获得物理页面，如果成功，则表明该peer_memory_client能够处理这块物理内存，最后调用nvidia_p2p_put_pages解除VA到PA的映射。acquire主要功能就是表示这块注册的peer_memory内存能够由该peer_memory_client来handle（成功返回1，失败则返回0）。
nv_mem_get_pages中调用了nvidia_p2p_get_pages通过传入虚拟地址（nv_mem_context->page_virt_start），获得物理地址，并将结果保存在nv_mem_context->page_table。
nv_dma_map中调用nvidia_p2p_dma_map_pages获取dma_address，page_size等，并填入sg_table中(用于scatter gather DMA)
