DAY96:阅读Stream Association Examples

Associating data with a stream allows fine-grained control over CPU + GPU concurrency, but what data is visible to which streams must be kept in mind when using devices of compute capability lower than 6.x. Looking at the earlier synchronization example:

Here we explicitly associate y with host accessibility, thus enabling access at all times from the CPU. (As before, note the absence of cudaDeviceSynchronize() before the access.) Accesses to y by the GPU running kernel will now produce undefined results.

Note that associating a variable with a stream does not change the associating of any other variable. E.g. associating x with stream1 does not ensure that only x is accessed by kernels launched in stream1, thus an error is caused by this code:

Note how the access to y will cause an error because, even though x has been associated with a stream, we have told the system nothing about who can see y. The system therefore conservatively assumes that kernel might access it and prevents the CPU from doing so.

本文备注/经验分享：

在昨天的章节里我们引入了cudaStreamAttachMemAsync()函数，用来增强在1代Unified Memory上的并发性和细粒度访问效果。今天的章节则从几个正确和错误的用法中，让你对该函数有一个直观的认识，了解它的用法。 我们将逐个分析一下今天章节给出的范例代码。这些代码都是之前的那个x,y两个变量（用来模拟示两段Unified Memory上的效果），外加新引入的cudaStreamAttachMemAsync, 看看现在的CPU和GPU的访问效果变化。 首先看第一个例子，第一个例子增加了在stream1中，将y变量所在的内存区域，以attach host的方式，进行一次异步的AttachMemAsync操作。 这里有几点需要注意的： 1）点是最后的可选参数，现在被明确的标明了为cudaMemAttachHost, 还记得昨天的说法，强调这里默认的将是AttachSingle, 即和一个单独的GPU上的流进行绑定/关联。今日的章节在明确的要求对y进行attach host操作后，在1代的老Unified Memory平台上，GPU将丧失对它的访问能力，同时CPU将获得随时随刻的访问能力。这样，因为y变量被限定成CPU访问，则在kernel运行期间，y可以安全的被CPU访问，而GPU将只能使用剩下的其他的没有限定的变量/Unified Memory存储器区域，也就是剩下的那个x变量现在可以自由的在GPU上访问，y可以在CPU上自由访问。 本例子额外需要注意的是：这里使用的stream1这个流存在误导性，注意本例有两处使用了该流，一处是cudaStreamAttachMemAsync, 根据昨天的说法，实际上在没有指定attach single的时候（这里指定的是attach host），参数中的流只起到昨日说过的两个作用中的一个，也就是将这个异步操作（AttachMemAsync）在本流中进行命令发布，而和后续的Unified Memory访问的时候，GPU的Kernel使用哪个流无关的。你完全可以将本文的例子中的第一处stream1改成任何其他创建的流，同时在后面的kernel启动的时候，依然能正常的还用stream1. 这并不妨碍。本文混淆的这点，容易造成对初学者的误导。 其次，第一个例子还应当给出一个反面教材。如果我将kernel中的x = 10; 改成y = 20;会如何？kernel会挂掉？还是整个CPU上的进程会直接挂掉？实际上，只是kernel会执行失败，当y被改成和CPU关联后，在老的Unified Memory平台上，试图在kernel里访问它，只是会安全的返回一个错误信息（例如资源不可用之类的），并不会造成全面崩溃。这点将和下面的例子形成鲜明对比（稍后说）。这是第二点。 类似的，因为本文中的第一处stream1只是用来当成一个发布命令的命令队列来用，所以这里的cudaDeviceSynchronize()可以改成任何流同步之类的操作，等该流中的这个异步操作完成即可。都是安全的。 第三点，本文的这个例子，也包括我刚才给出的反例，实际上只对老Unified Memory平台有效。在新的Pascal+的卡，和支持新Unified Memory的平台上（不是Windows一般都无问题），即使是attach到host的一段Unified Memory，例如刚才的反例中的y = 20，实际上依然可以在新Unified Memory下成功执行的。kernel可以即使在这种情况下，依然能成功的将y赋值成20的。用户需要明确这一点（因为考虑到现在大部分人都在用x64 + linux了，应当手头也都是至少Pascal起的新卡了） （你很难看到这里的例子描述的行为。所以需要强调） 类似的，本文还给出了另外一个例子，演示了默认的attach global的效果。注意昨天的说法：一段Unified Memory有3种attach的方式，一种是attach到GPU，也就是attach global，这种是默认的。也就是大家最常见的GPU在干活的时候，CPU一访问就挂的行为。另外一种是使用了cudaStreamAttachMemAsync的默认效果：attach single，将某段unified memory访问限制细分到具体某个流上，只要该流中没有活动的任务，则该unified memory就可以被CPU访问，也就是昨天说的主要内容。这种方式，很大程度的增强了并发性和细粒度访问性。 其次，则是刚才你看到的第一个例子的attach host，关联到CPU的行为，这种GPU不能访问，CPU可以随时访问。 而今天章节的第二个例子将演示，限定x变量到某个细分的流中，剩下的没有被改动过的y变量（依然说默认的attach global的效果），我们来看一下： 首先这里你看到cudaStreamAttachMemAsync的调用变得短了： cudaStreamAttachMemAsync(stream1, &x); 而刚才的例子1中，则是： cudaStreamAttachMemAsync(stream1, &y, 0, cudaMemAttachHost); 这种短的写法使用了后两个默认参数，其中倒数第二个0，只能写成0，这个我们之前说过。而第二个被省略的参数，则是cudaMemAttachSingle, 也就是本例中，要求x和stream1进行关联。 这是第一点需要注意的。 第二个需要注意的则是演示了，每次执行attach mem的操作的时候，只能关联一段unified memory，也就是昨天说的，如果你有多段Unified Memory，多个Unified Memory上的变量，或者数组，需要多次手工操作。 本章节的例2在启动了使用x变量的kernel后（该变量被attach到该kernel启动所在的流中---但是实际上这个并不影响什么，例子不是太好），然后试图在CPU上访问没有被执行过attach mem操作关联过的y，根据我们之前的说法，y此时将保持默认的attach global的经典特性，也就是只要GPU在忙碌，CPU将不能访问。在常规的Linux平台上，该代码很大概率将直接挂掉CPU进程。 （你会看到segmentation fault或者bus error之类的提示，也就是SIGSEGV或者SIGBUS信号会干掉你的进程，壮烈牺牲） 还有其他能举出的各自组合的例子，3种关联方式，多个独立的Unified Memory之间的任意一种组合，都值得举出作为例子。但本章节给出的这两种，算是稍微抛砖引玉了。读者可以自行编写其他组合的例子，实地的跑一下，增加印象。 手册不能写完所有可能，总得给读者留有思考的余地，也总得给其他编CUDA书的人发挥的余地。这两个例子已经很好了.

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