GPU地址空间的相关概念

转载于：http://blog.youkuaiyun.com/xiewen_bupt/article/details/47208903

知识点回顾：虚拟地址

一个进程要对内存的某存储单元进行读写操作时，需要该内存单元的内存位置，这个时候，通过物理寻址计算出该内存位置。然而不同的进程可能修改相同位置的内存导致程序间的相互影响，因此，计算机系统引入了虚拟地址。

1. 每个进程都有自己的虚拟地址空间，相互之间无法读取其他进程的内存。
2. 在大多数的系统中，虚拟地址空间分为许多页（>=4KB）。
3. 虚拟地址与物理地址的映射表称为页表。这个转换过程称之为 程序的再定位。
4. 32位操作系统，每个进程都有4GB大小空间的内存地址，因此一般来说，物理内存是不够用的。这个时候就引入了虚拟内存（如在磁盘中开辟空间）。

CUDA的地址空间

CUDA也使用了虚拟地址空间，为的是多个CUDA程序不会相互破坏GPU内存数据。然而，GPU没有给对内存的扩展，因此，虚拟地址通过页表转化后的物理地址必须对应物理内存。

在CUDA 1.0中，CPU与GPU各自独立拥有各自的虚拟地址空间和独立内存，必须显示的将数据拷贝才能使用对方内存中的数据。在这里，GPU可以在CPU内存中分配分页锁定内存（page-locked memory），该内存是分配在CPU物理内存中的，是DMA映射的。当数据从分页锁定内存拷贝到GPU内存中时，CUDA驱动程序会自动加速该内存拷贝的操作。

另外一点，异步拷贝操作要求CPU段的内存必须是分页锁定的，需要分页锁定内存起址信息，以确保内存复制完成之前，操作系统不会取消映射或者移动物理内存。

映射分页锁定内存 mapped memory

CUDA 2.2 提供了可映射分页锁定内存的特性。CPU端的分页锁定内存可以映射到GPU地址空间中，也就是说GPU端的页表包含了GPU虚拟地址与CPU物理地址的映射。这意味着在GPU端可以直接读取CPU端内存中的数据。这种方式我们称之为zero-copy。

需要注意几点：

1. CPU端和GPU端的虚拟地址空间仍然是独立的。只不过在GPU端，页表中存在了关于CPU端分页锁定内存的映射关系。
2. 由于存在不同的虚拟地址空间，因此这段映射的分页锁定内存存在两种形式：一种是，在CPU的指针，由cudaMallocHost返回；另外一种，是在GPU端的指针，需要通过 cudaHostGetDevicePointer()获得（记得之前需要调用cudaSetDeviceFlags()函数来设置cudaDeviceMapHost标志）。
3. 在多卡的情况下，每个GPU都会有一个上下文。则如果我们将分页锁定内存设置为“可分享的”，即cudaHostAllocPortabled标志，那么这个映射将会存在所有GPU的CUDA上下文中。
4. 随着统一虚拟地址（UVA）概念的出现，这种方式将不再推荐了。具体的见下一小节的讲解。

统一虚拟寻址 unified virtual addressing

在此之前我们讲到的两种情况，CPU端和GPU端都是独立的虚拟地址空间。在CUDA 4.0中，引入了统一虚拟寻址UVA的特性。顾名思义，在支持UVA的情况下（支持的系统会自动启动UVA），CPU内存和GPU内存的虚拟地址空间是一个整体了，不再是分开的。

几点编程上的便利：

1. 对于CPU端的分页非锁定内存和GPU端的内存，CUDA可以推断出来该内存指针位于那个设备上面。因此，cudaMemcpy()函数最后的形参（拷贝方向）不需要指定具体的拷贝方向了，填写一个cudaMemcpyDefault即可，剩下的就交给CUDA来辨别拷贝方向了。
2. 对于CPU端分页锁定内存，在UVA模式下，默认被映射和可分享的。而且CPU端和GPU端读写该内存的指针是相同的，只需要一个指针即可。一点例外：使用cudaHostRegister转化的CPU端分页锁定内存指针，设备端的指针还需要cudaHostGetDevicePointer()函数获得。

点对点寻址 peer-to-peer

在多卡环境下编程时，我们会碰到GPU与GPU之间的数据拷贝，使用的函数为（非点对点映射，只能称之为点对点内存拷贝）：

[cpp]  view plain  copy

1. __host__ ​cudaError_t cudaMemcpyPeer ( void* dst, int  dstDevice, const void* src, int  srcDevice, size_t count )  
2. __host__ ​cudaError_t cudaMemcpyPeerAsync ( void* dst, int  dstDevice, const void* src, int  srcDevice, size_t count, cudaStream_t stream = 0 )  

在这里的两个函数，指定了源设备和目的设备的ID，指定了拷贝的两个指针，指定了拷贝的大小。

