DAY 94:阅读Explicit Synchronization and Logical GPU Activity

K.2.2.2. Explicit Synchronization and Logical GPU Activity

Note that explicit synchronization is required even if kernel runs quickly and finishes before the CPU touches y in the above example. Unified Memory uses logical activity to determine whether the GPU is idle. This aligns with the CUDA programming model, which specifies that a kernel can run at any time following a launch and is not guaranteed to have finished until the host issues a synchronization call.

Any function call that logically guarantees the GPU completes its work is valid. This includes cudaDeviceSynchronize(); cudaStreamSynchronize() and cudaStreamQuery() (provided it returns cudaSuccess and not cudaErrorNotReady) where the specified stream is the only stream still executing on the GPU; cudaEventSynchronize() and cudaEventQuery() in cases where the specified event is not followed by any device work; as well as uses of cudaMemcpy() and cudaMemset() that are documented as being fully synchronous with respect to the host.

Dependencies created between streams will be followed to infer completion of other streams by synchronizing on a stream or event. Dependencies can be created viacudaStreamWaitEvent() or implicitly when using the default (NULL) stream.

It is legal for the CPU to access managed data from within a stream callback, provided no other stream that could potentially be accessing managed data is active on the GPU. In addition, a callback that is not followed by any device work can be used for synchronization: for example, by signaling a condition variable from inside the callback; otherwise, CPU access is valid only for the duration of the callback(s).

There are several important points of note:

• It is always permitted for the CPU to access non-managed zero-copy data while the GPU is active.
• The GPU is considered active when it is running any kernel, even if that kernel does not make use of managed data. If a kernel might use data, then access is forbidden, unless device property concurrentManagedAccess is 1.
• There are no constraints on concurrent inter-GPU access of managed memory, other than those that apply to multi-GPU access of non-managed memory.
• There are no constraints on concurrent GPU kernels accessing managed data.

Note how the last point allows for races between GPU kernels, as is currently the case for non-managed GPU memory. As mentioned previously, managed memory functions identically to non-managed memory from the perspective of the GPU. The following code example illustrates these points:

本文备注/经验分享：

