DAY95:阅读Managing Data Visibility and Concurrent CPU

K.2.2.3. Managing Data Visibility and Concurrent CPU + GPU Access with Streams

Until now it was assumed that for SM architectures before 6.x: 1) any active kernel may use any managed memory, and 2) it was invalid to use managed memory from the CPU while a kernel is active. Here we present a system for finer-grained control of managed memory designed to work on all devices supporting managed memory, including older architectures with concurrentManagedAccess equal to 0.

The CUDA programming model provides streams as a mechanism for programs to indicate dependence and independence among kernel launches. Kernels launched into the same stream are guaranteed to execute consecutively, while kernels launched into different streams are permitted to execute concurrently. Streams describe independence between work items and hence allow potentially greater efficiency through concurrency.

Unified Memory builds upon the stream-independence model by allowing a CUDA program to explicitly associate managed allocations with a CUDA stream. In this way, the programmer indicates the use of data by kernels based on whether they are launched into a specified stream or not. This enables opportunities for concurrency based on program-specific data access patterns. The function to control this behaviour is:

The cudaStreamAttachMemAsync() function associates length bytes of memory starting from ptr with the specified stream. (Currently, length must always be 0 to indicate that the entire region should be attached.) Because of this association, the Unified Memory system allows CPU access to this memory region so long as all operations in stream have completed, regardless of whether other streams are active. In effect, this constrains exclusive ownership of the managed memory region by an active GPU to per-stream activity instead of whole-GPU activity.

Most importantly, if an allocation is not associated with a specific stream, it is visible to all running kernels regardless of their stream. This is the default visibility for a cudaMallocManaged() allocation or a __managed__ variable; hence, the simple-case rule that the CPU may not touch the data while any kernel is running.

By associating an allocation with a specific stream, the program makes a guarantee that only kernels launched into that stream will touch that data. No error checking is performed by the Unified Memory system: it is the programmer’s responsibility to ensure that guarantee is honored.

In addition to allowing greater concurrency, the use of cudaStreamAttachMemAsync() can (and typically does) enable data transfer optimizations within the Unified Memory system that may affect latencies and other overhead.

本文备注/经验分享：

我们从阅读Unified Memory到现在，知道在不支持concurrentManagedAccess的卡和平台上，一旦GPU在忙碌，则CPU就不能访问所有的Unified Memory上的缓冲区、数组、变量等。而本章节则提供了一个函数，可以将Unified Memory切分成多个细小的部分，使得即使在老卡老平台上，也能一定程度的允许CPU和GPU并发访问Unified Memory，很大程度的提高了这些平台上的Unified Memory的使用方便性。 该函数则是本章节介绍的cudaStreamAttachMemAsync(). 注意本章节的文字部分有错误。该函数最后的默认参数是cudaMemAttachSingle, 也就是常数值4， 而不是这里的代码中的第四行的flags = 0 我们具体的来看下本函数。 首先该函数接受一个stream作为参数，这个参数的stream有两个作用，我们稍后说。 其次，该函数接受一段Unified Memory区域作为参数，也就是这里的中间两个参数，分别是该段Unified Memory的首地址（例如你动态的从cudaMallocManaged中得到的地址），和需要被操作的该区域中的数据部分的长度。 如同本章节所说，其中第二个长度信息，目前并不支持将某个分配得到的Unified Memory的再次切分一部分进行操作，只能这部分的整体，所以长度参数必须使用特殊值0. 这里需要注意一下。但即使如此，该函数也允许我们通过多次调用它，分别操作/控制多个分配得到的Unified Memory区域了。 最后的一个重要的参数是刚才提到过的，有问题的flag参数，该参数是可选的（含有默认的等于多少多少字样的，例如这里的=0, 代表可选，虽然这里的0是错误的）。 目前该参数有3种值，后续的章节会说。但这里需要强调的是，默认的值为cudaMemAttachSingle. 在今天章节和后续章节的很大一部分范例代码中，我们都将使用这个默认值（attach single），因为是默认值，很多时候并不具体的写出它，用户在读手册的时候需要注意。 好了，在有了默认标志（attach single）的前提下，该函数的第一个参数同时具有两个作用。 首先我们看到该函数是Async（异步的），则第一个参数stream的第一个作用是，指定了一个流来发布该操作命令， （如果你已经忘记了异步的是什么意思的话，请参考我们之前的章节，这里简单的说一下就是，该操作不等完成，就立刻返回控制权给CPU。CPU可以继续干其他的活，只有当CPU进行了一次和GPU（具体到这里，是和GPU上的一个流）同步后，才能确保该操作的确完成了，这种方式很多时候可以让CPU和GPU在同时忙碌，CUDA大部分以Async结尾的函数都具有这个效果） 好了，第一个stream参数的第一个作用我们说完了（在该stream中发布异步操作命令）。 然后第二个作用则和最后的一个参数有关，在我们刚才修正过的默认的最后一个参数的前提下（attach single的那个），是该操作将会将指定的unified memory区域，和GPU上的该stream进行绑定。 因为你知道我们之前的章节说，说过一个GPU设备上可以建立多个流，例如8个流。 这样如果我们有多个Unified Memory区域，例如也是8个好了，可以分别将每个区域绑定在一个流中，这样实际上将全局的，只要GPU忙碌，CPU就不能碰任何这8个区域中的任何一个区域，改成了只有这8个流中的某个在忙碌，则CPU不能碰这8个流中的某个绑定过的区域，从整体的不能动，细分成了1/8. 这样很大的程度的提高了，在老卡上的Unified Memory的细粒度控制，和更好的CPU/GPU并行访问性，往往能提高性能。 回到本章节的原文说法： “在这种关联下，无论其他流中是否有任务，只有某流中的任务完成，则CPU就能访问某流中绑定的Unified Memory”。 这是该操作的第二个作用，也是最大的作用。（还有其他作用，毕竟最后一个参数可选3种。另外的其他作用我们后续再说） 注意本章节的强调： 等效于，这种操作将全局的GPU上的限制，切分到了每个流中。当我们没有进行这种按流切分的时候，就是我们之前的说法，GPU全局存在一个总限制。 （你可以想象成限制在全局的一个默认流中） 此时的做法就是我们之前说过的，任何一个kernel在进行，则CPU都不能碰（否则会挂）。本章节强调了，默认分配出来的（动态或者静态，还记得这两种分配方式吗？），都是全局限制的，只有当你通过本章节的函数cudaStreamAttachMemAsync细分后，才会取消全局的限制，只限制在一个流中。 注意，如果用户在要求了细分限制，从全局变成了某个流，最后却违背了自己的做法，例如不做这个流中访问该段数据，则会导致未定义的后果。 所以虽然这种做法提高了并发性和细粒度的控制性，但也对程序员带来了更高的要求。如同很多CUDA的特性一样，这种做法是可选的，你可以不使用本章节的函数，依然能够完成任务，但是用了后，带来了更好的性能（不仅仅来自CPU和GPU的并发性，还有数据移动上的其他好处，以后再说），但是也有更大的限制。用户需要注意了。

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