当系统支持UVA时，我们就不必指定两个设备了，CUDA可以通过指针推断出来源于那个设备。因此，只需简化使用cudaMemcpy()函数就可以了。

[cpp]  view plain  copy

1. __host__ ​cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )  

之前的这两种GPU之间数据拷贝其实是以CPU内存为媒介进行拷贝的。当我们使用nvvp软件查看GPU之间数据传输的timeline时，我们就发现整个拷贝包含两个步骤：

1. SrcDevice to CPU (D2H)
2. CPU to dstDevice (H2D)

这个是很耗时的。其实有一种更快的点对点寻址方式，peer-to-peer拷贝。这个特性不是默认打开的，通常我们需要这两步骤：

第一步，判断设备A是否可以通过P2P方式访问设备B：

[cpp]  view plain  copy

1. __host__ ​cudaError_t cudaDeviceCanAccessPeer ( int* canAccessPeer, int  device, int  peerDevice )  

第二步，如果可以，那么将P2P方式打开：

[cpp]  view plain  copy

1. __host__ ​cudaError_t cudaDeviceEnablePeerAccess ( int  peerDevice, unsigned int  flags )  

注意，这两步只是满足了设备A对设备B内存P2P访问，而设备B却无法对设备A P2P访问，因为这个是非对称的，仍需要使设备B对设备A做同样的操作。

当我们使用nvvp软件查看GPU之间数据传输的timeline时，我们就会发现，这是GPU之间的拷贝只包含了一个P2P的过程了。

在最后的测试用例中有相关的测试代码。

NVIDIA GPUDirect

什么是GPUDirect？ GPUDirect是GPU与其他设备通信的技术概括，包含了一系列的技术特性。我们在这里列举一下：

1. 2010年，GPU可以与处于PCIE总线上的设备通过共享pinned memory来进行通信，取消了一次CPU内存中分页锁定内存向另外一段内存的拷贝过程。

2. 2011年，增加了P2P特性，也就是我们刚才讲的GPU之间P2P的通信。因为取消了CPU内存的“媒介”功能，使得速度提升。

3. 2013年，引入了RDMA的特性，这个特性将使得GPU与第三方存储设备之间可以直接数据通信，同样取消了CPU内存的“媒介”功能。

测试用例

这是一个P2P数据拷贝的简单测试代码，第一种拷贝没有打开P2P开关，第二种拷贝打开了P2P开关。

[cpp]  view plain  copy

1. #include <stdio.h>  
2. #include <stdlib.h>  
3. #include <cuda_runtime.h>  
4. #include <helper_functions.h>  
5. #include <helper_cuda.h>  
6. #include <helper_string.h>  
7.   
8. const int num = 1 << 25;  
9. const int num_gpu = 2;  
10.   
11. int main()  
12. {  
13.     int id_gpu;  
14.     float *data_d[num_gpu];  
15.     //malloc      
16.     for(id_gpu = 0; id_gpu < num_gpu; id_gpu++){  
17.         checkCudaErrors( cudaSetDevice(id_gpu) );  
18.         checkCudaErrors( cudaMalloc((void **)&data_d[id_gpu], num * sizeof(float)) );  
19.     }  
20.     //copy data from device 0  to device 1.  
21.     checkCudaErrors( cudaSetDevice(0) );  
22.     checkCudaErrors( cudaMemcpy(data_d[1], data_d[0], num * sizeof(float), cudaMemcpyDefault) );  
23.           
24.     //p2p copy data.  
25.     for (int i = 0; i < num_gpu; i++) {  
26.         int p2p;  
27.         for (int j = i+1; j < num_gpu; j++) {  
28.             checkCudaErrors( cudaDeviceCanAccessPeer(&p2p, i, j) );  
29.             if (p2p) {  
30.                 printf("P2P support between device %d and %d.\n", i, j);  
31.                 checkCudaErrors( cudaSetDevice(i) );  
32.                 checkCudaErrors( cudaDeviceEnablePeerAccess(j, 0) );  
33.                 checkCudaErrors( cudaSetDevice(j) );  
34.                 checkCudaErrors( cudaDeviceEnablePeerAccess(i, 0) );  
35.             } else {  
36.                 printf("No P2P support between device %d and %d.\n", i, j);  
37.             }  
38.         }  
39.     }  
40.   
41.     checkCudaErrors( cudaSetDevice(0) );  
42.     checkCudaErrors( cudaMemcpy(data_d[1], data_d[0], num * sizeof(float), cudaMemcpyDefault) );  
43.       
44.     //cudaFree  
45.     for(id_gpu = 0; id_gpu < num_gpu; id_gpu++){  
46.         checkCudaErrors( cudaSetDevice(id_gpu) );  
47.         checkCudaErrors( cudaFree(data_d[id_gpu]) );  
48.     }  
49.   
50.     checkCudaErrors( cudaDeviceReset() );     
51.     return 0;  
52. }  

测试结果截图（nvvp）：