linux内核 6中的统一内存模型，CUDA( 三 ) 转自《Linux阅马场》白嘉庆CUDA介绍

将具有复杂数据结构的代码移植到GPU上曾经是一项艰巨的任务，但是统一内存模型使此操作变得非常容易。我希望统一内存模型能够为CUDA程序员带来巨大的生产力提升。 UnifiedMemorywithC++
统一内存模型确实在C++数据结构中大放异彩。 C++通过带有拷贝构造函数（copyconstructors）的类来简化深度复制问题。拷贝构造函数是一个知道如何创建类所对应对象的函数，拷贝构造函数为对象的成员分配空间并从其他对象复制值过来。 C++还允许new和delete这俩个内存管理运算符被重载。这意味着我们可以创建一个基类，我们将其称为Managed ，它在重载的new运算符内部使用cudaMallocManaged() ，如以下代码所示。 classManaged{public:void*operatornew(size_tlen){void*ptr;cudaMallocManaged(&ptr,len);cudaDeviceSynchronize();returnptr;}voidoperatordelete(void*ptr){cudaDeviceSynchronize();cudaFree(ptr);}};
然后，我们可以让String类继承Managed类，并实现一个拷贝构造函数，该拷贝构造函数为需要拷贝的字符串分配统一内存。 //Derivingfrom“Managed”allowspass-by-referenceclassString:publicManaged{intlength;char*data;public://Unifiedmemorycopyconstructorallowspass-by-valueString(constString&s){length=s.length;cudaMallocManaged(&data,length);memcpy(data,s.data,length);}//...};
同样，我们使我们的dataElem类也继承Managed 。 //Note“managed”onthisclass,too.//C++nowhandlesourdeepcopiesclassdataElem:publicManaged{public:intprop1;intprop2;Stringname;};
通过这些更改， C++的类将在统一内存中分配空间，并自动处理深度复制。我们可以像分配任何C++的对象那样在统一内存中分配一个dataElem 。 dataElem*data=https://pcff.toutiao.jxnews.com.cn/p/20200703/newdataElem;
请注意，您需要确保树中的每个类都继承自Managed ，否则您的内存映射中会有一个漏洞。实际上，任何你想在CPU和GPU之间分享的内容都应该继承Managed 。如果你倾向于对所有程序都简单地使用统一内存模型，你可以在全局重载new和delete ，但这只在这种情况下有作用——你的程序中没有仅被CPU访问的数据（即程序中的所有数据都被GPU访问），因为只有CPU数据时没有必要迁移数据。
现在，我们可以选择将对象传递给内核函数了。如在C++中一样，我们可以按值传递或按引用传递，如以下示例代码所示。 //Pass-by-referenceversion__global__voidkernel_by_ref(dataElem&data){...}//Pass-by-valueversion__global__voidkernel_by_val(dataElemdata){...}intmain(void){dataElem*data=https://pcff.toutiao.jxnews.com.cn/p/20200703/newdataElem;...//passdatatokernelbyreferencekernel_by_ref<<
>>(*data);//passdatatokernelbyvalue--thiswillcreateacopykernel_by_val<<
>>(*data);}
多亏了统一内存模型，深度复制、按值传递和按引用传递都可以正常工作。统一内存模型为在GPU上运行C++代码提供了巨大帮助。
这篇文章的例子可以在Github上找到。统一内存模型的光明前景
CUDA6中关于统一内存模型的最令人兴奋的事情之一就是它仅仅是个开始。我们针对统一内存模型有一个包括性能提升与特性的长远规划。我们的第一个发行版旨在使CUDA编程更容易，尤其是对于初学者而言。从CUDA6开始， cudaMemcpy()不再是必需的。通过使用cudaMallocManaged() ，您可以拥有一个指向数据的指针，并且可以在CPU和GPU之间共享复杂的C/C++数据结构。这使编写CUDA程序变得容易得多，因为您可以直接编写内核，而不是编写大量数据管理代码并且要维护在主机和设备之间所有重复的数据。您仍然可以自由使用cudaMemcpy()（特别是cudaMemcpyAsync()）来提高性能，但现在这不是一项要求，而是一项优化。