找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 134|回复: 2

DAY13:

[复制链接]
发表于 2018-5-16 15:04:48 | 显示全部楼层 |阅读模式
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.

回复

使用道具 举报

 楼主| 发表于 2018-5-16 15:48:55 | 显示全部楼层
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。例如我知道我就能请你吃一次饭,以后没有机会了,那么果断这一次送来就送来吧。
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-5-16 16:03:59 | 显示全部楼层
sisiy 发表于 2018-5-16 15:48
Unified Virtual Address Space 统一的虚拟地址空间,这个空间包含:进程的传统Host虚拟地址空间,所有卡的 ...

event是目前唯一的能跨卡,跨context,跨stream的同步机制。所以能在进程间共享它也很重要。虽然说,完全可以不共享它,而是在单个CUDA进程内,例如等待某cuda event后,设定一个host上的,例如mutex来完成同步,但是能直接在多个进程间共享cuda event有的时候还是很方便的。而且直接使用它可能性能会更好,比迂回的转换成host上的同步对象的方式。例如考虑一个host进程A上的流a中,需要等待event 2,才能继续执行,而event 2来自host B上的流b中此时直接使用cuda event,可以在host B进程上的流b中的event 2发生后,host进程A中的流a就可以继续执行了,而传统的方式将CUDA同步后,转换成Host上的同步对象,再用这个继续让进程2继续,会带来额外的延迟之类的开销。直接用不仅仅简单方便,还可能性能会更好。 CUDA event这个共享了其实很方便,特别是对于一些对Host不熟悉,专心GPU的kernel开发的员工来说,
这样他可以继续使用传统的他熟悉的CUDA。而不需要去纠缠host上的各种通信方式了。还是有很大意义的。
回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表