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

转自《Linux阅马场》白嘉庆CUDA介绍
CUDA（ComputeUnifiedDeviceArchitecture ，统一计算设备架构）是由NVIDIA公司于2006年所推出的一种并行计算技术，是该公司对于GPGPU（General-purposecomputingongraphicsprocessingunits ，图形处理单元上的通用计算）技术的正式命名。通过此技术，用户可在GPU上进行通用计算，而开发人员可以使用C语言来为CUDA架构编写程序。相比CPU ，拥有CUDA技术的GPU成本不高，但计算性能很突出。本文中提到的是2014年发布的CUDA6 ， CUDA6最重要的新特性就是支持统一内存模型(UnifiedMemory) 。
注：文中经常出现“主机和设备” ，本文的“主机”特指CPU、“设备”特指GPU 。 CUDA6中的统一内存模型
NVIDIA在CUDA6中引入了统一内存模型（UnifiedMemory），这是CUDA历史上最重要的编程模型改进之一。在当今典型的PC或群集节点中， CPU和GPU的内存在物理上是独立的，并通过PCI-Express总线相连。在CUDA6之前，这是程序员最需要注意的地方。 CPU和GPU之间共享的数据必须在两个内存中都分配，并由程序直接地在两个内存之间来回复制。这给CUDA编程带来了很大难度。

文章图片
统一内存模型创建了一个托管内存池（apoolofmanagedmemory），该托管内存池由CPU和GPU共享，跨越了CPU与GPU之间的鸿沟。 CPU和GPU都可以使用单指针访问托管内存。关键是系统会自动地在主机和设备之间迁移在统一内存中分配的数据，从而使那些看起来像CPU内存中的代码在CPU上运行，而另一些看起来像GPU内存中的代码在GPU上运行。
在本文中，我将向您展示统一内存模型如何显著简化GPU加速型应用程序中的内存管理。下图显示了一个非常简单的示例。两种代码都从磁盘加载文件，对其中的字节进行排序，然后在释放内存之前使用CPU上已排序的数据。右侧的代码使用CUDA和统一内存模型在GPU上运行。和左边代码唯一的区别是，右边代码由GPU来启动一个内核（并在启动后进行同步），并使用新的APIcudaMallocManaged()在统一内存模型中为加载的文件分配空间。

文章图片
如果您曾经编程过CUDAC/C++ ，那么毫无疑问，右侧的代码会为您带来震撼。请注意，我们只分配了一次内存，并且只有一个指针指向主机和设备上的可访问数据。我们可以直接地将文件的内容读取到已分配的内存，然后就可以将内存的指针传递给在设备上运行的CUDA内核。然后，在等待内核处理完成之后，我们可以再次从CPU访问数据。 CUDA运行时隐藏了所有复杂性，自动将数据迁移到访问它的地方。统一内存模型提供了什么
统一内存模型为程序员提供了两大捷径简化编程、简化内存模型
统一内存模型通过使设备内存管理（devicememorymanagement）成为一项可选的优化，而不是一项硬性的要求，从而降低了CUDA平台上并行编程的门槛。借助统一内存模型，程序员现在可以直接开发并行的CUDA内核，而不必担心分配和复制设备内存的细节。这将降低在CUDA平台上编程的学习成本，也使得将现有代码移植到GPU的工作变得容易。但这些好处不仅有利于初学者。我在本文后面的示例中将展示统一内存模型如何使复杂的数据结构更易于与设备代码一起使用，以及它与C++结合时的强大威力。通过数据局部性原理提高性能
通过在CPU和GPU之间按需迁移数据，统一内存模型可以满足GPU上本地数据的性能需求，同时还提供了易于使用的全局共享数据。这个功能的复杂细节被CUDA驱动程序和运行时隐藏了，以确保应用程序代码更易于编写。迁移的关键是从每个处理器获得全部带宽。 250GB/s的GDDR5内存对于保证开普勒（Kepler）GPU的计算吞吐量至关重要。