DAY13：CUDA C Runtime之统一虚拟地址空间及进程间通信

3.2.7. Unified Virtual Address Space【统一虚拟地址空间】

When the application is run as a 64-bit process, a single address space is used for the host and all the devices of compute capability 2.0 and higher. All host memory allocations made via CUDA API calls and all device memory allocations on supported devices are within this virtual address range. As a consequence:

· The location of any memory on the host allocated through CUDA, or on any of the devices which use the unified address space, can be determined from the value of the pointer usingcudaPointerGetAttributes().

· When copying to or from the memory of any device which uses the unified address space, the cudaMemcpyKind parameter of cudaMemcpy*() can be set to cudaMemcpyDefault to determine locations from the pointers. This also works for host pointers not allocated through CUDA, as long as the current device uses unified addressing.

· Allocations via cudaHostAlloc() are automatically portable (see Portable Memory) across all the devices for which the unified address space is used, and pointers returned bycudaHostAlloc() can be used directly from within kernels running on these devices (i.e., there is no need to obtain a device pointer via cudaHostGetDevicePointer() as described in Mapped Memory.

Applications may query if the unified address space is used for a particular device by checking that the unifiedAddressing device property (see Device Enumeration) is equal to 1.

3.2.8. Interprocess Communication【进程间通信】

Any device memory pointer or event handle created by a host thread can be directly referenced by any other thread within the same process. It is not valid outside this process however, and therefore cannot be directly referenced by threads belonging to a different process.

To share device memory pointers and events across processes, an application must use the Inter Process Communication API, which is described in detail in the reference manual. The IPC API is only supported for 64-bit processes on Linux and for devices of compute capability 2.0 and higher.

Using this API, an application can get the IPC handle for a given device memory pointer using cudaIpcGetMemHandle(), pass it to another process using standard IPC mechanisms (e.g., interprocess shared memory or files), and use cudaIpcOpenMemHandle() to retrieve a device pointer from the IPC handle that is a valid pointer within this other process. Event handles can be shared using similar entry points.

An example of using the IPC API is where a single master process generates a batch of input data, making the data available to multiple slave processes without requiring regeneration or copying.

本文备注/经验分享：

Unified Virtual Address Space 统一的虚拟地址空间，这个空间包含：进程的传统Host虚拟地址空间，所有卡的虚拟地址空间。也就是CPU + GPU（多个）。用人话说就是，将你分配的普通malloc(), 每个卡上的cudaMalloc()出来的，这些得到的分配出来的缓冲区地址，都在同一个64-bit的进程虚拟地址空间内。可以直接使用一个普通的指针Type *p指向，而不是每个分配的指针只在每个设备上才有意义。以前的我们会往往遇到这种情况: 我在CPU上分配到地址int *p是0x12345678，然后在GPU上分配到的地址也是0x12345678，用户必须明确的知道这个地址是在哪里有效的，才能用它。（因为以前不是统一编址的，大家各自为战）所以你会看到以前cudaMemcpy之类的函数，指定了目标地址，源地址，传输大小等信息后， 却需要额外的添加一个类似cudaMemcpyHostToDevice这种参数告诉CUDA Runtime，源地址是从Host来的，目标地址是在设备（卡）上的。现在统一编址后，不需要用户维护这个信息了，直接CUDA就能知道，哦，这个地址是卡1上的，这个地址是卡2上的，这个地址是卡3的， 这个地址是内存，这个地址是自动管理的（unified memory）...类似这种，方便了很多。也为以后实现很多功能打下了基础。这个是个老特性，从Fermi开始的，但是有了这个基础，我们现在用Pascal，跨卡P2P Access（你还记得这个是什么吗，昨天才讲过？） ， 直接卡1上的kernel，能够使用一个指针p，而p指向的内容却在卡2上，没有这个基础，P2P Access无法实现。类似的，这还为其他特性，例如现在的unified memory，假设你有一个链表，非常巨大， CPU想负责一部分适合它处理的里面的节点数据，GPU想处理一部分它想处理的，以前的写法只能是每个节点标记一下，例如： 本节点标记为是内存上，必须用CPU处理，本节点链接到的下一个节点是在GPU上，这个下一个节点的指向的指针必须GPU有效，CPU不能处理，云云的。 现在统一编址后，可以直接获取某个节点在哪里，甚至通过UVA + Unified memory，程序员偶尔不小心用CPU处理了某个应当GPU处理的节点（或者反过来），也不要紧， Runtime/Driver自动给你迁移了位置，处理起来很方便。 再比如，以前很多显卡没有显存，（很多笔记本的集成的N卡，虽然支持CUDA。但没有显存）， 用户以前都用zero-copy，但是zero-copy以前有个问题，同样的一段缓冲区，例如100MB，它在CPU上的地址，和在GPU上的地址是不同的， 用户必须同时保存两份指针信息，一个指针是host上有效的，一个指针是GPU上有效的。用错了，程序就挂了。当Fermi开始，引入了UVA后， 这两个地址变成了同样的值，用户知道int *p可以在host上用，也可以直接在GPU上用，不仅仅简单了很多，还减少了很大的出错可能。很是方便的。 这个是一个巨大的基础改进。当年Fermi引入的和UVA同样的改进还有一部分，叫Generic Addressing。UVA是全局的（卡，CPU，多卡），Generic Addressing是卡内部的，pre-fermi的时候，卡内部的地址也不是统一的，local memory, shared memory, global memory是分裂的，一个指针必须需要在编译时刻知道指向哪里，否则不能使用。 就像DOS时代的segment一样难用。fermi起，将卡内，卡间（系统内）都统一了。一个指针可以打天下了。相当方便和给力。 这其实主要是为了易用性，对性能其实无提升的。

Interprocess Communication 可以将一个context（或者你理解成的使用了CUDA的进程）内分配的显存，共享给另外一个context（或者你理解成另外一个进程）用。这个特性需要Linux的。Windows下不能用。 你可以将它理解成CUDA版的CPU上的共享内存机制。 CPU上有IPC机制，可以在进程间共享一些信息/东西。其中的一个重要的点是共享内存。A,B两个CPU上的进程，可以同时将一段内存映射到自己的地址空间。CUDA IPC的道理和这个类似，只不过变成了是显存。这样一些在显存中的数据，两个进程可以共享或者交换信息。手册这里提到CPU部分主要是因为两点： （1）CUDA IPC和CPU上的共享内存很像，只不过是共享显存。（2）CUDA IPC需要通过CPU上的普通IPC才能建立。因为需要交换一些数据（显存句柄）；如何通过CPU上的IPC机制来交换这个句柄信息，这里没说。因为这个是常规的OS上的能力，正常人都应该直接掌握，而非CUDA的一部分。 没有CUDA IPC，常规做法是：Host进程1 cudaMemcpy 显存到内存， Host进程1和Host进程2通过Host上的IPC机制（例如pipe之类的）传输内容，Host进程2在将得到的内存复制到显存，而有了CUDA IPC后，直接在host进程1和2之间传递一个很小的句柄，就可以直接共享这段显存了。节省了大量的显存 -> 内存 -> 显存的复制时间。 就如同我想请你吃饭，一种办法是我去餐厅，拿到饭，送给你，然后你再吃掉；另外一种办法是我去办理一张会员卡（句柄），很薄很轻，然后我将卡给你，你直接就可以去吃饭了。不用我搬运沉重的饭。 CUDA IPC在较多的数据量的时候很有用。就如同刚才的，吃N顿，累计1000元的饭，如果我每次都去搬运饭菜给你，很累。 CUDA IPC在较多的数据量的时候很有用。 一次办理一张1000元的卡，轻松解决问题。当然小数据量的情况下可以无视CUDA IPC。例如我知道我就能请你吃一次饭，以后没有机会了，那么果断这一次送来就送来吧。

有不明白的地方，请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