DAY90：阅读Data Migration and Coherency

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

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

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

K.1.3. Data Migration and Coherency

Unified Memory attempts to optimize memory performance by migrating data towards the device where it is being accessed (that is, moving data to host memory if the CPU is accessing it and to device memory if the GPU will access it). Data migration is fundamental to Unified Memory, but is transparent to a program. The system will try to place data in the location where it can most efficiently be accessed without violating coherency.

The physical location of data is invisible to a program and may be changed at any time, but accesses to the data’s virtual address will remain valid and coherent from any processor regardless of locality. Note that maintaining coherence is the primary requirement, ahead of performance; within the constraints of the host operating system, the system is permitted to either fail accesses or move data in order to maintain global coherence between processors.

GPU architectures of compute capability lower than 6.x do not support fine-grained movement of the managed data to GPU on-demand. Whenever a GPU kernel is launched all managed memory generally has to be transfered to GPU memory to avoid faulting on memory access. With compute capability 6.x a new GPU page faulting mechanism is introduced that provides more seamless Unified Memory functionality. Combined with the system-wide virtual address space, page faulting provides several benefits. First, page faulting means that the CUDA system software doesn’t need to synchronize all managed memory allocations to the GPU before each kernel launch. If a kernel running on the GPU accesses a page that is not resident in its memory, it faults, allowing the page to be automatically migrated to the GPU memory on-demand. Alternatively, the page may be mapped into the GPU address space for access over the PCIe or NVLink interconnects (mapping on access can sometimes be faster than migration). Note that Unified Memory is system-wide: GPUs (and CPUs) can fault on and migrate memory pages either from CPU memory or from the memory of other GPUs in the system.

K.1.4. GPU Memory Oversubscription

Devices of compute capability lower than 6.x cannot allocate more managed memory than the physical size of GPU memory.

Devices of compute capability 6.x extend addressing mode to support 49-bit virtual addressing. This is large enough to cover the 48-bit virtual address spaces of modern CPUs, as well as the GPU’s own memory. The large virtual address space and page faulting capability enable applications to access the entire system virtual memory, not limited by the physical memory size of any one processor. This means that applications can oversubscribe the memory system: in other words they can allocate, access, and share arrays larger than the total physical capacity of the system, enabling out-of-core processing of very large datasets. cudaMallocManaged will not run out of memory as long as there is enough system memory available for the allocation.

K.1.5. Multi-GPU Support

For devices of compute capability lower than 6.x managed memory allocation behaves identically to unmanaged memory allocated using cudaMalloc(): the current active device is the home for the physical allocation, and all other GPUs receive peer mappings to the memory. This means that other GPUs in the system will access the memory at reduced bandwidth over the PCIe bus. Note that if peer mappings are not supported between the GPUs in the system, then the managed memory pages are placed in CPU system memory (“zero-copy” memory), and all GPUs will experience PCIe bandwidth restrictions. See Managed Memory with Multi-GPU Programs on pre-6.x Architectures for details.

Managed allocations on systems with devices of compute capability 6.x are visible to all GPUs and can migrate to any processor on-demand.

本文备注/经验分享：

今天的章节则继续介绍Unified Memory的一些底层细节，以便让用户知道这种便利从何而来。此外，今天后两段落还将介绍新老架构上，两代Unified Memory的优缺点（3.x/5.x的Kepler/Maxwell上的老Unified Memory，和Pascal+上的新Unified Memory）。

首先，Unified Memory提供了全自动的数据移动：在维护了系统内部的访存一致性(Coherence)的前提下，Unified Memory将通过全自动的数据移动（或者数据映射），来显示各个存储器的访存效果最优化。例如说，在没有全面的NVLink的机器上，CPU访存可能需要将数据移动到内存中，然后访问。而在有全面的CPU<--->GPU的NVLink的机器上（例如某POWER？），CPU可能会直接通过NVLink访问某GPU的显存。这些都是有可能的。请注意本章节强调了“一致性”是Unified Memory最重要的目标，不能出现为了Unified Memory的方便性，而让CPU或者GPU访问到错误的，或者不一致的数据。