今天的章节继续昨天1代Unified Memory（和部分特殊情况）的，不能CPU和GPU同时访问Unified Memory的限制说的。 昨天简单的提到说，在这种老的1代的Unified Memory的平台上，　CPU必须等待GPU完成kernel后，才能继续访问。而今天的内容则详细列举了什么是GPU完成了它身上的任务。几乎所有的可能情况在本章节都进行了说明。用户基本上应当认为本章节列出的情况是照顾到了所有的方面（Exhaustive） 我来说一下本章节的具体内容。 首先说，CPU上的代码，必须先明确的确定GPU已经完成了它身上的活后，才能访问，例如昨天的例子中，kernel<<<>>>()启动后，CPU立刻要求了一条cudaDeviceSynchronize()同步操作，才能安全的访问Unified Memory。这种明确性，可以通过多种方法来确定（本章节马上就要逐个说）。而除了这些方法外的任何假设，都是不安全的。例如说，本章节上去说了，昨天有一个只有1条语句的kernel（x = 10），那么我在启动kernel后，应当这kernel执行的超快就结束了吧，我估计Unified Memory几乎是立等可用的，我马上Host访问它行不？不行。 那我估计我的kernel最多执行1ms，我等10ms后，立刻访问行不？不行。 也就是说，任何假设性的操作都是不可以的。必须从逻辑上保证GPU上的kernel的确已经完成了，这样才行。 而本章节后续说明了几乎所有的，从逻辑上能保证这点的做法。我们逐个来看看： 首先本章节列举了三大同步的方法，也就是设备同步，流同步，和事件同步。 这三大同步的方式实际上在我们之前的章节中都说过， 首先是设备同步，cudaDeviceSynchronize()，这个没有什么可说的，大家喜闻乐见。类似的还有cudaStreamSynchronize()如果当前设备上发布任务的只有这一个流，果断没有问题。 以及，对一个流中要求进行Record操作的Event，通过cudaEventSynchronize()对它进行同步，并且该event后续无其他操作了，也可以。 这些都是我们之前曾经说过，或者练习过的同步方式，没有什么可说的，欢迎查看之前的章节。 注意这里的Event同步还可以跨设备进行同步操作，也需要注意，之前也说过这些。 然后还有就是查询类的操作，cudaStreamQuery()和cudaEventQuery(), 如果查询到相关的流或者时间的状态为完成态（非cudaErrorNotReady, 这是进行中的状态），则在类似的能保证只有1个流中有任务，或者该event无后续跟随的任务，则显然也能知道设备上已经完成了执行了， 这些都是从逻辑上很显然的东西。通过使用这些逻辑上显然的保证的操作，你方可安全的在后续的CPU代码中，在老的1代的Unified Memory上用CPU继续使用Unified Memory。类似的，之前说过的用经典默认流（非每个Host线程单独所有的“默认流”---后者实际上只是一个普通流，详情见我们之前的章节），对这种经典默认流的同步操作，也可以，因为这等效于正常的设备同步。 本章节强调了这点，但实际上我们之前在多流的章节都说过，这里复习一下也挺好。 以及，如果程序使用了多流，通过在流间用event的插入互相维持逻辑顺序（请详情参考我们之前的多个任务，以graph的方式插入到不同的流中，用event再维护它们之间的逻辑性），则在该图（graph）的最后逻辑完成位置的节点（所对应的流）单独进行一次同步，也是可以安全后续CPU使用的。换句话说，你现在能用的所有安全的同步方式，都可以用。 本章节几乎穷尽性的列举出来了他们，也算是个好事(避免你没有看过之前的章节，或者忘记了）。然后后续段落还说了在流的callback中（回调，也是之前章节的内容，一个任务完成后，可以不用自己的代码查询或者同步，而是要求CUDA Runtime自动调用一段你要求的CPU代码），这种在本章节中也说了，此时GPU上无其他任务在执行，也将会是安全的。你看到，老的Unified Memory的限制还是很多的，但本章节的叙述还是很全面的。这样按照所有方面来，几乎可以平安的使用。（当然，能直接上Pascal+的卡更好，这种卡在64-bit的Linux下，几乎无任何限制，本章节和昨日的章节，几乎可以完全忘记） 然后本章节再次强调了几个要点： （1）CPU能访问到的其他存储器还有zero-copy memory(这是内存），和常规内存（例如你的malloc出来的），这些都是原本CPU能访问的东西。Unified Memory，即使是老的1代，也并未对原本CPU就有的能力进行任何限制。这些都是可以在GPU忙碌的时候，CPU依然可以保持访问的。 （2）再次强调，在1代的老Unified Memory上（无concurrentManagedAccess)的时候，只要GPU在执行kernel，无论该kernel是否用到具体的某段unified memory缓冲区，则CPU都不能碰。这个实际上是昨天说过的内容，手册又重复了。 （3）Unified Memory在GPU之间的访问，无限制。这些限制只是对CPU来说的。实际上，用户应该知道，老的卡的GPU无缺页异常处理能力，所以必须在GPU执行kernel前，就将数据移动到GPU能访问到的三大位置上（普通内存映射成的global memory, 本卡的显存构成的global memory, 和对方卡通过P2P访问的显存），因为GPU将无法再次缺页处理，将他们移动或者复制回来，所以只要记住这一条，上面的所有注意限制事项都可以自行的推导出来。 万变不离其宗。 最后本章节还有一个演示了几点的例子，可以自行看一下：

这例子演示了几个注意事项的点。 Point 1点，演示了CPU可以照常干它本来能干的事情。 点2则说明了，GPU在忙碌的时候，CPU不能访问。否则会挂掉（Linux这里的segmentation fault, 和Windows上的弹出的内存访问一场窗口挂掉） 点3则说明，该GPU上的其他流中的kernel，并不限制访问（因为刚才说了，限制只是在老的Unified Memory上的CPU对GPU来说的）。 类似的，点4说明，其他卡上（第二张卡，卡1，注意卡的编号从0开始）的kernel，也可以正常访问（因为第二张卡也不是CPU。两张卡和1张卡上的2个流道理类似）。 大致这些情况。