DAY89：阅读Unified Memory Programming

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

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

本文共计1071字，阅读时间30分钟

前情回顾：

DAY79：阅读 Compute Capabilities

DAY80：阅读Compute Capability 3.x

DAY81：阅读Compute Capability 5.x

DAY82：阅读Compute Capability 6.x

DAY83：阅读Compute Capability 7.x

DAY 84:阅读 Driver API和CUDA Context

DAY85:阅读Module

DAY86：阅读Kernel Execution

DAY87：阅读Interoperability between Runtime and Driver APIs

DAY88：阅读CUDA Environment Variables

K. Unified Memory Programming

K.1. Unified Memory Introduction

Unified Memory is a component of the CUDA programming model, first introduced in CUDA 6.0, that defines a managed memory space in which all processors see a single coherent memory image with a common address space.

Note: A processor refers to any independent execution unit with a dedicated MMU. This includes both CPUs and GPUs of any type and architecture.

The underlying system manages data access and locality within a CUDA program without need for explicit memory copy calls. This benefits GPU programming in two primary ways:

• GPU programming is simplified by unifying memory spaces coherently across all GPUs and CPUs in the system and by providing tighter and more straightforward language integration for CUDA programmers.
• Data access speed is maximized by transparently migrating data towards the processor using it.

In simple terms, Unified Memory eliminates the need for explicit data movement via the cudaMemcpy*() routines without the performance penalty incurred by placing all data into zero-copy memory. Data movement, of course, still takes place, so a program’s run time typically does not decrease; Unified Memory instead enables the writing of simpler and more maintainable code.

Unified Memory offers a “single-pointer-to-data” model that is conceptually similar to CUDA’s zero-copy memory. One key difference between the two is that with zero-copy allocations the physical location of memory is pinned in CPU system memory such that a program may have fast or slow access to it depending on where it is being accessed from. Unified Memory, on the other hand, decouples memory and execution spaces so that all data accesses are fast.

The term Unified Memory describes a system that provides memory management services to a wide range of programs, from those targeting the Runtime API down to those using the Virtual ISA (PTX). Part of this system defines the managed memory space that opts in to Unified Memory services.

Managed memory is interoperable and interchangeable with device-specific allocations, such as those created using the cudaMalloc() routine. All CUDA operations that are valid on device memory are also valid on managed memory; the primary difference is that the host portion of a program is able to reference and access the memory as well.

Note: Unified memory is not supported on discrete GPUs attached to Tegra.

K.1.1. System Requirements

Unified Memory has two basic requirements:

• a GPU with SM architecture 3.0 or higher (Kepler class or newer)
• a 64-bit host application and non-embedded operating system (Linux, Windows, macOS)

GPUs with SM architecture 6.x or higher (Pascal class or newer) provide additional Unified Memory features such as on-demand page migration and GPU memory oversubscription that are outlined throughout this document. Note that currently these features are only supported on Linux operating systems. Applications running on Windows (whether in TCC or WDDM mode) or macOS will use the basic Unified Memory model as on pre-6.x architectures even when they are running on hardware with compute capability 6.x or higher. See Data Migration and Coherency for details.

K.1.2. Simplifying GPU Programming

Unification of memory spaces means that there is no longer any need for explicit memory transfers between host and device. Any allocation created in the managed memory space is automatically migrated to where it is needed.

A program allocates managed memory in one of two ways: via the cudaMallocManaged() routine, which is semantically similar to cudaMalloc(); or by defining a global __managed__ variable, which is semantically similar to a __device__ variable. Precise definitions of these are found later in this document.

Note: On supporting platforms with devices of compute capability 6.x, Unified Memory will enable applications to allocate and share data using the default system allocator. This allows the GPU to access the entire system virtual memory without using a special allocator.

The following code examples illustrate how the use of managed memory can change the way in which host code is written. First, a simple program written without the benefit of unified memory:

This first example combines two numbers together on the GPU with a per-thread ID and returns the values in an array. Without managed memory, both host- and device-side storage for the return values is required (host_ret and ret in the example), as is an explicit copy between the two using cudaMemcpy().

Compare this with the Unified Memory version of the program, which allows direct access of GPU data from the host. Notice the cudaMallocManaged() routine, which returns a pointer valid from both host and device code. This allows ret to be used without a separate host_ret copy, greatly simplifying and reducing the size of the program.

Finally, language integration allows direct reference of a GPU-declared __managed__ variable and simplifies a program further when global variables are used.

Note the absence of explicit cudaMemcpy() commands and the fact that the return array ret is visible on both CPU and GPU.

It is worth a comment on the synchronization between host and device. Notice how in the non-managed example, the synchronous cudaMemcpy() routine is used both to synchronize the kernel (that is, to wait for it to finish running), and to transfer the data to the host. The Unified Memory examples do not call cudaMemcpy() and so require an explicit cudaDeviceSynchronize()before the host program can safely use the output from the GPU.

Note: An alternative here would be to set the environment variable CUDA_LAUNCH_BLOCKING=1, ensuring that all kernel launches complete synchronously. This simplifies the code by eliminating all explicit synchronization, but obviously has broader impact on execution behavior as a whole.

