专栏首页吉浦迅科技DAY13:CUDA C Runtime之统一虚拟地址空间及进程间通信

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上发帖

本文分享自微信公众号 - 吉浦迅科技(gpusolution)

原文出处及转载信息见文内详细说明,如有侵权,请联系 yunjia_community@tencent.com 删除。

原始发表时间:2018-05-16

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

我来说两句

0 条评论
登录 后参与评论

相关文章

  • NVIDIA Jetson OpenCV开发实战教程(中)

    从一个应用程序开始,它将图像显示为一个Mat对象,然后调整大小、旋转图像或检测“canny”的边缘,再显示结果。然后,为了忽略图像feather的高频边缘,模糊...

    GPUS Lady
  • CUDA程序媛之友谊的船

    北京海淀,太平洋吹来暖湿的季风,学霸正在疯长,又到了大学生们最忙碌的季节——配置CUDA环境。在导师眼中,GPU能为学生发毕业论文带来好运,值得为它冒险。现代社...

    GPUS Lady
  • DAY72:阅读Toolkit Support for Dynamic Parallelism

    我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第72天,我们正在讲解CUDA 动态并行,希望在接下来的28天里,您可以...

    GPUS Lady
  • 国外大牛教你,如何用Python开发一个简单的区块链数据结构| 建议收藏

    根据IEEE此前的一项调查,Python已成为最受开发者欢迎的语言之一。由于其对于技术小白天然友好的特性,以及不断更新的新功能。Python越来越受到国内外开发...

    区块链大本营
  • 【区块链实践】杭州互联网法院司法区块链解决方案1,司法区块链解决方案2,司法区块链可以解决哪些问题?3,为什么要运行司法区块链?4, 参考

    据杭州日报消息,2018年09月18日,杭州互联网法院宣布司法区块链正式上线运行,成为全国首家应用区块链技术定纷止争的法院。司法区块链让电子数据的生成、存储、传...

    辉哥
  • CUDA程序媛之友谊的船

    北京海淀,太平洋吹来暖湿的季风,学霸正在疯长,又到了大学生们最忙碌的季节——配置CUDA环境。在导师眼中,GPU能为学生发毕业论文带来好运,值得为它冒险。现代社...

    GPUS Lady
  • 用表情符号解释比特币 - Part 1

    比特币是一个存储和消费数字货币的革命性方式,并且有着变革其他领域的潜力。无须成为一个数学家或密码学家, 你就可以理解它是怎么回事。当开始看到整个系统是如何形成时...

    用户1558438
  • Android Gradle实用技巧(四) | 自动瘦身APK文件

    随着工程越来越大,功能越来越多,开发人员越来越多,代码越来越复杂,不可避免的会产生一些不在使用的资源,这类资源如果没有清理的话,会增加我们Apk的包大小,也会增...

    飞雪无情
  • 单价最高80万的电子宠物猫,上线一周交易量突破4400万,区块链技术与游戏结合真的前景无限?

    近期加拿大Axiom Zen工作室开发的一款名为CryptoKitties的电子宠物养成游戏,忽然间登上了全球各地的媒体。这款游戏的玩法比较简单,玩家可以从开发...

    企鹅号小编
  • 埃航黑匣子逐步破解,坠机矛头直指全自动化飞行软件

    场景描述:自动驾驶是航空技术的标配,对于长途飞行、飞行时的常规操作能都有效的配合飞行员完成,但还存在不少问题需要解决。

    HyperAI超神经

作者介绍

GPUS Lady

苏州吉浦迅科技有限公司联合创始人

苏州吉浦迅科技有限公司 · 联合创始人 (已认证)

精选专题

活动推荐

扫码关注云+社区

领取腾讯云代金券