江苏龙网

  1. 首页 > 资讯 > 科技 > >

linux内核|CUDA 6中的统一内存模型( 三 )


本文插图

链表是一种非常常见的数据结构 ， 但是由于它们本质上是由指针组成的嵌套数据结构 ， 因此在内存空间之间传递它们非常复杂 。 如果没有统一内存模型 ， 则无法在CPU和GPU之间分享链表 。 唯一的选择是在零拷贝内存（被pin住的主机内存）中分配链表 ， 这意味着GPU的访问受限于PCI-express性能 。 通过在统一内存模型中分配链表数据 ， 设备代码可以正常使用GPU上的指针 ， 从而发挥设备内存的全部性能 。 程序可以维护单链表 ， 并且无论在主机或设备中都可以添加和删除链表元素 。
将具有复杂数据结构的代码移植到GPU上曾经是一项艰巨的任务 ， 但是统一内存模型使此操作变得非常容易 。 我希望统一内存模型能够为CUDA程序员带来巨大的生产力提升 。 Unified Memory with C++
【linux内核|CUDA 6中的统一内存模型】统一内存模型确实在C++数据结构中大放异彩 。 C++通过带有拷贝构造函数（copy constructors）的类来简化深度复制问题 。 拷贝构造函数是一个知道如何创建类所对应对象的函数 ， 拷贝构造函数为对象的成员分配空间并从其他对象复制值过来 。 C++还允许 new和 delete这俩个内存管理运算符被重载 。 这意味着我们可以创建一个基类 ， 我们将其称为 Managed ， 它在重载的 new运算符内部使用 cudaMallocManaged() ， 如以下代码所示 。 class Managed {public:void *operator new(size_t len) {void *ptr;cudaMallocManaged(&ptr, len);cudaDeviceSynchronize();return ptr;}void operator delete(void *ptr) {cudaDeviceSynchronize();cudaFree(ptr);}};
然后 ， 我们可以让 String类继承 Managed类 ， 并实现一个拷贝构造函数 ， 该拷贝构造函数为需要拷贝的字符串分配统一内存 。 // Deriving from “Managed” allows pass-by-referenceclass String : public Managed {int length;char *data;public:// Unified memory copy constructor allows pass-by-valueString (const String &s) {length = s.length;cudaMallocManaged(&data, length);memcpy(data, s.data, length);}// ...};
同样 ， 我们使我们的 dataElem类也继承 Managed 。 // Note “managed” on this class, too.// C++ now handles our deep copiesclass dataElem : public Managed {public:int prop1;int prop2;String name;};
通过这些更改 ， C++的类将在统一内存中分配空间 ， 并自动处理深度复制 。 我们可以像分配任何C++的对象那样在统一内存中分配一个 dataElem 。 dataElem *data = http://news.hoteastday.com/a/new dataElem;
请注意 ， 您需要确保树中的每个类都继承自 Managed ， 否则您的内存映射中会有一个漏洞 。 实际上 ， 任何你想在CPU和GPU之间分享的内容都应该继承 Managed 。 如果你倾向于对所有程序都简单地使用统一内存模型 ， 你可以在全局重载 new和 delete ，但这只在这种情况下有作用——你的程序中没有仅被CPU访问的数据（即程序中的所有数据都被GPU访问） ， 因为只有CPU数据时没有必要迁移数据 。
现在 ， 我们可以选择将对象传递给内核函数了 。 如在C++中一样 ， 我们可以按值传递或按引用传递 ， 如以下示例代码所示 。 // Pass-by-reference version__global__ void kernel_by_ref(dataElem &data) { ... }// Pass-by-value version__global__ void kernel_by_val(dataElem data) { ... }int main(void) {dataElem *data = http://news.hoteastday.com/a/new dataElem;...// pass data to kernel by referencekernel_by_ref<<
>>(*data);// pass data to kernel by value -- this will create a copykernel_by_val<<
>>(*data);}
多亏了统一内存模型 ， 深度复制、按值传递和按引用传递都可以正常工作 。 统一内存模型为在GPU上运行C++代码提供了巨大帮助 。


推荐阅读


上一篇：龚树生教授|什么是声电联合刺激适合什么人？

下一篇：数码小AI资讯号|redmi K40倒计时，或首发LCD屏下指纹，跑分超50万！