DAY91：阅读Programming Model

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第91天，我们正在讲解Unified Memory Programming，希望在接下来的10天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。

关注微信公众号，可以看到之前的阅读

本文共计621字，阅读时间15分钟

K.2. Programming Model

K.2.1. Managed Memory Opt In

Most platforms require a program to opt in to automatic data management by either annotating a __device__ variable with the __managed__ keyword (see the Language Integration section) or by using a new cudaMallocManaged() call to allocate data.

Devices of compute capability lower than 6.x must always allocate managed memory on the heap, either with an allocator or by declaring global storage. It is not possible either to associate previously allocated memory with Unified Memory, or to have the Unified Memory system manage a CPU or a GPU stack pointer.

Starting with CUDA 8.0 and on supporting systems with devices of compute capability 6.x, memory allocated with the default OS allocator (e.g. malloc or new) can be accessed from both GPU code and CPU code using the same pointer. On these systems, Unified Memory is the default: there is no need to use a special allocator or the creation of a specially managed memory pool.

K.2.1.1. Explicit Allocation Using cudaMallocManaged()

Unified memory is most commonly created using an allocation function that is semantically and syntactically similar to the standard CUDA allocator, cudaMalloc(). The function description is as follows:

The cudaMallocManaged() function reserves size bytes of managed memory and returns a pointer in devPtr. Note the difference in cudaMallocManaged() behavior between various GPU architectures. By default, the devices of compute capability lower than 6.x allocate managed memory directly on the GPU. However, the devices of compute capability 6.x and greater do not allocate physical memory when calling cudaMallocManaged(): in this case physical memory is populated on first touch and may be resident on the CPU or the GPU. The managed pointer is valid on all GPUs and the CPU in the system, although program accesses to this pointer must obey the concurrency rules of the Unified Memory programming model (see Coherency and Concurrency). Below is a simple example, showing the use of cudaMallocManaged():

A program’s behavior is functionally unchanged when cudaMalloc() is replaced with cudaMallocManaged(); however, the program should go on to eliminate explicit memory copies and take advantage of automatic migration. Additionally, dual pointers (one to host and one to device memory) can be eliminated.

Device code is not able to call cudaMallocManaged(). All managed memory must be allocated from the host or at global scope (see the next section). Allocations on the device heap using malloc() in a kernel will not be created in the managed memory space, and so will not be accessible to CPU code.

本文备注/经验分享：