本文备注/经验分享：

如同我们之前曾经说过的一样，CUDA中存在很多特性，其实是可选的： 不使用它，你的代码照样可以编写，kernel可以照常运行；而使用了它，则在很多方面可以进行简化，你的CUDA生涯可以过的更好一点。 今天开始的Unified Memory，则也可以算是成这种。但需要说明的是，实际上存在两代的Unified Memory： 从Kepler开始的一代，这一代基本上算是一个可选特性，只是后面地方方便了一点（也包含Maxwell）。 而从Pascal+开始，Unified Memory则算是第二代了，功能得到了很大的增强，很多时候，已经不能将它看成可选特性了，有些场合的应用，必须通过Unified Memory才能完成（或者说，才能实际点的完成）。还记得我们在之前的计算能力章节（6.X）时候说过吗？Pascal基本上对Maxwell没有太多的变化，只是制程先进了。当时说过，99%完全一样的指令集，它们是一样的。那么Unified Memory，则是Pascal除了指令增强（增加了几条FP16和INT8的指令），NVENC增强之外的，第三个Pascal增强的版本。我们将陆续为用户介绍这两代Unified Memory的功能。 回到本章节，本章节上去的段落介绍了，Unified Memory是一种逻辑上的概念，它既不是显存，也不是系统的内存，而是一种系统上的任何处理器（CPU或者GPU）都可以访问，并能保证一致性的逻辑上的一种虚拟的存储器。这里还稍微略带说明了一点MMU的概念。 对于并非CS出身的用户，我这里简单介绍一下，MMU是一种虚拟化内存，或者你可以简单的认为能将内存访问功能扩展的一种硬件单元，现在的CPU和GPU，都集成了它。MMU也是现代的操作系统所存在的一个基本性依赖。（具体详情用户可以自行搜索）。简单的说，通过CPU和GPU各自内部集成的MMU单元，NV的CUDA为我们提供了一种，构建在显存和内存之上的，一种虚拟的存储器。这种虚拟的存储器就叫Unified Memory。 实现这个其实并不是很容易，竞争对手A家的多个Linux驱动版本中，分别撤销了这个功能（OCL的SVM Buffer），这充分体现了NV是一家软件功能的技术能力。用户只需要知道它是一种比显存和内存的直接使用都高级的存储器即可。会用就可以了，其实并不需要知道太多的技术实现细节。 （为何是虚拟的，这个道理很简单：你拆开机箱能看到内存条，你拆掉显卡散热片，能看到显存芯片。这两个是真的。显然并不存在一种叫Unified Memory的特定的硬件芯片，所以必然是用这两者构建而成的虚拟存储器）。那么我们为何要使用它呢？在之前的章节用户已经会分配显存了，会通过kernel里面的指针读写普通线性显存，也会通过kernel里面的纹理或者Surface来访问CUDA Array了， 为何今天用户还要麻烦的去学习一种新的叫Unified Memory的东西？类似的，用户早就熟悉了内存的使用，会在CPU上正常的使用它们，例如打开一个文件，将内容读取到内存，为何用户还要去继续折腾这些？ 今天章节的中间段落则描述了使用Unified Memory的诸多好处，用户可以看一下是否值得学习这种新的存储器： （1）使用更加简单。在用户学会了Unified Memory后，可以直接忘记cudaMemcp*()这种手工的移动，无论是从显存移动到内存，还是从内存移动到显存。Unified Memory给你提供的这种虚拟的存储器，可以直接从GPU上进行访问，也可以直接从CPU上进行访问。 任何一种处理器，均不需要用户将数据移动到它自己能直接访问到的地方，进行访问。例如说，GPU不需要用户将数据跨PCI-E传输到显存，再访问。 CPU也不再需要用户将数据手工跨PCI-E传输回来到内存，才能访问。 Unified Memory提供了一种方便的，能任何处理器都能直接访问到的存储器。这能解放不少懒人用户。甚至还能减少出错，因为用户不再需要一种手工的数据传输过程，期间容易写错的代码也被取消掉了。这是第一点好处。是否有点心动？ （2）第二点好处则是，可能会比你的手工数据移动，提供更好的性能。 底层的Unified Memory实现，可能会自动将一部分数据，放置到某个存储器更近的位置（例如部分放置到某卡的显存，部分放置到内存），这种自动的就近数据存放，有可能提升性能。 换句话说，不仅仅用起来更简单，而且用起来速度可能会更快。这样Unified Memory的这两个基本点，在一个具体的应用了GPU系统的单位，前者降低了人力成本的支出；后者则在不增加预算的情况下，提升了性能；或者说同样的性能的情况下，减少了相关的预算。 还是一定程度很有诱惑力的。我们继续看后续段落。 如同任何技术都不是凭空产生的一样，Unified Memory其实脱胎于以前的CUDA中就存在的一些旧的技术。主要有两点，一点是从Fermi起引入的Generic Addressing和UVA，还记得它们两个吗？前者是将一个卡内部的多种存储器（Global Memory，Local Memory, Shared Memory), 合并到同一个地址空间，可以让你用普通的一个指针打天下。后者则是可以将一个GPU系统内（例如一个4卡的双路CPU系统），将某CUDA应用所能访问到的所有处理器资源的存储器（所有卡的显存，和CPU的内存），都统一到同一个地址空间。 （具体细节请看我们之前的章节）。 则今天的Unified Memory，则在这种统一的地址编址的情况下，增加了刚才说的，能各个处理器（多个CPU和多个GPU），能互相访问这些被统一在一起编址的存储器的能力。就如同两家公司在合并，以前等于名义上进行了更名，员工还是各自的；现在则是彻底的融合成和一家公司，员工可以自由的流动。 读到这里，有些用户可以说，以前不是有zero-copy内存么？我可以分配一段内存，然后映射给显卡用，这些这种映射的内存，显然CPU依然还能用，也增加了GPU能用的功能，这不就是Unified Memory么？有什么区别？

