DAY97：阅读 Stream Attach With Multithreaded Host Programs

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

K.2.2.5. Stream Attach With Multithreaded Host Programs

The primary use for cudaStreamAttachMemAsync() is to enable independent task parallelism using CPU threads. Typically in such a program, a CPU thread creates its own stream for all work that it generates because using CUDA’s NULL stream would cause dependencies between threads.

The default global visibility of managed data to any GPU stream can make it difficult to avoid interactions between CPU threads in a multi-threaded program. FunctioncudaStreamAttachMemAsync() is therefore used to associate a thread’s managed allocations with that thread’s own stream, and the association is typically not changed for the life of the thread.

Such a program would simply add a single call to cudaStreamAttachMemAsync() to use unified memory for its data accesses:

In this example, the allocation-stream association is established just once, and then data is used repeatedly by both the host and device. The result is much simpler code than occurs with explicitly copying data between host and device, although the result is the same.

本文备注/经验分享：

今天的章节继续讲解cudaStreamAttachMemAsync的一个具体案例。如果你已经忘记该函数的主要用途之一是将Unified Memory 1.0上的老平台下的访问能给细分到每个流中，以便取得最大的CPU/GPU的并发访问性，则请回看一下我们上两个章节。 今天的本章节的例子，是非常实用的。在我们日常的多流的应用中（无论你是1个卡N个流，或者M张卡，一共M*N个流），常见有两种应用方式。一个则是CPU上的代码是单线程的，然后CPU反复在不同的流中发布传输或者计算任务；另外一种则是CPU上的代码是多线程的，每个线程只负责自己一个流中的任务。这两种方式都很常见。之前我们手册上的章节往往都是第二种方式。今天的本章节要说的例子则是第一种。 有人可能会问，不是现在都是多核，甚至多路CPU吗？为何还会有单个CPU线程（单个CPU线程显然只能同时在1个CPU核心上运行）控制多个流的方式？这种方式够用么？实际上答案往往出乎你的意料，够用。如同我们曾经多次在论坛说过的一样，一个多卡的GPU系统（例如8卡好了），上面的GPU们都是一些苦力，干重活（kernel）。而系统上的CPU则相当于CEO，只负责领导职务（发布命令或者说任务）和简单的处理即可。因此往往简单的CPU单线程 + GPU多卡多流海量多线程是一种常见的组合。 反正一个核心的CPU你也用不满。足够了。 那么话又说回来了，为何本章节突然给出了一个多线程的CPU的例子？（注意不是GPU多线程）既然CPU上一个核心都用不满，用那么多线程，分布在那么多的核心上干嘛？这里就说说一下CPU多线程的2个主要用途了。 一个用途是大家都知道的，尽量利用现在的机器的CPU的多核心。因为每个核心往往能运行1-2个CPU线程（2线程的需要CPU具有超线程（HT）技术）。你一个8核的机器，往往至少8个CPU线程，才能用满。 则另外一个用途则是，进行逻辑切分，简化逻辑。哪怕一个CPU上的处理，1个核心都用不满，我们依然可以上多个线程的，这样每个线程可以只负责一个确定的任务，编写代码的人往往不需要在多个任务之间跳来跳去，能很大程度的简化逻辑实现。 本章节的例子的主要用途就是第二种。实际上，这种每个CPU线程只负责一个流的做法，非常好用。 我们具体看一下代码： 代码只给出了一个run_task()函数，这是因为我们常见的CPU上的多线程方式，实现的是否时候往往将每个线程所需要执行的共同代码，写成一个函数。就如同这里的一样。 需要注意的是函数原型，常见的Linux和Windows上的pthread_create()和CreateThread()之类的函数和这里给出的run_task()函数的原型不兼容。 你如果需要实际的练习本文例子，还需要将本文的代码单独再包裹一下才能用。 例如pthread_create需要一个void * fun(void *)的形式的函数。你需要处理一下这里的参数和返回值才能用。这个就不多说了。 好了。我们继续往下看，注意这里的函数上来单独建立了一个stream, 给每个线程私用，该stream在线程结束前（也就是本函数返回前），会被cudaStreamDestroy掉。 这样每个stream都是每个线程执行的时候单独所有的。 然后注意这里动态分配了unified memory，通过省略后两个参数的cudaStreamAttachMemAsync, 将分配的unified memory限定为本stream（中的kernel）私有。 注意虽然这里的分配时候的unified memory所拿到的指针/缓冲区首地址，只能在每个线程内部可见，后面也的确用cudaFree释放了，是线程私有的， 但是你必须手工的调用一次cudaStreamAttachMemAsync一次，额外的限定特定的流专用。这是因为无论指针的可见性如何，是否是本线程单独可见，缓冲区默认总是所有的GPU中的流都能用的。 然后注意这里使用的同步方式（等待这个attach的异步操作完成）是用的cudaStreamSynchronize(), 我们之前章节的例子是用的cudaDeviceSynchronize(), 这里这样用是为了尽量减少对其他线程的干扰。我们之前也说过，你可以用设备同步，流同步，事件同步（前提是该事件后面没有其他任务了），这些都可以的。根据需要选用。本例子使用的是流同步。三大同步方式到现在（手册快结束）你应当已经用的非常熟练了吧。 类似的还有下面的同步kernel的时候，也是流同步，这个是常规的，没有unified memory之前我们也这样做，第二次同步就不多说了。 然后完成这些准备工作后，该CPU线程具体的干活部分（从完成attach mem后，到最后的三句释放资源前），将变得非常简单。

for(int i=0; i<N; i++) {

transform<<< 100, 256, 0, stream >>>(in, data, length);

cudaStreamSynchronize(stream);

host_process(data, length); // CPU uses managed data.

convert<<< 100, 256, 0, stream >>>(out, data, length);

}

每个CPU线程都在直接的反复黑用这unified memory（这里用一个for来演示频繁使用），然后CPU等待GPU完成后就立刻正常使用。 看上去是如此的方便，就像Pascal+的Unified Memory 2.0平台那样。 但是实际上，该代码可以安全的在老卡上跑（例如你的K80！例如你手上的老Maxwell卡！）。 这样用这种实例中的方式，我们几乎可以没有太大影响的，在老卡的平台上，成功的方便的多流多CPU线程的独立使用Unified Memory，避免最初的章节中说过的，一旦GPU在忙碌，CPU就不能访问它们的影响。值得借鉴。 如同本章节末尾说的一样，只需要简单的在开头多进行一次attach操作，后面就可以相当简化的反复使用，规避了反复的cudaMemcpy过去，再从显存cudaMemcpy之类的复制回来。 注意本文用了much simpler, 的确如此。简单的多。

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