我们之前的关于Unified Memory的章节，简单介绍了它的作用（简化，便捷，以及，可能的更好的性能和更广泛的用途---例如之前说过的超量分配），从今天开始，我们看一下Unified Memory的具体使用，使用中的具体注意事项，以及，对它使用的优化处理。后续三个方面的章节将结束本手册的全部内容。 今天首先说了，如何使用cudaMallocManaged()进行Unified Memory的分配，这个函数的用法和作用（原文是语法和语义）上，很接近普通的，之前你所熟悉的常规cudaMalloc()分配。对于计算能力3.X和5.X的设备，也就是具有一代Unified Memory支持的设备，这种分配只能通过这里说的cudaMallocManaged(), 以及，下次章节要详细介绍的__managed__的方式来进行分配；而具有计算能力6.X+的，并且在支持的硬件和软件平台上，其实这里主要说的是POWER系列，例如你能买到的DGX上，具有系统级别的Unified Memory：任何你能从CPU上的正常分配方式， 无论是malloc/new之类的，都将自动变成Unified Memory。 这点还是很诱惑人的，请试想你映射了一个文件作为Unified Memory（命名映射），或者从系统整体的虚拟内存中进行了分配（物理内存+swap分区/文件构成，也就是匿名映射），都可以直接作为Unified Memory来使用。前者能构成很多人梦想中的GPU几乎能访问anything，只要这东西能映射成内存； 后者则在之前我们说的，较高的计算能力和支持的平台上，显存构成了内存的缓存，而现在则加上了内存的缓存/虚拟内存子系统，更多了一级支持。 请注意本章节说了很多操作系统方面的名词，对于非CS出身的用户，我建议自行至少去搜索看一下：什么是堆（Heap）、栈（Stack），以及，虚拟内存方面的东西。避免你在看读本章节的内容的时候感到迷惑。然后本章节具体讲解了一下cudaMallocManaged()的用法，请看代码 cudaMallocManaged(void **devPtr, size_t size, unsigned int flags=0); 这个函数和普通的: cudaMalloc (void **devPtr, size_t size); 非常的接近，请注意这里的uint的flags = 0; 这是一个可选的参数， C++编译器，和部分的C编译器支持可选参数。但是这里的flags我们本章暂时不做解释（可以选择attach host/global的两种分配），后面会说。 在不使用该可选参数的情况下，你会看到（代码main函数中）： 用： cudaMallocManaged(void **devPtr, size_t size); 替换： cudaMalloc(void **devPtr, size_t size); 也就是这种情况下，两者几乎完全一样。只是函数名多了一个尾巴， 你就从普通的global memory的分配，只能GPU用，变成了系统中的CPU和GPU都能用的unified memory了，非常简便。（第三个可选参数的问题，我们下次说） 本文指出，这种简单的替换，因为可以让你直接省略掉后续的复制之类的过程（CPU能直接访问了么！），所以必须遵循一定的并发访问原则（后续章节说明），注意该原则在一代Unified Memory和二代Unified Memory上是不同的。简单的说，一代Unified Memory不允许CPU和GPU并发访问；二代的则可以。 在目前的你已经读取到的章节内容中，你可以简单的认为，GPU在访问的时候，CPU就不能访问，而在后续的章节中，你会逐渐看到，如何在1代Unified Memory，通过流来限定部分Unified Memory的作用范围，从而能部分的让CPU和GPU访问；以及，更进一步的，在二代Unified Memory中，能够自由的让CPU和GPU并发访问。这是一个逐渐的过程。但是在目前的认识内，只要在GPU访问的同时，CPU正确的别碰这些Unified Memory，那么则只要简单的替换掉刚才说过的cudaMalloc函数，你就已经能够基本正常的使用Unified Memory。 这是一个CUDA的显著特点，入门可以只需要5分钟就能用，更好的使用则可能需要50分钟。（但是这5分钟就可以让你快速的使用CUDA并利用它挣钱了）。 然后本章节的末尾强调了，那么在这种只有简单的初步认知的情况下，Unified Memory已经给你带来了巨大的好处： （1）你现在不需要将数据复制来复制去了。很多时候，无辜的复制会降低性能，而又很多时候，全自动的在CPU上进行缺页处理自动传输，性能反而可能会更好。同时用户也减少了必须的代码的编写，增强了程序的正确性，降低了出错可能。 然后第(2)点则是，用户现在不再需要维护两份缓冲区（的指针）了，原本因为复制为复制用户所准备的双份缓冲区，现在已经不需要了，则降低了用户代码逻辑的复杂性，让它变得简洁。所有从事本行业的人可能都知道一点，当代码的逻辑变得复杂后，虽然每一个片段你都可以轻松读懂它，甚至整个项目的代码的每一行。但是从整体上，你将很难去把握它，而今天的这么简单的用途的第(2)点，就可以降低你这方面的成本。很是诱人的。 写一个复杂度为A的软件，可能100个人中有99.9个人都能做到；而当复杂度变成2A的时候，却只有33人了；当复杂度变成4A的时候，可能只有3个人；而当复杂度变成10A后，可能任何常规手段都无法简单的将这个任务完成。Unified Memory的基本使用则可以快速的在逻辑等效的前提下，将复杂度10A的东西，变成只有，例如2.5A这么多。还是很不错的。 最后本章节还说明了一下，设备端的分配（例如new或者malloc，或者动态并行时候的cudaMalloc, 还记得它们吗？），不能当成unified memory使用。 我建议用户回头看一下之前说过在设备端的分配问题（当时说了，这是另外一个设备上的独立的小堆（heap））。如果用户不愿意看一下，也可以直接简单的记住本章节的最后的关于这点的说法和结论。

有不明白的地方，请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