今天章节的后续段落，则说明了它和Zero-Copy Memory（映射的内存成的Global memory)的具体异同： （1）相同的是，它们都提供了一种统一的，能被CPU和GPU都访问到的存储器（或者说地址/指针）。GPU在访它们的时候，都不需要手工的通过cudaMemcpy*进行数据传输/复制，能就地访问（zero-copy特性）。 （2）不同的是，之前的实现是总是用的内存作为后备的存储介质，例如一个4卡的系统，4张卡访问这种zero-copy存储器的时候，将总是4张卡在跨过PCI-E在访问内存，可能会造成性能上的瓶颈。而Unified Memory则能将数据的存放位置，放置在一个高效的地方，例如经常被CPU访问的，放置在内存，经常被卡1访问的，放置在卡1的显存，经常被卡2访问的，放置在卡2的显存。甚至根据情况，如果使用了很大量的数据，可能部分数据在卡2上，部分数据在卡1上，两个卡之间进行自动的P2P Access。甚至还可能其中部分非常少量使用的数据部分，放置在内存---毕竟后者这个容量巨大，虽然速度的确慢了很多。 所以这比传统的Zero-Copy Memory（纯内存后备的），性能上能好很多。特别的，Pascal起，具有增强的Unified Memory，允许显卡在使用了Unified Memory的情况下，进行超量分配，什么叫超量分配？例如说，我有一张6GB的卡（GTX1060 6GB版），我需要处理一段数据，大小为8GB，这超过了我的显存容量。我可以考虑直接购买一张更大的卡显存的卡，但这会增加硬件成本。我也可以考虑应用更加复杂的算法，每次移动部分数据到显存里，处理完毕后，移动出来，然后将下一片数据移动进行。 但这会增加人工实现的成本，甚至有些算法无法有这种变通的处理办法。而有了Unified Memory后，在Pascal上，满足一定的条件下，可以直接分配8GB的Unified Memory，Unified Memory将自动使用显存（6GB）缓冲这8GB的数据，从而达到逻辑上，你有8GB的这么大的等效缓存；同时只有最频繁使用的6GB的部分，被放置在显存上，其他被少量访问的部分，放置在内存上。这样允许你以将近显存的访问速度（和这种虚拟的cache的具体效率有关），得到比自己显存大很多的容量。这比以前只能使用zero-copy内存，8GB全部在内存上，频繁的通过PCI-E传输要快的多。这是一个很显著的优势。 回到今天的章节的后半部分。 你看有这么多好的优点，是否开始感到心动？至少那些偷懒的用户，也应当开始心动了。但是今天后面的段落说，想用它们需要有一定的条件。 （1）只能在64-bit系统下用。同时Windows上只有兼容性支持（性能降低到类似zero-copy memory的程度，效果很差），几乎等于没有。也就是说，系统需要你使用Linux（64-bit）才可以有效的利用。 （2）至少需要一张计算能力3.0+的卡，6.0+的更好（后者具有更强大的Unified Memory能力和效果，这个后续我们会说到）。但是考虑到大部分的单位都在使用Linux + 较新的卡，所以这一般不是个问题。甚至还有更强大的IBM POWER系统，NVLink能直接服务CPU<--->GPU间（内存和显存间）的数据移动（如果Unified Memory自动管理的时候认为需要存在这种移动，下同），或者NVLink能服务GPU<--->GPU间的传输。具有更好的效果，此外在这种系统上，Unified Memory能直接访问任何存储器，例如我直接内存映射一个文件（将文件的映射成内存的一部分），此时GPU将具有直接访问磁盘上的文件的能力，很是强大 （这部分信息来自每次NV宣传Unified Memory时候的宣传内存。然后我们并没有真正的检验过它。这里提到是因为NV说的多，同时的确让人心动）。 此外，本文还说，在TX2这种平台上，额外的插入一张N卡，则Unified Memory并不支持这张额外插入的独立N卡的（有人这样做么？）这点同样未经验证。 这是今天的大致内容，对Unified Memory的基本介绍章节。

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

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