按照本章节的原话说，一致性是要比性能更重要，这也是显然的。再快的实现，一旦结果是错误的，将无意义。请注意非CS出身的用户，一致(coherency)的问题在多个处理器同时能访存的系统上都存在，

建议自行搜索了解相关信息.

此外，在有了一致性保证的基础上，Unified Memory所提供的全自动的数据迁移服务，是对用户程序（或者说代码）透明的，用户不需要手工的写怎么去做，这种迁移就将自动的发生，将自动的将属于放置到某处理器所能访到的最佳位置（例如某GPU自己的显存，也可能是其他地方），而不需要用户去操心。这样，有了一致性和全自动的数据移动，就构建成了Unified Memory所提供的这种所有的CPU和GPU都能访问到的，虚拟的，统一的存储器结构---它的方便性昨天你已经看到了，现在你知道它是大致怎么实现的了。

然后今天章节还介绍一下老Unified Memory的缺点，请注意老Unified Memory所能提供的便利和数据优化位置存储，自动迁移，依然都是有的，这里说的缺点，只是说老一代的卡（精确的说，两代，Kepler和Maxwell）具有不好的方面。主要缺点在于，老卡的MMU（内存管理单元，昨天的章节已经让你搜索它的基本功能了，不要让名字迷惑---这很大程度的是一种内存虚拟化单元），不具有细粒度（例如按照页面大小）的Page Fault能力。这里又引入了一个新的操作系统课程名词，缺页异常/错误（Page Fault），

对于非CS专业的用户来说，可以简单的理解成，这是实现现代的、高级的虚拟内存服务的基础，（具体细节你需要搜索一下，然后阅读一下，这是一个大话题），

老鸟玩家以前在DOS时代玩游戏见到的DOS/4GW，还是现在的Windows/Linux上提供的虚拟内存/交换文件/交换分区，都依赖于他。这不过这些是CPU上的。你看缺页异常处理能力，在CPU上能干这么多的事情。我们今天章节里面遇到的Page Fault Handling能力，则是对GPU说的。老一代的卡不具有精细的缺页异常处理能力，所以在老的Unified Memory上，在kernel启动前，要么（1）它所需要访存的数据只能提前被整体迁移到自己的显存上准备好，要么（2）能在支持P2P Access的平台上，迁移到另外一张卡的显存上，要么（3）存放在内存上，使用类似的Zero-Copy Memory的访问（内存映射的显存的方式）访问。但是无论这三种的哪一种方式为你提供Unified Memory的底层存储，都存在问题。如果提前将数据全部移动到卡的显存上，则显然你的kernel能用的Unified Memory大小无法超过当前卡的显存容量，同时这还增加了数据整体移动的成本。这是显然的。

而如果将Unified Memory的数据放到伙伴卡的显存上，通过P2P Access访问，则还存在访问速度慢（老卡只能用PCI-E走数据访问伙伴卡的显存），甚至有些系统不能支持P2P Access的情况（例如：当前芯片组或者CPU集成的PCI-E交换机不能支持的情况下）。

而放置到内存，则多卡访问的时候内存瓶颈或者PCI-E瓶颈都可能是个性能上的灾难。所以你看，虽然老卡（Kepler/Maxwell）上能用，但是还存在诸多缺点的。用户应当认为，老卡上的Unified Memory，在适当轻量使用的情况下，Just Works，但性能不会太好。不仅仅如此，老卡因为需要提前数据准备好位置（在Kernel启动前），还会导致CPU和GPU无法同时访问数据的情况，请考虑kernel正在运行中，数据在显存里，kernel还没结束，突然CPU说，我想访问一下里面的数据。。。这就尴尬了（这种情况下在老卡上会直接挂掉）。不过幸好时代在进步，Pascal起，虽然作为通用计算并无太大的变化（增加了几条通用的FP16或者INT8指令），但是在Unified Memory这种CUDA的辅助特性上，则进行了很强大的增强。6.X+的GPU的Unified Memory进化到了第二代，上面所遇到的所有问题难点都消失了。新卡的MMU单元，具有完善的缺页异常和硬件/Driver处理能力。

数据现在可以任意存在在内存，自己卡的显存（你即将运行kernel的卡），或者伙伴卡的显存上，

kernel启动前，不需要关心数据具体在哪里，只要kernel在执行过程中，SM发现数据不在显存，它会动态的细粒度的将数据访问到，或者迁移到自己的显存访问到。并不需要整体在启动前迁移好或者准备好数据。刚才你已经知道，哪怕是全自动的迁移或者移动/复制数据，这成本是很大的。例如某kernel需要使用1GB的缓冲区，但是实际上该kernel某次运行只用了里面的256MB，但是具体哪256MB的数据和问题的动态变化有关，不能提前知道，则之前老卡不仅仅需要将这256MB自动的复制或者移动到自己的显存，剩下的1GB的768MB数据也都需要被整体移动，这浪费了3/4的传输。现在的新卡上的新MMU的细粒度缺页能力，改善了这一点。同时，之前因为往往需要数据整体移动，导致Unified Memory不能超过自己的显存容量的问题也解决了，

因为随时可以细粒度的小部分小部分的移动，某卡完全可以自己，例如3GB的显存，使用6GB的容量，反正又不需要一次性的移动完数据，我用到哪些移动哪些就是了。

（这是GPU硬件和显卡驱动驱动，显卡驱动聪明） Unified Memory的改进是硬件和驱动上的，kernel不需要改动。这样，就带来了今天章节说的，Pascal+上的超量分配能力： 我一张只有3GB显存的卡，在使用了Unified Memory的情况下，我可以分配8GB的大小，具体细节可以看一下本章节，这里我简单的介绍一下。 在这种情况下，8GB的实际总数量可能分布在内存中一部分，伙伴卡中一部分，以及，本卡中一部分。而实际上，本卡中的具体某时某刻，里面的3GB，是总8GB的具体哪一部分，则是动态变化的。所以实际上，这时候，在超量分配的情况下，本卡的小容量显存，是总的大容量的Unified Memory的一个动态的缓冲，你可以认为此时，本卡的3GB显存变成了, 整个GPU系统（例如一个双路CPU+4卡GPU的服务器）的缓存，根据你已经学过的计算机课程（好吧，似乎没有照顾非CS专业出身的用户。。）， 此时等效的，将变成，在增强的二代Unified Memory的情况下，此时一张卡的显存将变成整体系统的大容量存储器的一个缓存。例如，你可以看成是一个3GB的L3 Cache。此时根据缓存的一定特性，理想状态下（注意是理想状态，和访存的模型有关），你等效于能最好情况以将近这3GB显存的访问速度，使用将近整个系统的大容量。还是非常诱人的特性。这也是竞争对手A家从Vega卡还是（4代GCN），所提供的特性， 竞争对手A家在这个情况下，叫自己的HBM显存为“High Bandwidth Cache”，可见一斑。实际上，这个特性非常重要（注意，Windows下用不了）。历年来NV的GTC总是会宣传一下。 此图演示了CPU在使用大容量存储器，老卡（K40）在使用较大容量的显存，以及，新卡（P100）在NVLink和PCI-E，以及附带的Unified Memory的Hints操作下（这个我们后续章节说），的性能和应用情况。 首先大家看到蓝色的是CPU，该U在使用从1.4GB到58.6GB的working set（即数据工作集，具体概念请参考操作系统课程）的情况下，都能正常（内存较大么）工作，但是性能也是最低的（因为CPU最慢）。而K40（Kepler，老卡，1代Unified Memory）则在自己的显存容量限制内，能正常工作。而到了P100（新卡，2代Unified Memory)，则出现了较好的现象（橙色，综合，绿色），它能在自己的显存容量限制内工作（16GB，P100），性能不错。 而随着问题的规模的增加，例如数据增加到28.9到58.6这么大，超出了P100的16GB显存容量限制，P100依然成功通过新的Unified Memory超量分配机制，正常的运行了kernel。此时依然比直接能用大内存的CPU，还是快很多的。这是一个很迷人的特性。也解决了之前的一个难题：内存容量大，但是CPU慢；显存容量小，但是GPU快。如今通过Unified Memory在新卡（P100）上的超量分配能力，你使用的Unified Memory容量大如内存，却依然能发挥GPU的高性能。

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

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